From 8a18d132fc49534c963f72218f49a9082ae6441d Mon Sep 17 00:00:00 2001 From: thebhatman Date: Sun, 1 Dec 2019 11:00:58 +0300 Subject: [PATCH] Port Swish and Mish layers --- .../dnn/include/opencv2/dnn/all_layers.hpp | 12 ++ modules/dnn/src/init.cpp | 2 + modules/dnn/src/layers/elementwise_layers.cpp | 162 ++++++++++++++++++ modules/dnn/src/opencl/activations.cl | 12 ++ modules/dnn/test/test_halide_layers.cpp | 2 +- 5 files changed, 189 insertions(+), 1 deletion(-) diff --git a/modules/dnn/include/opencv2/dnn/all_layers.hpp b/modules/dnn/include/opencv2/dnn/all_layers.hpp index da2f382bb1..363dcc01cb 100644 --- a/modules/dnn/include/opencv2/dnn/all_layers.hpp +++ b/modules/dnn/include/opencv2/dnn/all_layers.hpp @@ -462,6 +462,18 @@ CV__DNN_EXPERIMENTAL_NS_BEGIN static Ptr create(const LayerParams ¶ms); }; + class CV_EXPORTS SwishLayer : public ActivationLayer + { + public: + static Ptr create(const LayerParams ¶ms); + }; + + class CV_EXPORTS MishLayer : public ActivationLayer + { + public: + static Ptr create(const LayerParams ¶ms); + }; + class CV_EXPORTS SigmoidLayer : public ActivationLayer { public: diff --git a/modules/dnn/src/init.cpp b/modules/dnn/src/init.cpp index 7f6c831f70..07db56eb0b 100644 --- a/modules/dnn/src/init.cpp +++ b/modules/dnn/src/init.cpp @@ -103,6 +103,8 @@ void initializeLayerFactory() CV_DNN_REGISTER_LAYER_CLASS(PReLU, ChannelsPReLULayer); CV_DNN_REGISTER_LAYER_CLASS(Sigmoid, SigmoidLayer); CV_DNN_REGISTER_LAYER_CLASS(TanH, TanHLayer); + CV_DNN_REGISTER_LAYER_CLASS(Swish, SwishLayer); + CV_DNN_REGISTER_LAYER_CLASS(Mish, MishLayer); CV_DNN_REGISTER_LAYER_CLASS(ELU, ELULayer); CV_DNN_REGISTER_LAYER_CLASS(BNLL, BNLLLayer); CV_DNN_REGISTER_LAYER_CLASS(AbsVal, AbsLayer); diff --git a/modules/dnn/src/layers/elementwise_layers.cpp b/modules/dnn/src/layers/elementwise_layers.cpp index e89c539f68..2ea1302b82 100644 --- a/modules/dnn/src/layers/elementwise_layers.cpp +++ b/modules/dnn/src/layers/elementwise_layers.cpp @@ -534,6 +534,152 @@ struct TanHFunctor int64 getFLOPSPerElement() const { return 1; } }; +struct SwishFunctor +{ + typedef SwishLayer Layer; + + bool supportBackend(int backendId, int) + { + return backendId == DNN_BACKEND_OPENCV || + backendId == DNN_BACKEND_HALIDE; + } + + void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const + { + for( int cn = cn0; cn < cn1; cn++, srcptr += planeSize, dstptr += planeSize ) + { + for( int i = 0; i < len; i++ ) + { + float x = srcptr[i]; + dstptr[i] = x / (1.0f + exp(-x)); + } + } + } + +#ifdef HAVE_OPENCL + bool applyOCL(InputArrayOfArrays inps, OutputArrayOfArrays outs, OutputArrayOfArrays internals) + { + std::vector inputs; + std::vector outputs; + + inps.getUMatVector(inputs); + outs.getUMatVector(outputs); + String buildopt = oclGetTMacro(inputs[0]); + + for (size_t i = 0; i < inputs.size(); i++) + { + UMat& src = inputs[i]; + UMat& dst = outputs[i]; + + ocl::Kernel kernel("SwishForward", ocl::dnn::activations_oclsrc, buildopt); + kernel.set(0, (int)src.total()); + kernel.set(1, ocl::KernelArg::PtrReadOnly(src)); + kernel.set(2, ocl::KernelArg::PtrWriteOnly(dst)); + + size_t gSize = src.total(); + CV_Assert(kernel.run(1, &gSize, NULL, false)); + } + + return true; + } +#endif + +#ifdef HAVE_HALIDE + void attachHalide(const Halide::Expr& input, Halide::Func& top) + { + Halide::Var x("x"), y("y"), c("c"), n("n"); + top(x, y, c, n) = input / (1.0f + exp(-input)); + } +#endif // HAVE_HALIDE + +#ifdef HAVE_INF_ENGINE + InferenceEngine::Builder::Layer initInfEngineBuilderAPI() + { + CV_Error(Error::StsNotImplemented, ""); + } +#endif // HAVE_INF_ENGINE + + bool tryFuse(Ptr&) { return false; } + + void getScaleShift(Mat&, Mat&) const {} + + int64 getFLOPSPerElement() const { return 3; } + +}; + +struct MishFunctor +{ + typedef MishLayer Layer; + + bool supportBackend(int backendId, int) + { + return backendId == DNN_BACKEND_OPENCV || + backendId == DNN_BACKEND_HALIDE; + } + + void apply(const float* srcptr, float* dstptr, int len, size_t planeSize, int cn0, int cn1) const + { + for( int cn = cn0; cn < cn1; cn++, srcptr += planeSize, dstptr += planeSize ) + { + for( int i = 0; i < len; i++ ) + { + float x = srcptr[i]; + dstptr[i] = x * tanh(log(1.0f + exp(x))); + } + } + } + +#ifdef HAVE_OPENCL + bool applyOCL(InputArrayOfArrays inps, OutputArrayOfArrays outs, OutputArrayOfArrays internals) + { + std::vector inputs; + std::vector outputs; + + inps.getUMatVector(inputs); + outs.getUMatVector(outputs); + String buildopt = oclGetTMacro(inputs[0]); + + for (size_t i = 0; i < inputs.size(); i++) + { + UMat& src = inputs[i]; + UMat& dst = outputs[i]; + + ocl::Kernel kernel("MishForward", ocl::dnn::activations_oclsrc, buildopt); + kernel.set(0, (int)src.total()); + kernel.set(1, ocl::KernelArg::PtrReadOnly(src)); + kernel.set(2, ocl::KernelArg::PtrWriteOnly(dst)); + + size_t gSize = src.total(); + CV_Assert(kernel.run(1, &gSize, NULL, false)); + } + + return true; + } +#endif + +#ifdef HAVE_HALIDE + void attachHalide(const Halide::Expr& input, Halide::Func& top) + { + Halide::Var x("x"), y("y"), c("c"), n("n"); + top(x, y, c, n) = input * tanh(log(1.0f + exp(input))); + } +#endif // HAVE_HALIDE + +#ifdef HAVE_INF_ENGINE + InferenceEngine::Builder::Layer initInfEngineBuilderAPI() + { + CV_Error(Error::StsNotImplemented, ""); + } +#endif // HAVE_INF_ENGINE + + bool tryFuse(Ptr&) { return false; } + + void getScaleShift(Mat&, Mat&) const {} + + int64 getFLOPSPerElement() const { return 3; } + +}; + struct SigmoidFunctor { typedef SigmoidLayer Layer; @@ -1111,6 +1257,22 @@ Ptr TanHLayer::create(const LayerParams& params) return l; } +Ptr SwishLayer::create(const LayerParams& params) +{ + Ptr l(new ElementWiseLayer()); + l->setParamsFrom(params); + + return l; +} + +Ptr MishLayer::create(const LayerParams& params) +{ + Ptr l(new ElementWiseLayer()); + l->setParamsFrom(params); + + return l; +} + Ptr SigmoidLayer::create(const LayerParams& params) { Ptr l(new ElementWiseLayer()); diff --git a/modules/dnn/src/opencl/activations.cl b/modules/dnn/src/opencl/activations.cl index ff9d2401d3..b900e6add6 100644 --- a/modules/dnn/src/opencl/activations.cl +++ b/modules/dnn/src/opencl/activations.cl @@ -95,6 +95,18 @@ __kernel void SigmoidForward(const int count, __global const T* in, __global T* out[index] = 1.0f / (1.0f + exp(-in[index])); } +__kernel void SwishForward(const int count, __global const T* in, __global T* out) { + int index = get_global_id(0); + if(index < count) + out[index] = in[index] / (1.0f + exp(-in[index])); +} + +__kernel void MishForward(const int count, __global const T* in, __global T* out) { + int index = get_global_id(0); + if(index < count) + out[index] = in[index] * tanh(log(1.0f + exp(in[index]))); +} + __kernel void BNLLForward(const int n, __global const T* in, __global T* out) { int index = get_global_id(0); if (index < n) { diff --git a/modules/dnn/test/test_halide_layers.cpp b/modules/dnn/test/test_halide_layers.cpp index 0cee157b15..a9eff9faa2 100644 --- a/modules/dnn/test/test_halide_layers.cpp +++ b/modules/dnn/test/test_halide_layers.cpp @@ -583,7 +583,7 @@ TEST_P(NoParamActivation, Accuracy) testInPlaceActivation(lp, backendId, targetId); } INSTANTIATE_TEST_CASE_P(Layer_Test_Halide, NoParamActivation, Combine( -/*type*/ Values("TanH", "Sigmoid", "AbsVal", "BNLL"), +/*type*/ Values("TanH", "Sigmoid", "AbsVal", "BNLL", "Swish", "Mish"), dnnBackendsAndTargetsWithHalide() ));