From 5458cc794168f7dc96390f9e58291ff146da67bd Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Tue, 16 Jun 2026 01:27:05 +0530 Subject: [PATCH 1/4] extended gpu support for selu --- core/inc/SOFIE/ROperator.hxx | 3 ++- core/inc/SOFIE/ROperator_Selu.hxx | 40 +++++++++++++++++++++++++++++++ core/src/RModel_ALPAKA.cxx | 3 ++- 3 files changed, 44 insertions(+), 2 deletions(-) diff --git a/core/inc/SOFIE/ROperator.hxx b/core/inc/SOFIE/ROperator.hxx index f6a8d43e..9a9eb209 100644 --- a/core/inc/SOFIE/ROperator.hxx +++ b/core/inc/SOFIE/ROperator.hxx @@ -39,7 +39,8 @@ enum class OperatorKind { UNARY_ABS=23, CLIP=24, NOT=25, - POOL=26 + POOL=26, + SELU=27 }; inline const char* toString(OperatorKind kind) { diff --git a/core/inc/SOFIE/ROperator_Selu.hxx b/core/inc/SOFIE/ROperator_Selu.hxx index 5bec42cc..3db298d0 100644 --- a/core/inc/SOFIE/ROperator_Selu.hxx +++ b/core/inc/SOFIE/ROperator_Selu.hxx @@ -25,6 +25,7 @@ public: fNX(UTILITY::Clean_name(nameX)), fNY(UTILITY::Clean_name(nameY)){ fInputTensorNames = { fNX }; fOutputTensorNames = { fNY }; + fKind = OperatorKind::SELU; } std::vector TypeInference(std::vector input) override { @@ -59,6 +60,45 @@ public: } std::vector GetStdLibs() override { return { std::string("cmath") };} + + std::string Generate_GPU_Kernel_ALPAKA(std::string /*opName*/) override { + std::string op; + op = "\n//---- SELU_KERNEL_ALPAKA//\n"; + op += "struct SeluKernel {\n"; + op += SP + "template\n"; + op += SP + "ALPAKA_FN_ACC void operator()(TAcc const& acc, T const* __restrict__ data, T* __restrict__ out, std::size_t numElements) const {\n"; + op += SP + SP + "const auto idx = alpaka::getIdx(acc)[0];\n"; + op += SP + SP + "if (idx < numElements) {\n"; + op += SP + SP + SP + "T x = data[idx];\n"; + op += SP + SP + SP + "T inner = T(1.6732632423543772848170429916717) * (exp(x) - T(1));\n"; + op += SP + SP + SP + "out[idx] = T(1.0507009873554804934193349852946) * ((x > T(0) ? x : T(0)) + (inner < T(0) ? inner : T(0)));\n"; + op += SP + SP + "}\n"; + op += SP + "}\n"; + op += "};\n"; + return op; + } + + std::string Generate_GPU_Kernel_Definitions_ALPAKA(std::string /*opName*/) override { + return SP + "SeluKernel seluKernel;\n"; + } + + std::string Generate_GPU_ALPAKA(std::string OpName) override { + OpName = "op_" + OpName; + if (fShape.empty()) { + throw std::runtime_error("SOFIE Selu called to Generate_GPU_ALPAKA without being initialized"); + } + std::stringstream out; + std::string length = ConvertDimShapeToLength(fShape); + out << "\n//------ SELU_GPU_ALPAKA\n"; + out << SP << "auto const elementsPerThread_" << fNX << " = Vec::all(static_cast(1));\n"; + out << SP << "auto const elementsPerGrid_" << fNX << " = Vec::all(Idx{" << length << "});\n"; + out << SP << "auto const workDiv_" << fNX << " = sofie_workdiv(elementsPerGrid_" << fNX << ");\n"; + out << SP << "auto task_" << OpName << " = alpaka::createTaskKernel(workDiv_" << fNX + << ", seluKernel, alpaka::getPtrNative(deviceBuf_" << fNX + << "), alpaka::getPtrNative(deviceBuf_" << fNY << "), static_cast(" << length << "));\n"; + out << SP << "alpaka::enqueue(queue, task_" << OpName << ");\n"; + return out.str(); + } }; }//SOFIE diff --git a/core/src/RModel_ALPAKA.cxx b/core/src/RModel_ALPAKA.cxx index 842d3ff8..429b732e 100644 --- a/core/src/RModel_ALPAKA.cxx +++ b/core/src/RModel_ALPAKA.cxx @@ -576,7 +576,8 @@ void RModel::GenerateSessionCode_GPU_ALPAKA() { SOFIE::OperatorKind::UNARY_SIN, SOFIE::OperatorKind::UNARY_COS, SOFIE::OperatorKind::UNARY_ABS, - SOFIE::OperatorKind::NOT + SOFIE::OperatorKind::NOT, + SOFIE::OperatorKind::SELU }; bool OpNeedsBlas = false; From a9c330c8293d328714c578974adfa8be303ffb3e Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Tue, 16 Jun 2026 03:03:11 +0530 Subject: [PATCH 2/4] add gtest for selu gpu support --- test/alpaka/TestAlpakaElementwiseUnary.cxx | 35 +++++++++++++++++++++ test/input_models/Selu.onnx | Bin 0 -> 90 bytes test/input_models/references/Selu.ref.hxx | 5 +++ 3 files changed, 40 insertions(+) create mode 100644 test/input_models/Selu.onnx create mode 100644 test/input_models/references/Selu.ref.hxx diff --git a/test/alpaka/TestAlpakaElementwiseUnary.cxx b/test/alpaka/TestAlpakaElementwiseUnary.cxx index ff9d729a..ec9f06c4 100644 --- a/test/alpaka/TestAlpakaElementwiseUnary.cxx +++ b/test/alpaka/TestAlpakaElementwiseUnary.cxx @@ -16,6 +16,8 @@ #include "Softplus_FromONNX_GPU_ALPAKA.hxx" #include "Elu_FromONNX_GPU_ALPAKA.hxx" #include "input_models/references/Elu.ref.hxx" +#include "Selu_FromONNX_GPU_ALPAKA.hxx" +#include "input_models/references/Selu.ref.hxx" TEST_F(SofieAlpakaTest, Sin) { @@ -357,3 +359,36 @@ TEST_F(SofieAlpakaTest, Elu) } } +TEST_F(SofieAlpakaTest, Selu) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + + std::vector input({1.0f, -2.0f, 3.0f, 0.5f, -1.0f, 2.0f}); + + auto input_h = alpaka::allocBuf(host, Ext1D::all(Idx{input.size()})); + float* input_ptr = reinterpret_cast(alpaka::getPtrNative(input_h)); + for (Idx i = 0; i < input.size(); ++i) input_ptr[i] = input[i]; + + auto input_d = alpaka::allocBuf(device, Ext1D::all(Idx{input.size()})); + alpaka::memcpy(queue, input_d, input_h); + alpaka::wait(queue); + + constexpr size_t nOut = sizeof(Selu_ExpectedOutput::outputs) / sizeof(float); + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{nOut})); + + { + SOFIE_Selu::Session session; + auto result = session.infer(input_d); + alpaka::wait(queue); + cudaDeviceSynchronize(); + alpaka::memcpy(queue, result_h, result); + alpaka::wait(queue); + } + + float* res_ptr = reinterpret_cast(alpaka::getPtrNative(result_h)); + float* correct = Selu_ExpectedOutput::outputs; + for (size_t i = 0; i < nOut; ++i) { + EXPECT_LE(std::abs(res_ptr[i] - correct[i]), TOLERANCE) << "i=" << i; + } +} + diff --git a/test/input_models/Selu.onnx b/test/input_models/Selu.onnx new file mode 100644 index 0000000000000000000000000000000000000000..ce999694f1c939033beef5471c0eb0a0dda71551 GIT binary patch literal 90 zcmd Date: Fri, 14 Aug 2026 03:14:31 +0530 Subject: [PATCH 3/4] feat: make Selu alpha and gamma configurable via ONNX attributes --- core/inc/SOFIE/ROperator_Selu.hxx | 21 +++++++++++++-------- parsers/src/ParseSelu.cxx | 15 ++++++++++++++- test/input_models/Selu.onnx | Bin 90 -> 124 bytes test/input_models/references/Selu.ref.hxx | 2 +- 4 files changed, 28 insertions(+), 10 deletions(-) diff --git a/core/inc/SOFIE/ROperator_Selu.hxx b/core/inc/SOFIE/ROperator_Selu.hxx index 3db298d0..4167049f 100644 --- a/core/inc/SOFIE/ROperator_Selu.hxx +++ b/core/inc/SOFIE/ROperator_Selu.hxx @@ -5,8 +5,6 @@ #include "SOFIE/ROperator.hxx" #include "SOFIE/RModel.hxx" -#include - namespace SOFIE{ template @@ -15,13 +13,16 @@ class ROperator_Selu final : public ROperator private: + float falpha = 1.67326319217681884765625f; //ONNX spec default + float fgamma = 1.05070102214813232421875f; //ONNX spec default std::string fNX; std::string fNY; std::vector fShape; public: ROperator_Selu(){} - ROperator_Selu(std::string nameX, std::string nameY): + ROperator_Selu(float alpha, float gamma, std::string nameX, std::string nameY): + falpha(alpha), fgamma(gamma), fNX(UTILITY::Clean_name(nameX)), fNY(UTILITY::Clean_name(nameY)){ fInputTensorNames = { fNX }; fOutputTensorNames = { fNY }; @@ -53,8 +54,10 @@ public: } std::stringstream out; std::string length = ConvertDimShapeToLength(fShape); + out << "\t" << "constexpr float " << OpName << "_alpha = " << std::setprecision(std::numeric_limits::max_digits10) << falpha << ";\n"; + out << "\t" << "constexpr float " << OpName << "_gamma = " << std::setprecision(std::numeric_limits::max_digits10) << fgamma << ";\n"; out << "\t" << "for (int id = 0; id < " << length << " ; id++){\n"; - out << "\t\t" << "tensor_" << fNY << "[id] = 1.0507009873554804934193349852946 * (std::max(float(0.0), tensor_" << fNX << "[id]) + std::min(0.0, 1.6732632423543772848170429916717 * (std::exp(" << "tensor_" << fNX << "[id]" <<")-1)));\n"; + out << "\t\t" << "tensor_" << fNY << "[id] = " << OpName << "_gamma * (std::max(0.0f, tensor_" << fNX << "[id]) + std::min(0.0f, " << OpName << "_alpha * (std::exp(" << "tensor_" << fNX << "[id]" <<")-1)));\n"; out << "\t}\n"; return out.str(); } @@ -66,12 +69,12 @@ public: op = "\n//---- SELU_KERNEL_ALPAKA//\n"; op += "struct SeluKernel {\n"; op += SP + "template\n"; - op += SP + "ALPAKA_FN_ACC void operator()(TAcc const& acc, T const* __restrict__ data, T* __restrict__ out, std::size_t numElements) const {\n"; + op += SP + "ALPAKA_FN_ACC void operator()(TAcc const& acc, T const* __restrict__ data, T* __restrict__ out, std::size_t numElements, T alpha, T gamma) const {\n"; op += SP + SP + "const auto idx = alpaka::getIdx(acc)[0];\n"; op += SP + SP + "if (idx < numElements) {\n"; op += SP + SP + SP + "T x = data[idx];\n"; - op += SP + SP + SP + "T inner = T(1.6732632423543772848170429916717) * (exp(x) - T(1));\n"; - op += SP + SP + SP + "out[idx] = T(1.0507009873554804934193349852946) * ((x > T(0) ? x : T(0)) + (inner < T(0) ? inner : T(0)));\n"; + op += SP + SP + SP + "T inner = alpha * (exp(x) - T(1));\n"; + op += SP + SP + SP + "out[idx] = gamma * ((x > T(0) ? x : T(0)) + (inner < T(0) ? inner : T(0)));\n"; op += SP + SP + "}\n"; op += SP + "}\n"; op += "};\n"; @@ -95,7 +98,9 @@ public: out << SP << "auto const workDiv_" << fNX << " = sofie_workdiv(elementsPerGrid_" << fNX << ");\n"; out << SP << "auto task_" << OpName << " = alpaka::createTaskKernel(workDiv_" << fNX << ", seluKernel, alpaka::getPtrNative(deviceBuf_" << fNX - << "), alpaka::getPtrNative(deviceBuf_" << fNY << "), static_cast(" << length << "));\n"; + << "), alpaka::getPtrNative(deviceBuf_" << fNY << "), static_cast(" << length << "), static_cast(" + << std::setprecision(std::numeric_limits::max_digits10) << falpha << "), static_cast(" + << std::setprecision(std::numeric_limits::max_digits10) << fgamma << "));\n"; out << SP << "alpaka::enqueue(queue, task_" << OpName << ");\n"; return out.str(); } diff --git a/parsers/src/ParseSelu.cxx b/parsers/src/ParseSelu.cxx index 5a37b296..59166e29 100644 --- a/parsers/src/ParseSelu.cxx +++ b/parsers/src/ParseSelu.cxx @@ -17,9 +17,22 @@ ParserFuncSignature ParseSelu = [](RModelParser_ONNX &parser, const onnx::NodePr std::unique_ptr op; + float attr_alpha = 1.67326319217681884765625f; + float attr_gamma = 1.05070102214813232421875f; + + for (int_t i = 0; i < nodeproto.attribute_size(); i++) { + std::string attribute_name = nodeproto.attribute(i).name(); + if (attribute_name == "alpha") + attr_alpha = nodeproto.attribute(i).f(); + else if (attribute_name == "gamma") + attr_gamma = nodeproto.attribute(i).f(); + } + std::string output_name = nodeproto.output(0); switch (input_type) { - case ETensorType::FLOAT: op.reset(new ROperator_Selu(input_name, output_name)); break; + case ETensorType::FLOAT: + op.reset(new ROperator_Selu(attr_alpha, attr_gamma, input_name, output_name)); + break; default: throw std::runtime_error("TMVA::SOFIE - Unsupported - Operator Selu does not yet support input type " + std::to_string(static_cast(input_type))); diff --git a/test/input_models/Selu.onnx b/test/input_models/Selu.onnx index ce999694f1c939033beef5471c0eb0a0dda71551..951e9465805079c84d90d1382d2b9c04a28b5f59 100644 GIT binary patch delta 67 zcma#5vE|^fD&jKdV$IAeC@m3U%P%bf(n>7BsX3)u{9LSwIRzPsq6`cS4ht9=K?3QC Oxw#+#2av!-X?p-r7Z8d7 delta 33 mcmb=4lIGyB3g8muV$IAeC@m3U%P%bf(n>7BsX3(+ZS4V$)(L+A diff --git a/test/input_models/references/Selu.ref.hxx b/test/input_models/references/Selu.ref.hxx index 33c8ff21..20d20307 100644 --- a/test/input_models/references/Selu.ref.hxx +++ b/test/input_models/references/Selu.ref.hxx @@ -1,5 +1,5 @@ // Auto-generated SELU reference - DO NOT EDIT #pragma once namespace Selu_ExpectedOutput { - float outputs[] = {1.05070102f, -1.52016652f, 3.15210295f, 0.52535051f, -1.11133075f, 2.10140204f}; + float outputs[] = {3.00000000f, -5.18798828f, 9.00000000f, 1.50000000f, -3.79272366f, 6.00000000f}; } // namespace Selu_ExpectedOutput From 2d996ca26797fb0cbd4680033f0efdbff3778191 Mon Sep 17 00:00:00 2001 From: Harsh Chauhan Date: Fri, 14 Aug 2026 13:56:47 +0530 Subject: [PATCH 4/4] test: rename Selu gtest and model to SeluNonDefaultCoeffs --- core/inc/SOFIE/ROperator_Selu.hxx | 4 ++-- test/alpaka/TestAlpakaElementwiseUnary.cxx | 12 ++++++------ .../{Selu.onnx => SeluNonDefaultCoeffs.onnx} | Bin .../{Selu.ref.hxx => SeluNonDefaultCoeffs.ref.hxx} | 5 ++--- 4 files changed, 10 insertions(+), 11 deletions(-) rename test/input_models/{Selu.onnx => SeluNonDefaultCoeffs.onnx} (100%) rename test/input_models/references/{Selu.ref.hxx => SeluNonDefaultCoeffs.ref.hxx} (50%) diff --git a/core/inc/SOFIE/ROperator_Selu.hxx b/core/inc/SOFIE/ROperator_Selu.hxx index 4167049f..bbeb1ae8 100644 --- a/core/inc/SOFIE/ROperator_Selu.hxx +++ b/core/inc/SOFIE/ROperator_Selu.hxx @@ -13,8 +13,8 @@ class ROperator_Selu final : public ROperator private: - float falpha = 1.67326319217681884765625f; //ONNX spec default - float fgamma = 1.05070102214813232421875f; //ONNX spec default + float falpha = 1.67326319217681884765625f; + float fgamma = 1.05070102214813232421875f; std::string fNX; std::string fNY; std::vector fShape; diff --git a/test/alpaka/TestAlpakaElementwiseUnary.cxx b/test/alpaka/TestAlpakaElementwiseUnary.cxx index ec9f06c4..61f714bd 100644 --- a/test/alpaka/TestAlpakaElementwiseUnary.cxx +++ b/test/alpaka/TestAlpakaElementwiseUnary.cxx @@ -16,8 +16,8 @@ #include "Softplus_FromONNX_GPU_ALPAKA.hxx" #include "Elu_FromONNX_GPU_ALPAKA.hxx" #include "input_models/references/Elu.ref.hxx" -#include "Selu_FromONNX_GPU_ALPAKA.hxx" -#include "input_models/references/Selu.ref.hxx" +#include "SeluNonDefaultCoeffs_FromONNX_GPU_ALPAKA.hxx" +#include "input_models/references/SeluNonDefaultCoeffs.ref.hxx" TEST_F(SofieAlpakaTest, Sin) { @@ -359,7 +359,7 @@ TEST_F(SofieAlpakaTest, Elu) } } -TEST_F(SofieAlpakaTest, Selu) +TEST_F(SofieAlpakaTest, SeluNonDefaultCoeffs) { constexpr float TOLERANCE = DEFAULT_TOLERANCE; @@ -373,11 +373,11 @@ TEST_F(SofieAlpakaTest, Selu) alpaka::memcpy(queue, input_d, input_h); alpaka::wait(queue); - constexpr size_t nOut = sizeof(Selu_ExpectedOutput::outputs) / sizeof(float); + constexpr size_t nOut = sizeof(SeluNonDefaultCoeffs_ExpectedOutput::outputs) / sizeof(float); auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{nOut})); { - SOFIE_Selu::Session session; + SOFIE_SeluNonDefaultCoeffs::Session session; auto result = session.infer(input_d); alpaka::wait(queue); cudaDeviceSynchronize(); @@ -386,7 +386,7 @@ TEST_F(SofieAlpakaTest, Selu) } float* res_ptr = reinterpret_cast(alpaka::getPtrNative(result_h)); - float* correct = Selu_ExpectedOutput::outputs; + float* correct = SeluNonDefaultCoeffs_ExpectedOutput::outputs; for (size_t i = 0; i < nOut; ++i) { EXPECT_LE(std::abs(res_ptr[i] - correct[i]), TOLERANCE) << "i=" << i; } diff --git a/test/input_models/Selu.onnx b/test/input_models/SeluNonDefaultCoeffs.onnx similarity index 100% rename from test/input_models/Selu.onnx rename to test/input_models/SeluNonDefaultCoeffs.onnx diff --git a/test/input_models/references/Selu.ref.hxx b/test/input_models/references/SeluNonDefaultCoeffs.ref.hxx similarity index 50% rename from test/input_models/references/Selu.ref.hxx rename to test/input_models/references/SeluNonDefaultCoeffs.ref.hxx index 20d20307..6d7afe2b 100644 --- a/test/input_models/references/Selu.ref.hxx +++ b/test/input_models/references/SeluNonDefaultCoeffs.ref.hxx @@ -1,5 +1,4 @@ -// Auto-generated SELU reference - DO NOT EDIT #pragma once -namespace Selu_ExpectedOutput { +namespace SeluNonDefaultCoeffs_ExpectedOutput { float outputs[] = {3.00000000f, -5.18798828f, 9.00000000f, 1.50000000f, -3.79272366f, 6.00000000f}; -} // namespace Selu_ExpectedOutput +} // namespace SeluNonDefaultCoeffs_ExpectedOutput