add alpha parameter to ELU

diff --git a/modules/dnn/include/opencv2/dnn/all_layers.hpp b/modules/dnn/include/opencv2/dnn/all_layers.hpp
index c70ce2b..0bec9e3 100644 (file)
--- 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<ELULayer> create(const LayerParams &params);
     };
 

diff --git a/modules/dnn/src/cuda/activations.cu b/modules/dnn/src/cuda/activations.cu
index c38fa03..0980b5d 100644 (file)
--- a/modules/dnn/src/cuda/activations.cu
+++ b/modules/dnn/src/cuda/activations.cu
@@ -119,8 +119,8 @@ void sigmoid(const Stream& stream, Span<T> output, View<T> input) {
 }
 
 template <class T>
-void elu(const Stream& stream, Span<T> output, View<T> input) {
-    generic_op<T, ELUFunctor<T>>(stream, output, input);
+void elu(const Stream& stream, Span<T> output, View<T> input, T alpha) {
+    generic_op<T, ELUFunctor<T>>(stream, output, input, {alpha});
 }
 
 template <class T>
@@ -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<float>(const Stream&, Span<float>, View<float>);
 template void swish<float>(const Stream&, Span<float>, View<float>);
 template void mish<float>(const Stream&, Span<float>, View<float>);
 template void sigmoid<float>(const Stream&, Span<float>, View<float>);
-template void elu<float>(const Stream&, Span<float>, View<float>);
+template void elu<float>(const Stream&, Span<float>, View<float>, float);
 template void abs<float>(const Stream& stream, Span<float> output, View<float> input);
 template void bnll<float>(const Stream&, Span<float>, View<float>);
 template void ceil<float>(const Stream&, Span<float>, View<float>);

diff --git a/modules/dnn/src/cuda/functors.hpp b/modules/dnn/src/cuda/functors.hpp
index 04b545a..98ae175 100644 (file)
--- a/modules/dnn/src/cuda/functors.hpp
+++ b/modules/dnn/src/cuda/functors.hpp
@@ -169,16 +169,20 @@ struct SigmoidFunctor {
 template <class T>
 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 <class T>

diff --git a/modules/dnn/src/cuda4dnn/kernels/activations.hpp b/modules/dnn/src/cuda4dnn/kernels/activations.hpp
index 0fcf7da..d7c471a 100644 (file)
--- 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<T> output, csl::View<T> input);
 
     template <class T>
-    void elu(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);
+    void elu(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input, T alpha);
 
     template <class T>
     void abs(const csl::Stream& stream, csl::Span<T> output, csl::View<T> input);

diff --git a/modules/dnn/src/cuda4dnn/primitives/activation.hpp b/modules/dnn/src/cuda4dnn/primitives/activation.hpp
index a179db2..77a7970 100644 (file)
--- 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 T>
     class ELUOp final : public BaseOp<ELUOp, T> {
     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<T> output, csl::TensorView<T> input) const
         {
-            kernels::elu<T>(stream, output, input);
+            kernels::elu<T>(stream, output, input, alpha);
         }
 
     private:
         csl::Stream stream;
+        T alpha;
     };
 
     template <class T>

diff --git a/modules/dnn/src/layers/elementwise_layers.cpp b/modules/dnn/src/layers/elementwise_layers.cpp
index 56e82cc..7cec0d5 100644 (file)
--- 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<SigmoidFunctor>::ocl_kernel
 struct ELUFunctor : public BaseDefaultFunctor<ELUFunctor>
 {
     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<ELUFunctor>
 
     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<BackendNode> initCUDA(int target, csl::Stream stream)
     {
-        return make_cuda_node<cuda4dnn::ELUOp>(target, stream);
+        return make_cuda_node<cuda4dnn::ELUOp>(target, stream, alpha);
     }
 #endif
 
@@ -1012,7 +1020,7 @@ struct ELUFunctor : public BaseDefaultFunctor<ELUFunctor>
     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<ELUFunctor>
 #ifdef HAVE_DNN_NGRAPH
     std::shared_ptr<ngraph::Node> initNgraphAPI(const std::shared_ptr<ngraph::Node>& node)
     {
-        return std::make_shared<ngraph::op::Elu>(node, 1.0);
+        return std::make_shared<ngraph::op::Elu>(node, alpha);
     }
 #endif  // HAVE_DNN_NGRAPH
 
@@ -1856,8 +1864,10 @@ Ptr<SigmoidLayer> SigmoidLayer::create(const LayerParams& params)
 
 Ptr<ELULayer> ELULayer::create(const LayerParams& params)
 {
-    Ptr<ELULayer> l(new ElementWiseLayer<ELUFunctor>(ELUFunctor()));
+    float alpha = params.get<float>("alpha", 1.0f);
+    Ptr<ELULayer> l(new ElementWiseLayer<ELUFunctor>(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 bc2a105..e110160 100644 (file)
--- 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);
   }
 }