diff --git a/modules/dnn/include/opencv2/dnn/all_layers.hpp b/modules/dnn/include/opencv2/dnn/all_layers.hpp index 2acb41076d..dae8701970 100644 --- a/modules/dnn/include/opencv2/dnn/all_layers.hpp +++ b/modules/dnn/include/opencv2/dnn/all_layers.hpp @@ -794,6 +794,26 @@ CV__DNN_INLINE_NS_BEGIN static Ptr create(const LayerParams ¶ms); }; + class CV_EXPORTS SignLayer : public ActivationLayer + { + public: + static Ptr create(const LayerParams ¶ms); + }; + + class CV_EXPORTS ShrinkLayer : public ActivationLayer + { + public: + float bias; + float lambd; + static Ptr create(const LayerParams ¶ms); + }; + + class CV_EXPORTS ReciprocalLayer : public ActivationLayer + { + public: + static Ptr create(const LayerParams ¶ms); + }; + /* Layers used in semantic segmentation */ class CV_EXPORTS CropLayer : public Layer diff --git a/modules/dnn/src/cuda/activations.cu b/modules/dnn/src/cuda/activations.cu index f5dafcea7f..ed34d57e0b 100644 --- a/modules/dnn/src/cuda/activations.cu +++ b/modules/dnn/src/cuda/activations.cu @@ -248,6 +248,21 @@ void selu(const Stream& stream, Span output, View input, T alpha, T gamma) generic_op>(stream, output, input, {alpha, gamma}); } +template +void sign(const Stream& stream, Span output, View input) { + generic_op>(stream, output, input); +} + +template +void shrink(const Stream& stream, Span output, View input, T bias, T lambd) { + generic_op>(stream, output, input, {bias, lambd}); +} + +template +void reciprocal(const Stream& stream, Span output, View input) { + generic_op>(stream, output, input); +} + template void thresholdedrelu(const Stream& stream, Span output, View input, T alpha) { generic_op>(stream, output, input, {alpha}); @@ -312,6 +327,9 @@ template void selu<__half>(const Stream&, Span<__half>, View<__half>, __half, __ template void thresholdedrelu<__half>(const Stream&, Span<__half>, View<__half>, __half); template void power<__half>(const Stream&, Span<__half>, View<__half>, __half, __half, __half); template void exp<__half>(const Stream&, Span<__half>, View<__half>, __half, __half); +template void sign<__half>(const Stream&, Span<__half>, View<__half>); +template void shrink<__half>(const Stream&, Span<__half>, View<__half>, __half, __half); +template void reciprocal<__half>(const Stream&, Span<__half>, View<__half>); #endif @@ -351,6 +369,9 @@ template void selu(const Stream&, Span, View, float, float) template void thresholdedrelu(const Stream&, Span, View, float); template void power(const Stream&, Span, View, float, float, float); template void exp(const Stream&, Span, View, float, float); +template void sign(const Stream&, Span, View); +template void shrink(const Stream&, Span, View, float, float); +template void reciprocal(const Stream&, Span, View); template static void launch_vectorized_axiswise_relu(const Stream& stream, Span output, View input, std::size_t inner_size, View slope) { diff --git a/modules/dnn/src/cuda/functors.hpp b/modules/dnn/src/cuda/functors.hpp index 640c7c8ad6..378df82f26 100644 --- a/modules/dnn/src/cuda/functors.hpp +++ b/modules/dnn/src/cuda/functors.hpp @@ -726,6 +726,50 @@ struct DivFunctor { CUDA4DNN_DEVICE T operator()(T x, T y) { return x / y; } }; +template +struct SignFunctor { + struct Params { + CUDA4DNN_HOST_DEVICE Params() {} + }; + + CUDA4DNN_DEVICE SignFunctor() : SignFunctor(Params{}) { } + + CUDA4DNN_DEVICE T operator()(T value) { + return value > T(0) ? T(1) : (value < T(0) ? T(-1) : T(0)); + } +}; + +template +struct ShrinkFunctor { + struct Params { + CUDA4DNN_HOST_DEVICE Params() : bias(0), lambd(0.5) { } + CUDA4DNN_HOST_DEVICE Params(T bias_, T lambd_) : bias(bias_), lambd(lambd_) { } + T bias, lambd; + }; + + CUDA4DNN_DEVICE ShrinkFunctor() : bias(0), lambd(0.5) { } + CUDA4DNN_DEVICE ShrinkFunctor(const Params& params) : bias{params.bias}, lambd{params.lambd} { } + + CUDA4DNN_DEVICE T operator()(T value) { + return value > lambd ? value - bias : (value < -lambd ? value + bias : T(0)); + } + + T bias, lambd; +}; + +template +struct ReciprocalFunctor { + struct Params { + CUDA4DNN_HOST_DEVICE Params() {} + }; + + CUDA4DNN_DEVICE ReciprocalFunctor() : ReciprocalFunctor(Params{}) { } + + CUDA4DNN_DEVICE T operator()(T value) { + return T(1.0f)/value; + } +}; + }}}} /* namespace cv::dnn::cuda4dnn::kernels */ #endif /* OPENCV_DNN_SRC_CUDA_FUNCTORS_HPP */ diff --git a/modules/dnn/src/cuda4dnn/kernels/activations.hpp b/modules/dnn/src/cuda4dnn/kernels/activations.hpp index ef1f6da3e6..6958b93d5e 100644 --- a/modules/dnn/src/cuda4dnn/kernels/activations.hpp +++ b/modules/dnn/src/cuda4dnn/kernels/activations.hpp @@ -123,6 +123,14 @@ namespace cv { namespace dnn { namespace cuda4dnn { namespace kernels { template void exp(const csl::Stream& stream, csl::Span output, csl::View input, T normScale, T normShift); + template + void sign(const csl::Stream& stream, csl::Span output, csl::View input); + + template + void shrink(const csl::Stream& stream, csl::Span output, csl::View input, T bias, T lambd); + + template + void reciprocal(const csl::Stream& stream, csl::Span output, csl::View input); }}}} /* namespace cv::dnn::cuda4dnn::kernels */ #endif /* OPENCV_DNN_SRC_CUDA4DNN_KERNELS_ACTIVATIONS_HPP */ diff --git a/modules/dnn/src/cuda4dnn/primitives/activation.hpp b/modules/dnn/src/cuda4dnn/primitives/activation.hpp index 39ebf513a7..564202e8c0 100644 --- a/modules/dnn/src/cuda4dnn/primitives/activation.hpp +++ b/modules/dnn/src/cuda4dnn/primitives/activation.hpp @@ -584,6 +584,52 @@ namespace cv { namespace dnn { namespace cuda4dnn { const T normScale, normShift; }; + template + class ShrinkOp final : public BaseOp { + public: + ShrinkOp(csl::Stream stream_, T bias_, T lambd_) + : stream(std::move(stream_)), bias{ bias_ }, lambd{ lambd_ } { } + + void calculate(csl::TensorSpan output, csl::TensorView input) const + { + kernels::shrink(stream, output, input, bias, lambd); + } + + private: + csl::Stream stream; + const T bias, lambd; + }; + + template + class SignOp final : public BaseOp { + public: + SignOp(csl::Stream stream_) + : stream(std::move(stream_)) { } + + void calculate(csl::TensorSpan output, csl::TensorView input) const + { + kernels::sign(stream, output, input); + } + + private: + csl::Stream stream; + }; + + template + class ReciprocalOp final : public BaseOp { + public: + ReciprocalOp(csl::Stream stream_) + : stream(std::move(stream_)) { } + + void calculate(csl::TensorSpan output, csl::TensorView input) const + { + kernels::reciprocal(stream, output, input); + } + + private: + csl::Stream stream; + }; + }}} /* namespace cv::dnn::cuda4dnn */ #endif /* OPENCV_DNN_SRC_CUDA4DNN_PRIMITIVES_ACTIVATION_HPP */ diff --git a/modules/dnn/src/init.cpp b/modules/dnn/src/init.cpp index 86ceba382e..6979d1864d 100644 --- a/modules/dnn/src/init.cpp +++ b/modules/dnn/src/init.cpp @@ -130,6 +130,8 @@ void initializeLayerFactory() CV_DNN_REGISTER_LAYER_CLASS(HardSwish, HardSwishLayer); CV_DNN_REGISTER_LAYER_CLASS(Sin, SinLayer); CV_DNN_REGISTER_LAYER_CLASS(Sinh, SinhLayer); + CV_DNN_REGISTER_LAYER_CLASS(Sign, SignLayer); + CV_DNN_REGISTER_LAYER_CLASS(Shrink, ShrinkLayer); CV_DNN_REGISTER_LAYER_CLASS(Softplus, SoftplusLayer); CV_DNN_REGISTER_LAYER_CLASS(Softsign, SoftsignLayer); CV_DNN_REGISTER_LAYER_CLASS(Tan, TanLayer); @@ -144,6 +146,7 @@ void initializeLayerFactory() CV_DNN_REGISTER_LAYER_CLASS(Silence, BlankLayer); CV_DNN_REGISTER_LAYER_CLASS(Const, ConstLayer); CV_DNN_REGISTER_LAYER_CLASS(Arg, ArgLayer); + CV_DNN_REGISTER_LAYER_CLASS(Reciprocal, ReciprocalLayer); CV_DNN_REGISTER_LAYER_CLASS(Crop, CropLayer); CV_DNN_REGISTER_LAYER_CLASS(Eltwise, EltwiseLayer); diff --git a/modules/dnn/src/layers/elementwise_layers.cpp b/modules/dnn/src/layers/elementwise_layers.cpp index 0accbe0fbb..eb2ecce3ce 100644 --- a/modules/dnn/src/layers/elementwise_layers.cpp +++ b/modules/dnn/src/layers/elementwise_layers.cpp @@ -2270,6 +2270,96 @@ struct ChannelsPReLUFunctor : public BaseFunctor int64 getFLOPSPerElement() const { return 1; } }; +struct SignFunctor : public BaseDefaultFunctor +{ + typedef SignLayer Layer; + + bool supportBackend(int backendId, int) + { + return backendId == DNN_BACKEND_OPENCV || + backendId == DNN_BACKEND_CUDA; + } + + inline float calculate(float x) const + { + return x > 0 ? 1 : (x < 0 ? -1 : 0); + } + +#ifdef HAVE_CUDA + Ptr initCUDA(int target, csl::Stream stream) + { + return make_cuda_node(target, stream); + } +#endif + + int64 getFLOPSPerElement() const { return 1; } +}; + +template<> +const char* const SignFunctor::BaseDefaultFunctor::ocl_kernel_name = "SignForward"; + + +struct ShrinkFunctor : public BaseDefaultFunctor +{ + typedef ShrinkLayer Layer; + float bias; + float lambd; + + explicit ShrinkFunctor(float bias_ = 0.0f, float lambd_ = 0.5f) : bias(bias_), lambd(lambd_) {} + + bool supportBackend(int backendId, int) + { + return backendId == DNN_BACKEND_OPENCV || + backendId == DNN_BACKEND_CUDA; + } + + inline float calculate(float x) const + { + return x > lambd ? x - bias : (x < -lambd ? x + bias : 0); + } + +#ifdef HAVE_CUDA + Ptr initCUDA(int target, csl::Stream stream) + { + return make_cuda_node(target, stream); + } +#endif + + int64 getFLOPSPerElement() const { return 1; } +}; + +template<> +const char* const ShrinkFunctor::BaseDefaultFunctor::ocl_kernel_name = "ShrinkForward"; + +struct ReciprocalFunctor : public BaseDefaultFunctor +{ + typedef ReciprocalLayer Layer; + + bool supportBackend(int backendId, int) + { + return backendId == DNN_BACKEND_OPENCV || + backendId == DNN_BACKEND_CUDA; + } + + inline float calculate(float x) const + { + return 1.0/x; + } + +#ifdef HAVE_CUDA + Ptr initCUDA(int target, csl::Stream stream) + { + return make_cuda_node(target, stream); + } +#endif + + int64 getFLOPSPerElement() const { return 1; } +}; + +template<> +const char* const ReciprocalFunctor::BaseDefaultFunctor::ocl_kernel_name = "ReciprocalForward"; + + #define ACTIVATION_CREATOR_FOR(_Layer, _Functor, ...) \ Ptr<_Layer> _Layer::create() { \ return return Ptr<_Layer>( new ElementWiseLayer<_Functor>(_Functor()) ); } @@ -2611,5 +2701,32 @@ Ptr ChannelsPReLULayer::create(const LayerParams& params) return l; } +Ptr SignLayer::create(const LayerParams& params) +{ + Ptr l(new ElementWiseLayer()); + l->setParamsFrom(params); + + return l; +} + +Ptr ReciprocalLayer::create(const LayerParams& params) +{ + Ptr l(new ElementWiseLayer()); + l->setParamsFrom(params); + + return l; +} + +Ptr ShrinkLayer::create(const LayerParams& params) +{ + float bias = params.get("bias", 0.f); + float lambd = params.get("lambd", 0.5f); + Ptr l(new ElementWiseLayer(ShrinkFunctor(bias, lambd))); + l->setParamsFrom(params); + l->bias = bias; + l->lambd = lambd; + + return l; +} } } diff --git a/modules/dnn/src/onnx/onnx_importer.cpp b/modules/dnn/src/onnx/onnx_importer.cpp index 6cf0fe8f57..2a440a1284 100644 --- a/modules/dnn/src/onnx/onnx_importer.cpp +++ b/modules/dnn/src/onnx/onnx_importer.cpp @@ -3675,8 +3675,8 @@ void ONNXImporter::buildDispatchMap_ONNX_AI(int opset_version) std::vector simpleLayers{"Acos", "Acosh", "Asin", "Asinh", "Atan", "Atanh", "Ceil", "Celu", "Cos", "Cosh", "Dropout", "Erf", "Exp", "Floor", "HardSigmoid", "HardSwish", - "Identity", "Log", "Round", "Selu", "Sigmoid", "Sin", "Sinh", "Softmax", - "Softplus", "Softsign", "Sqrt", "Tan", "ThresholdedRelu"}; + "Identity", "Log", "Round", "Reciprocal", "Selu", "Sign", "Sigmoid", "Sin", "Sinh", "Softmax", + "Softplus", "Softsign", "Shrink", "Sqrt", "Tan", "ThresholdedRelu"}; for (const auto& name : simpleLayers) { dispatch[name] = &ONNXImporter::parseSimpleLayers; diff --git a/modules/dnn/src/opencl/activations.cl b/modules/dnn/src/opencl/activations.cl index 3e99f18570..d105623403 100644 --- a/modules/dnn/src/opencl/activations.cl +++ b/modules/dnn/src/opencl/activations.cl @@ -306,3 +306,26 @@ __kernel void ThresholdedReluForward(const int n, __global T* in, __global T* ou if(index < n) out[index] = (in[index] > alpha ? in[index] : 0.f); } + +__kernel void ShrinkForward(const int n, __global T* in, __global T* out, + const KERNEL_ARG_DTYPE bias, + const KERNEL_ARG_DTYPE lambd) +{ + int index = get_global_id(0); + if(index < n) + out[index] = in[index] < -lambd ? in[index] + bias : (in[index] > lambd ? in[index] - bias : 0.f); +} + +__kernel void SignForward(const int n, __global T* in, __global T* out) +{ + int index = get_global_id(0); + if(index < n) + out[index] = in[index] > 0.f ? 1.0f : (in[index] < 0.f) ? -1.0f : 0.0f); +} + +__kernel void ReciprocalForward(const int n, __global T* in, __global T* out) +{ + int index = get_global_id(0); + if(index < n) + out[index] = 1.0f/in[index]; +} \ No newline at end of file diff --git a/modules/dnn/test/test_onnx_conformance_layer_parser_denylist.inl.hpp b/modules/dnn/test/test_onnx_conformance_layer_parser_denylist.inl.hpp index eef421491c..49db810c7f 100644 --- a/modules/dnn/test/test_onnx_conformance_layer_parser_denylist.inl.hpp +++ b/modules/dnn/test/test_onnx_conformance_layer_parser_denylist.inl.hpp @@ -337,8 +337,6 @@ "test_range_float_type_positive_delta_expanded", "test_range_int32_type_negative_delta", "test_range_int32_type_negative_delta_expanded", -"test_reciprocal", -"test_reciprocal_example", "test_reduce_sum_default_axes_keepdims_example", "test_reduce_sum_default_axes_keepdims_random", "test_reduce_sum_do_not_keepdims_example", @@ -479,9 +477,6 @@ "test_shape_start_1_end_2", "test_shape_start_1_end_negative_1", "test_shape_start_negative_1", -"test_shrink_hard", -"test_shrink_soft", -"test_sign", "test_simple_rnn_batchwise", "test_simple_rnn_defaults", "test_simple_rnn_with_initial_bias",