extended gpu support for selu

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) {

core/inc/SOFIE/ROperator_Selu.hxx

@@ -5,8 +5,6 @@
 #include "SOFIE/ROperator.hxx"
 #include "SOFIE/RModel.hxx"
 
-#include <sstream>
-
 namespace SOFIE{
 
 template <typename T>
@@ -15,16 +13,20 @@ 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<Dim> 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<ETensorType> TypeInference(std::vector<ETensorType> 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<float>::max_digits10) << falpha << ";\n";
+      out << "\t" << "constexpr float " << OpName << "_gamma = " << std::setprecision(std::numeric_limits<float>::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<std::string> 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<typename TAcc, typename T>\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<alpaka::Grid, alpaka::Threads>(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<Idx>(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<Acc>(workDiv_" << fNX
+         << ", seluKernel, alpaka::getPtrNative(deviceBuf_" << fNX
+         << "), alpaka::getPtrNative(deviceBuf_" << fNY << "), static_cast<Idx>(" << length << "), static_cast<float>("
+         << std::setprecision(std::numeric_limits<float>::max_digits10) << falpha << "), static_cast<float>("
+         << std::setprecision(std::numeric_limits<float>::max_digits10) << fgamma << "));\n";
+      out << SP << "alpaka::enqueue(queue, task_" << OpName << ");\n";
+      return out.str();
+   }
 };
 
 }//SOFIE

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;

parsers/src/ParseSelu.cxx

@@ -17,9 +17,22 @@ ParserFuncSignature ParseSelu = [](RModelParser_ONNX &parser, const onnx::NodePr
 
    std::unique_ptr<ROperator> 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<float>(input_name, output_name)); break;
+   case ETensorType::FLOAT:
+      op.reset(new ROperator_Selu<float>(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<int>(input_type)));

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<float> input({1.0f, -2.0f, 3.0f, 0.5f, -1.0f, 2.0f});
+
+    auto input_h = alpaka::allocBuf<float, Idx>(host, Ext1D::all(Idx{input.size()}));
+    float* input_ptr = reinterpret_cast<float*>(alpaka::getPtrNative(input_h));
+    for (Idx i = 0; i < input.size(); ++i) input_ptr[i] = input[i];
+
+    auto input_d = alpaka::allocBuf<float, Idx>(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<float, Idx>(host, Ext1D::all(Idx{nOut}));
+
+    {
+        SOFIE_Selu::Session<alpaka::TagGpuCudaRt> 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<float*>(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;
+    }
+}
+

test/input_models/references/Selu.ref.hxx

@@ -0,0 +1,5 @@
+// Auto-generated SELU reference - DO NOT EDIT
+#pragma once
+namespace Selu_ExpectedOutput {
+   float outputs[] = {3.00000000f, -5.18798828f, 9.00000000f, 1.50000000f, -3.79272366f, 6.00000000f};
+} // namespace Selu_ExpectedOutput