From 4995aecd62e0a65cfb9a0f75ccbe2e16d43f6c5f Mon Sep 17 00:00:00 2001 From: Smirnov Egor Date: Tue, 30 Nov 2021 14:43:18 +0300 Subject: [PATCH] add alpha parameter to ELU --- .../dnn/include/opencv2/dnn/all_layers.hpp | 2 ++ modules/dnn/src/cuda/activations.cu | 8 ++++---- modules/dnn/src/cuda/functors.hpp | 12 +++++++---- .../dnn/src/cuda4dnn/kernels/activations.hpp | 2 +- .../src/cuda4dnn/primitives/activation.hpp | 5 +++-- modules/dnn/src/layers/elementwise_layers.cpp | 20 ++++++++++++++----- modules/dnn/src/opencl/activations.cl | 5 +++-- 7 files changed, 36 insertions(+), 18 deletions(-) diff --git a/modules/dnn/include/opencv2/dnn/all_layers.hpp b/modules/dnn/include/opencv2/dnn/all_layers.hpp index c70ce2b50a..0bec9e35e2 100644 --- a/modules/dnn/include/opencv2/dnn/all_layers.hpp +++ b/modules/dnn/include/opencv2/dnn/all_layers.hpp @@ -545,6 +545,8 @@ CV__DNN_INLINE_NS_BEGIN class CV_EXPORTS ELULayer : public ActivationLayer { public: + float alpha; + static Ptr create(const LayerParams ¶ms); }; diff --git a/modules/dnn/src/cuda/activations.cu b/modules/dnn/src/cuda/activations.cu index c38fa0346f..0980b5dd46 100644 --- a/modules/dnn/src/cuda/activations.cu +++ b/modules/dnn/src/cuda/activations.cu @@ -119,8 +119,8 @@ void sigmoid(const Stream& stream, Span output, View input) { } template -void elu(const Stream& stream, Span output, View input) { - generic_op>(stream, output, input); +void elu(const Stream& stream, Span output, View input, T alpha) { + generic_op>(stream, output, input, {alpha}); } template @@ -187,7 +187,7 @@ template void tanh<__half>(const Stream&, Span<__half>, View<__half>); template void swish<__half>(const Stream&, Span<__half>, View<__half>); template void mish<__half>(const Stream&, Span<__half>, View<__half>); template void sigmoid<__half>(const Stream&, Span<__half>, View<__half>); -template void elu<__half>(const Stream&, Span<__half>, View<__half>); +template void elu<__half>(const Stream&, Span<__half>, View<__half>, __half); template void abs<__half>(const Stream& stream, Span<__half> output, View<__half> input); template void bnll<__half>(const Stream&, Span<__half>, View<__half>); template void ceil<__half>(const Stream&, Span<__half>, View<__half>); @@ -207,7 +207,7 @@ template void tanh(const Stream&, Span, View); template void swish(const Stream&, Span, View); template void mish(const Stream&, Span, View); template void sigmoid(const Stream&, Span, View); -template void elu(const Stream&, Span, View); +template void elu(const Stream&, Span, View, float); template void abs(const Stream& stream, Span output, View input); template void bnll(const Stream&, Span, View); template void ceil(const Stream&, Span, View); diff --git a/modules/dnn/src/cuda/functors.hpp b/modules/dnn/src/cuda/functors.hpp index 04b545acaf..98ae175ce8 100644 --- a/modules/dnn/src/cuda/functors.hpp +++ b/modules/dnn/src/cuda/functors.hpp @@ -169,16 +169,20 @@ struct SigmoidFunctor { template struct ELUFunctor { struct Params { - CUDA4DNN_HOST_DEVICE Params() { } + CUDA4DNN_HOST_DEVICE Params() : alpha(1) { } + CUDA4DNN_HOST_DEVICE Params(T alpha_) : alpha(alpha_) { } + T alpha; }; - CUDA4DNN_DEVICE ELUFunctor() { } - CUDA4DNN_DEVICE ELUFunctor(const Params& params) { } + CUDA4DNN_DEVICE ELUFunctor() : ELUFunctor(Params{}) { } + CUDA4DNN_DEVICE ELUFunctor(const Params& params) : alpha{params.alpha} { } CUDA4DNN_DEVICE T operator()(T value) { using csl::device::expm1; - return value >= T(0) ? value : expm1(value); + return value >= T(0) ? value : alpha * expm1(value); } + + T alpha; }; template diff --git a/modules/dnn/src/cuda4dnn/kernels/activations.hpp b/modules/dnn/src/cuda4dnn/kernels/activations.hpp index 0fcf7dab8a..d7c471a5ec 100644 --- a/modules/dnn/src/cuda4dnn/kernels/activations.hpp +++ b/modules/dnn/src/cuda4dnn/kernels/activations.hpp @@ -34,7 +34,7 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { void sigmoid(const csl::Stream& stream, csl::Span output, csl::View input); template - void elu(const csl::Stream& stream, csl::Span output, csl::View input); + void elu(const csl::Stream& stream, csl::Span output, csl::View input, T alpha); template void abs(const csl::Stream& stream, csl::Span output, csl::View input); diff --git a/modules/dnn/src/cuda4dnn/primitives/activation.hpp b/modules/dnn/src/cuda4dnn/primitives/activation.hpp index a179db2da5..77a79703fe 100644 --- a/modules/dnn/src/cuda4dnn/primitives/activation.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/activation.hpp @@ -156,15 +156,16 @@ namespace cv { namespace dnn { namespace cuda4dnn { template class ELUOp final : public BaseOp { public: - ELUOp(csl::Stream stream_) : stream(std::move(stream_)) { } + ELUOp(csl::Stream stream_, T alpha_) : stream(std::move(stream_)), alpha(alpha_) { } void calculate(csl::TensorSpan output, csl::TensorView input) const { - kernels::elu(stream, output, input); + kernels::elu(stream, output, input, alpha); } private: csl::Stream stream; + T alpha; }; template diff --git a/modules/dnn/src/layers/elementwise_layers.cpp b/modules/dnn/src/layers/elementwise_layers.cpp index 56e82cc3d1..7cec0d5f7b 100644 --- a/modules/dnn/src/layers/elementwise_layers.cpp +++ b/modules/dnn/src/layers/elementwise_layers.cpp @@ -987,6 +987,9 @@ const char* const SigmoidFunctor::BaseDefaultFunctor::ocl_kernel struct ELUFunctor : public BaseDefaultFunctor { typedef ELULayer Layer; + float alpha; + + explicit ELUFunctor(float alpha_ = 1.f) : alpha(alpha_) {} bool supportBackend(int backendId, int) { @@ -998,13 +1001,18 @@ struct ELUFunctor : public BaseDefaultFunctor inline float calculate(float x) const { - return x >= 0.f ? x : exp(x) - 1.f; + return x >= 0.f ? x : alpha * (exp(x) - 1.f); + } + + inline void setKernelParams(ocl::Kernel& kernel) const + { + kernel.set(3, alpha); } #ifdef HAVE_CUDA Ptr initCUDA(int target, csl::Stream stream) { - return make_cuda_node(target, stream); + return make_cuda_node(target, stream, alpha); } #endif @@ -1012,7 +1020,7 @@ struct ELUFunctor : public BaseDefaultFunctor void attachHalide(const Halide::Expr& input, Halide::Func& top) { Halide::Var x("x"), y("y"), c("c"), n("n"); - top(x, y, c, n) = select(input >= 0.0f, input, exp(input) - 1); + top(x, y, c, n) = select(input >= 0.0f, input, alpha * (exp(input) - 1)); } #endif // HAVE_HALIDE @@ -1026,7 +1034,7 @@ struct ELUFunctor : public BaseDefaultFunctor #ifdef HAVE_DNN_NGRAPH std::shared_ptr initNgraphAPI(const std::shared_ptr& node) { - return std::make_shared(node, 1.0); + return std::make_shared(node, alpha); } #endif // HAVE_DNN_NGRAPH @@ -1856,8 +1864,10 @@ Ptr SigmoidLayer::create(const LayerParams& params) Ptr ELULayer::create(const LayerParams& params) { - Ptr l(new ElementWiseLayer(ELUFunctor())); + float alpha = params.get("alpha", 1.0f); + Ptr l(new ElementWiseLayer(ELUFunctor(alpha))); l->setParamsFrom(params); + l->alpha = alpha; return l; } diff --git a/modules/dnn/src/opencl/activations.cl b/modules/dnn/src/opencl/activations.cl index bc2a105aba..e110160c06 100644 --- a/modules/dnn/src/opencl/activations.cl +++ b/modules/dnn/src/opencl/activations.cl @@ -131,13 +131,14 @@ __kernel void PowForward(const int n, __global const T* in, __global T* out, out[index] = pow(shift + scale * in[index], power); } -__kernel void ELUForward(const int n, __global const T* in, __global T* out) +__kernel void ELUForward(const int n, __global const T* in, __global T* out, + const KERNEL_ARG_DTYPE alpha) { int index = get_global_id(0); if (index < n) { T src = in[index]; - out[index] = (src >= 0.f) ? src : exp(src) - 1; + out[index] = (src >= 0.f) ? src : alpha * (exp(src) - 1); } }