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 9ce2fdb96f..7cec0d5f7b 100644 --- a/modules/dnn/src/layers/elementwise_layers.cpp +++ b/modules/dnn/src/layers/elementwise_layers.cpp @@ -1012,7 +1012,7 @@ struct ELUFunctor : public BaseDefaultFunctor #ifdef HAVE_CUDA Ptr initCUDA(int target, csl::Stream stream) { - return make_cuda_node(target, stream); + return make_cuda_node(target, stream, alpha); } #endif