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..bbeb1ae8 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,16 +13,20 @@ class ROperator_Selu final : public ROperator private: + float falpha = 1.67326319217681884765625f; + float fgamma = 1.05070102214813232421875f; 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 }; + fKind = OperatorKind::SELU; } std::vector TypeInference(std::vector input) override { @@ -52,13 +54,56 @@ 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(); } 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, 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 = 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"; + 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 << "), 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(); + } }; }//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; 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/alpaka/TestAlpakaElementwiseUnary.cxx b/test/alpaka/TestAlpakaElementwiseUnary.cxx index ff9d729a..61f714bd 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 "SeluNonDefaultCoeffs_FromONNX_GPU_ALPAKA.hxx" +#include "input_models/references/SeluNonDefaultCoeffs.ref.hxx" TEST_F(SofieAlpakaTest, Sin) { @@ -357,3 +359,36 @@ TEST_F(SofieAlpakaTest, Elu) } } +TEST_F(SofieAlpakaTest, SeluNonDefaultCoeffs) +{ + 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(SeluNonDefaultCoeffs_ExpectedOutput::outputs) / sizeof(float); + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{nOut})); + + { + SOFIE_SeluNonDefaultCoeffs::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 = 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/SeluNonDefaultCoeffs.onnx b/test/input_models/SeluNonDefaultCoeffs.onnx new file mode 100644 index 00000000..951e9465 Binary files /dev/null and b/test/input_models/SeluNonDefaultCoeffs.onnx differ diff --git a/test/input_models/references/SeluNonDefaultCoeffs.ref.hxx b/test/input_models/references/SeluNonDefaultCoeffs.ref.hxx new file mode 100644 index 00000000..6d7afe2b --- /dev/null +++ b/test/input_models/references/SeluNonDefaultCoeffs.ref.hxx @@ -0,0 +1,4 @@ +#pragma once +namespace SeluNonDefaultCoeffs_ExpectedOutput { + float outputs[] = {3.00000000f, -5.18798828f, 9.00000000f, 1.50000000f, -3.79272366f, 6.00000000f}; +} // namespace SeluNonDefaultCoeffs_ExpectedOutput