diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index a99f6d4..39d6590 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -21,6 +21,8 @@ set(sources_headers SOFIE/RModelProfilerGPU.hxx SOFIE/ROperator.hxx SOFIE/ROperator_BasicUnary.hxx + ROperator_HardSigmoid.hxx + ROperator_HardSwish.hxx SOFIE/ROperator_BasicBinary.hxx SOFIE/ROperator_BasicNary.hxx SOFIE/ROperator_BatchNormalization.hxx diff --git a/core/inc/SOFIE/ROperator.hxx b/core/inc/SOFIE/ROperator.hxx index aa20352..5993395 100644 --- a/core/inc/SOFIE/ROperator.hxx +++ b/core/inc/SOFIE/ROperator.hxx @@ -39,7 +39,10 @@ enum class OperatorKind { UNARY_ABS=23, CLIP=24, NOT=25, - POOL=26 + POOL=26, + HARDSIGMOID=27, + HARDSWISH=28, + SOFTPLUS=29 }; inline const char* toString(OperatorKind kind) { @@ -52,6 +55,9 @@ inline const char* toString(OperatorKind kind) { case OperatorKind::BATCHNORM: return "BATCHNORM"; case OperatorKind::CONV: return "CONV"; case OperatorKind::UNDEFINED: return "UNDEFINED"; + case OperatorKind::HARDSIGMOID:return "HARDSIGMOID"; + case OperatorKind::HARDSWISH: return "HARDSWISH"; + case OperatorKind::SOFTPLUS: return "SOFTPLUS"; default: return "UNKNOWN"; } } diff --git a/core/inc/SOFIE/ROperator_BasicUnary.hxx b/core/inc/SOFIE/ROperator_BasicUnary.hxx index dfe6714..30e815a 100644 --- a/core/inc/SOFIE/ROperator_BasicUnary.hxx +++ b/core/inc/SOFIE/ROperator_BasicUnary.hxx @@ -65,7 +65,7 @@ struct UnaryOpTraits { template struct UnaryOpTraits { static std::string Name() { return "Softplus"; } - static std::string Op(const std::string &X) { return "std::log(std::exp(" + X + ") + 1)"; } + static std::string Op(const std::string &X) { return "((" + X + " >= 0x1.4000000000000p+4f) ? " + X + " : std::log1p(std::exp(" + X + ")))"; } }; template @@ -121,6 +121,9 @@ public: case EBasicUnaryOperator::kAbs: fKind = OperatorKind::UNARY_ABS; break; + case EBasicUnaryOperator::kSoftplus: + fKind = OperatorKind::SOFTPLUS; + break; } fInputTensorNames = { fNX }; fOutputTensorNames = { fNY }; diff --git a/core/inc/SOFIE/ROperator_HardSigmoid.hxx b/core/inc/SOFIE/ROperator_HardSigmoid.hxx new file mode 100644 index 0000000..3c5308a --- /dev/null +++ b/core/inc/SOFIE/ROperator_HardSigmoid.hxx @@ -0,0 +1,101 @@ +#ifndef SOFIE_ROPERATOR_HARDSIGMOID +#define SOFIE_ROPERATOR_HARDSIGMOID + +#include +#include +#include + +#include + +namespace SOFIE { + +template +class ROperator_HardSigmoid final : public ROperator +{ + +private: + + std::string fNX; + std::string fNY; + std::vector fShape; + float fAlpha; + float fBeta; + +public: + ROperator_HardSigmoid(){} + ROperator_HardSigmoid(std::string nameX, std::string nameY, float alpha, float beta): + fNX(UTILITY::Clean_name(nameX)), fNY(UTILITY::Clean_name(nameY)), fAlpha(alpha), fBeta(beta){ + fInputTensorNames = { fNX }; + fOutputTensorNames = { fNY }; + fKind = OperatorKind::HARDSIGMOID; + } + + std::vector TypeInference(std::vector input) override { + return input; + } + + std::vector> ShapeInference(std::vector> input) override { + return input; + } + + void Initialize(RModel& model) override { + if (!model.CheckIfTensorAlreadyExist(fNX)){ + throw std::runtime_error("SOFIE HardSigmoid Op Input Tensor " + fNX + " is not found in model"); + } + fShape = model.GetTensorShape(fNX); + model.AddIntermediateTensor(fNY, model.GetTensorType(fNX), fShape); + } + + std::string Generate(std::string OpName) override { + OpName = "op_" + OpName; + if (fShape.empty()){ + throw std::runtime_error("SOFIE HardSigmoid operator called to Generate without being initialized first"); + } + std::stringstream out; + size_t length = ConvertShapeToLength(fShape); + + // HardSigmoid: y = max(0, min(1, alpha * x + beta)) + out << "\n//------ HardSigmoid\n"; + out << SP << "for (int id = 0; id < " << length << " ; id++){\n"; + out << SP << SP << "tensor_" << fNY << "[id] = std::fmax(0x0p+0f, std::fmin(0x1p+0f, " + << fAlpha << "f * tensor_" << fNX << "[id] + " << fBeta << "f));\n"; + out << SP << "}\n"; + return out.str(); + } + std::string Generate_GPU_Kernel_ALPAKA(std::string /*opName*/) override { + std::string op = "\n//------ HARDSIGMOID_KERNEL_ALPAKA\n"; + op += SP + "struct HardSigmoidKernel{\n"; + op += SP + SP + "template\n"; + op += SP + SP + "ALPAKA_FN_ACC void operator()(TAcc const & acc, T const * data, T * out, std::size_t numElements, T const alpha, T const beta) const {\n"; + op += SP + SP + SP + "const auto idx = alpaka::getIdx(acc)[0];\n"; + op += SP + SP + SP + "if (idx < numElements) {\n"; + op += SP + SP + SP + SP + "T x = data[idx];\n"; + op += SP + SP + SP + SP + "T h = alpha * x + beta;\n"; + op += SP + SP + SP + SP + "out[idx] = (h < T(0)) ? T(0) : ((h > T(1)) ? T(1) : h);\n"; + op += SP + SP + SP + "}\n"; + op += SP + SP + "}\n"; + op += SP + "};\n"; + return op; + } + + std::string Generate_GPU_Kernel_Definitions_ALPAKA(std::string /*opName*/) override { + return SP + "HardSigmoidKernel hardSigmoidKernel;\n"; + } + + std::string Generate_GPU_ALPAKA(std::string OpName) override { + std::stringstream out; + auto length = ConvertShapeToLength(fShape); + out << "\n//------ op_" << OpName << "_ALPAKA\n"; + out << SP << "auto const elementsPerThread_" << fNX << " = alpaka::Vec::all(static_cast(1));\n"; + out << SP << "auto const elementsPerGrid_" << fNX << " = alpaka::Vec::all(static_cast(" << length << "));\n"; + out << SP << "auto const workDiv_" << fNX << " = sofie_workdiv(elementsPerGrid_" << fNX << ");\n"; + out << SP << "auto task_op_" << OpName << " = alpaka::createTaskKernel(workDiv_" << fNX << ", hardSigmoidKernel, alpaka::getPtrNative(deviceBuf_" << fNX << "), alpaka::getPtrNative(deviceBuf_" << fNY << "), static_cast(" << length << "), static_cast(" << fAlpha << "), static_cast(" << fBeta << "));\n"; + out << SP << "alpaka::enqueue(queue, task_op_" << OpName << ");\n"; + return out.str(); + } + std::vector GetStdLibs() override { return { std::string("cmath") };} +}; + +} // namespace SOFIE + +#endif \ No newline at end of file diff --git a/core/inc/SOFIE/ROperator_HardSwish.hxx b/core/inc/SOFIE/ROperator_HardSwish.hxx new file mode 100644 index 0000000..8dd1738 --- /dev/null +++ b/core/inc/SOFIE/ROperator_HardSwish.hxx @@ -0,0 +1,103 @@ +#ifndef SOFIE_ROPERATOR_HARDSWISH +#define SOFIE_ROPERATOR_HARDSWISH + +#include +#include +#include + +#include + +namespace SOFIE { + +template +class ROperator_HardSwish final : public ROperator +{ + +private: + + std::string fNX; + std::string fNY; + std::vector fShape; + +public: + ROperator_HardSwish(){} + ROperator_HardSwish(std::string nameX, std::string nameY): + fNX(UTILITY::Clean_name(nameX)), fNY(UTILITY::Clean_name(nameY)){ + fInputTensorNames = { fNX }; + fOutputTensorNames = { fNY }; + fKind = OperatorKind::HARDSWISH; + } + + std::vector TypeInference(std::vector input) override { + return input; + } + + std::vector> ShapeInference(std::vector> input) override { + return input; + } + + void Initialize(RModel& model) override { + if (!model.CheckIfTensorAlreadyExist(fNX)){ + throw std::runtime_error("SOFIE HardSwish Op Input Tensor " + fNX + " is not found in model"); + } + fShape = model.GetTensorShape(fNX); + model.AddIntermediateTensor(fNY, model.GetTensorType(fNX), fShape); + } + + std::string Generate(std::string OpName) override { + OpName = "op_" + OpName; + if (fShape.empty()){ + throw std::runtime_error("SOFIE HardSwish operator called to Generate without being initialized first"); + } + std::stringstream out; + size_t length = ConvertShapeToLength(fShape); + + // HardSwish: y = x * max(0, min(1, x/6 + 0.5)) + // Split topology for debuggability + out << "\n//------ HardSwish\n"; + out << SP << "for (int id = 0; id < " << length << " ; id++){\n"; + out << SP << SP << "float h = 0x1.5555555555555p-3f * tensor_" << fNX << "[id] + 0x1p-1f;\n"; + out << SP << SP << "tensor_" << fNY << "[id] = tensor_" << fNX + << "[id] * std::fmax(0x0p+0f, std::fmin(0x1p+0f, h));\n"; + out << SP << "}\n"; + return out.str(); + } + + std::string Generate_GPU_Kernel_ALPAKA(std::string /*opName*/) override { + std::string op = "\n//------ HARDSWISH_KERNEL_ALPAKA\n"; + op += SP + "struct HardSwishKernel{\n"; + op += SP + SP + "template\n"; + op += SP + SP + "ALPAKA_FN_ACC void operator()(TAcc const & acc, T const * data, T * out, std::size_t numElements) const {\n"; + op += SP + SP + SP + "const auto idx = alpaka::getIdx(acc)[0];\n"; + op += SP + SP + SP + "if (idx < numElements) {\n"; + op += SP + SP + SP + SP + "T x = data[idx];\n"; + op += SP + SP + SP + SP + "T h = T(0x1.5555555555555p-3) * x + T(0.5);\n"; + op += SP + SP + SP + SP + "out[idx] = x * ((h < T(0)) ? T(0) : ((h > T(1)) ? T(1) : h));\n"; + op += SP + SP + SP + "}\n"; + op += SP + SP + "}\n"; + op += SP + "};\n"; + return op; + } + + std::string Generate_GPU_Kernel_Definitions_ALPAKA(std::string /*opName*/) override { + return SP + "HardSwishKernel hardSwishKernel;\n"; + } + + std::string Generate_GPU_ALPAKA(std::string OpName) override { + std::stringstream out; + auto length = ConvertShapeToLength(fShape); + out << "\n//------ op_" << OpName << "_ALPAKA\n"; + out << SP << "auto const elementsPerThread_" << fNX << " = alpaka::Vec::all(static_cast(1));\n"; + out << SP << "auto const elementsPerGrid_" << fNX << " = alpaka::Vec::all(static_cast(" << length << "));\n"; + out << SP << "auto const workDiv_" << fNX << " = sofie_workdiv(elementsPerGrid_" << fNX << ");\n"; + out << SP << "auto task_op_" << OpName << " = alpaka::createTaskKernel(workDiv_" << fNX << ", hardSwishKernel, alpaka::getPtrNative(deviceBuf_" << fNX << "), alpaka::getPtrNative(deviceBuf_" << fNY << "), static_cast(" << length << "));\n"; + out << SP << "alpaka::enqueue(queue, task_op_" << OpName << ");\n"; + return out.str(); + } + + std::vector GetStdLibs() override { return { std::string("cmath") };} +}; + +} // namespace SOFIE + +#endif \ No newline at end of file diff --git a/parsers/CMakeLists.txt b/parsers/CMakeLists.txt index 7174e90..22b3e71 100644 --- a/parsers/CMakeLists.txt +++ b/parsers/CMakeLists.txt @@ -33,6 +33,8 @@ target_include_directories(SOFIE_parsers set(sources_cxx src/RModelParser_ONNX.cxx src/ParseBasicUnary.cxx + ParseHardSigmoid.cxx + ParseHardSwish.cxx src/ParseBasicBinary.cxx src/ParseBasicIs.cxx src/ParseBatchNormalization.cxx diff --git a/parsers/src/ParseHardSigmoid.cxx b/parsers/src/ParseHardSigmoid.cxx new file mode 100644 index 0000000..625e496 --- /dev/null +++ b/parsers/src/ParseHardSigmoid.cxx @@ -0,0 +1,47 @@ +#include "SOFIE/RModelParser_ONNX.hxx" +#include "SOFIE/ROperator_HardSigmoid.hxx" +#include "onnx_proto3.pb.h" + +namespace SOFIE { + +ParserFuncSignature ParseHardSigmoid = [](RModelParser_ONNX &parser, const onnx::NodeProto &nodeproto) { + ETensorType input_type; + + // ONNX spec defaults: alpha=0.2, beta=0.5 + float alpha = 0.2f; + float beta = 0.5f; + + for (int_t i = 0; i < nodeproto.attribute_size(); i++) { + std::string attribute_name = nodeproto.attribute(i).name(); + if (attribute_name == "alpha") + alpha = nodeproto.attribute(i).f(); + else if (attribute_name == "beta") + beta = nodeproto.attribute(i).f(); + } + + auto input_name = nodeproto.input(0); + if (parser.IsRegisteredTensorType(input_name)) { + input_type = parser.GetTensorType(input_name); + } else { + throw std::runtime_error("TMVA::SOFIE ONNX Parser HardSigmoid op has input tensor " + input_name + + " but its type is not yet registered"); + } + + std::unique_ptr op; + std::string output_name = nodeproto.output(0); + + switch (input_type) { + case ETensorType::FLOAT: op.reset(new ROperator_HardSigmoid(input_name, output_name, alpha, beta)); break; + default: + throw std::runtime_error("TMVA::SOFIE - Unsupported - Operator HardSigmoid does not yet support input type " + + std::to_string(static_cast(input_type))); + } + + if (!parser.IsRegisteredTensorType(output_name)) { + parser.RegisterTensorType(output_name, input_type); + } + + return op; +}; + +} // namespace SOFIE \ No newline at end of file diff --git a/parsers/src/ParseHardSwish.cxx b/parsers/src/ParseHardSwish.cxx new file mode 100644 index 0000000..21fc398 --- /dev/null +++ b/parsers/src/ParseHardSwish.cxx @@ -0,0 +1,35 @@ +#include "SOFIE/RModelParser_ONNX.hxx" +#include "SOFIE/ROperator_HardSwish.hxx" +#include "onnx_proto3.pb.h" + +namespace SOFIE { + +ParserFuncSignature ParseHardSwish = [](RModelParser_ONNX &parser, const onnx::NodeProto &nodeproto) { + ETensorType input_type; + + auto input_name = nodeproto.input(0); + if (parser.IsRegisteredTensorType(input_name)) { + input_type = parser.GetTensorType(input_name); + } else { + throw std::runtime_error("TMVA::SOFIE ONNX Parser HardSwish op has input tensor " + input_name + + " but its type is not yet registered"); + } + + std::unique_ptr op; + std::string output_name = nodeproto.output(0); + + switch (input_type) { + case ETensorType::FLOAT: op.reset(new ROperator_HardSwish(input_name, output_name)); break; + default: + throw std::runtime_error("TMVA::SOFIE - Unsupported - Operator HardSwish does not yet support input type " + + std::to_string(static_cast(input_type))); + } + + if (!parser.IsRegisteredTensorType(output_name)) { + parser.RegisterTensorType(output_name, input_type); + } + + return op; +}; + +} // namespace SOFIE \ No newline at end of file diff --git a/parsers/src/RModelParser_ONNX.cxx b/parsers/src/RModelParser_ONNX.cxx index afb8b93..51729f2 100644 --- a/parsers/src/RModelParser_ONNX.cxx +++ b/parsers/src/RModelParser_ONNX.cxx @@ -200,6 +200,8 @@ RModelParser_ONNX::RModelParser_ONNX() noexcept : fOperatorsMapImpl(std::make_un RegisterOperator("Cos", ParseCos); RegisterOperator("Abs", ParseAbs); RegisterOperator("Softplus", ParseSoftplus); + RegisterOperator("HardSigmoid", ParseHardSigmoid); + RegisterOperator("HardSwish", ParseHardSwish); RegisterOperator("Atan", ParseAtan); RegisterOperator("Floor", ParseFloor); diff --git a/test/alpaka/TestAlpakaElementwiseUnary.cxx b/test/alpaka/TestAlpakaElementwiseUnary.cxx index ff9d729..15e0ad2 100644 --- a/test/alpaka/TestAlpakaElementwiseUnary.cxx +++ b/test/alpaka/TestAlpakaElementwiseUnary.cxx @@ -13,6 +13,10 @@ #include "input_models/references/Exp.ref.hxx" #include "input_models/references/Log.ref.hxx" #include "input_models/references/Neg.ref.hxx" +#include "HardSigmoid_FromONNX_GPU_ALPAKA.hxx" +#include "HardSwish_FromONNX_GPU_ALPAKA.hxx" +#include "input_models/references/HardSigmoid.ref.hxx" +#include "input_models/references/HardSwish.ref.hxx" #include "Softplus_FromONNX_GPU_ALPAKA.hxx" #include "Elu_FromONNX_GPU_ALPAKA.hxx" #include "input_models/references/Elu.ref.hxx" @@ -357,3 +361,86 @@ TEST_F(SofieAlpakaTest, Elu) } } + +TEST_F(SofieAlpakaTest, HardSigmoid) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + + std::vector input({ + -10.0f, -6.0f, -3.0f, -1.0f, 0.0f, 1.0f, 3.0f, 6.0f, 10.0f, 16.0f, 20.0f, + -4.52299976f, -2.11800003f, -0.68800002f, -0.15099999f, 0.09799999f, 0.43599998f, + 1.22300004f, 2.70499992f, 4.33099985f, 8.08300018f, 14.6459999f, 18.2789993f, 25.0f + }); + + auto A = alpaka::allocBuf(host, Ext1D::all(Idx{input.size()})); + float *A_ptr = reinterpret_cast(alpaka::getPtrNative(A)); + + for (Idx i = 0; i < input.size(); ++i) { + A_ptr[i] = input[i]; + } + + auto A_d = alpaka::allocBuf(device, Ext1D::all(Idx{input.size()})); + alpaka::memcpy(queue, A_d, A); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{input.size()})); + + { + SOFIE_HardSigmoid::Session session("HardSigmoid_FromONNX_GPU_ALPAKA.dat"); + auto result = session.infer(A_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 = HardSigmoid_ExpectedOutput::outputs; + + for (size_t i = 0; i < input.size(); ++i) { + EXPECT_LE(std::abs(res_ptr[i] - correct[i]), TOLERANCE); + } +} + +TEST_F(SofieAlpakaTest, HardSwish) +{ + constexpr float TOLERANCE = DEFAULT_TOLERANCE; + + std::vector input({ + -10.0f, -6.0f, -3.0f, -1.0f, 0.0f, 1.0f, 3.0f, 6.0f, 10.0f, 16.0f, 20.0f, + -4.52299976f, -2.11800003f, -0.68800002f, -0.15099999f, 0.09799999f, 0.43599998f, + 1.22300004f, 2.70499992f, 4.33099985f, 8.08300018f, 14.6459999f, 18.2789993f, 25.0f + }); + + auto A = alpaka::allocBuf(host, Ext1D::all(Idx{input.size()})); + float *A_ptr = reinterpret_cast(alpaka::getPtrNative(A)); + + for (Idx i = 0; i < input.size(); ++i) { + A_ptr[i] = input[i]; + } + + auto A_d = alpaka::allocBuf(device, Ext1D::all(Idx{input.size()})); + alpaka::memcpy(queue, A_d, A); + alpaka::wait(queue); + + auto result_h = alpaka::allocBuf(host, Ext1D::all(Idx{input.size()})); + + { + SOFIE_HardSwish::Session session("HardSwish_FromONNX_GPU_ALPAKA.dat"); + auto result = session.infer(A_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 = HardSwish_ExpectedOutput::outputs; + + for (size_t i = 0; i < input.size(); ++i) { + EXPECT_LE(std::abs(res_ptr[i] - correct[i]), TOLERANCE); + } +} + diff --git a/test/input_models/HardSigmoid.onnx b/test/input_models/HardSigmoid.onnx new file mode 100644 index 0000000..cf1b27e Binary files /dev/null and b/test/input_models/HardSigmoid.onnx differ diff --git a/test/input_models/HardSwish.onnx b/test/input_models/HardSwish.onnx new file mode 100644 index 0000000..336ccac --- /dev/null +++ b/test/input_models/HardSwish.onnx @@ -0,0 +1,12 @@ +  +sofie-test:F + +XY" HardSwishHardSwish_ModelZ +X + + +b +Y + + +B \ No newline at end of file diff --git a/test/input_models/references/HardSigmoid.ref.hxx b/test/input_models/references/HardSigmoid.ref.hxx new file mode 100644 index 0000000..6e032e0 --- /dev/null +++ b/test/input_models/references/HardSigmoid.ref.hxx @@ -0,0 +1,5 @@ +namespace HardSigmoid_ExpectedOutput { +float outputs[] = { +0.0f, 0.0f, 0.0f, 0.3f, 0.5f, 0.7f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.07639999389648433f, 0.36239999532699585f, 0.46980000138282774f, 0.5195999994874001f, 0.5871999979019165f, 0.7446000099182128f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f +}; +} diff --git a/test/input_models/references/HardSwish.ref.hxx b/test/input_models/references/HardSwish.ref.hxx new file mode 100644 index 0000000..4fd9a9c --- /dev/null +++ b/test/input_models/references/HardSwish.ref.hxx @@ -0,0 +1,5 @@ +namespace HardSwish_ExpectedOutput { +float outputs[] = { +-0.0f, -0.0f, -0.0f, -0.33333333333333337f, 0.0f, 0.6666666666666666f, 3.0f, 6.0f, 10.0f, 16.0f, 20.0f, -0.0f, -0.3113459937133788f, -0.2651093396574655f, -0.07169983022427559f, 0.050600665301442145f, 0.24968265989685062f, 0.8607882116788232f, 2.5720040597279876f, 4.330999851226807f, 8.083000183105469f, 14.645999908447266f, 18.27899932861328f, 25.0f +}; +}