From 2674cf0c48b32bd815c9dea73c63e21993380be0 Mon Sep 17 00:00:00 2001 From: thnkinbtfly Date: Mon, 14 Dec 2020 05:37:51 +0000 Subject: [PATCH] support swish activation --- include/tkDNN/Layer.h | 7 ++- include/tkDNN/NetworkRT.h | 1 + include/tkDNN/kernels.h | 1 + include/tkDNN/pluginsRT/ActivationSwishRT.h | 60 +++++++++++++++++++++ src/Activation.cpp | 3 ++ src/DarknetParser.cpp | 1 + src/NetworkRT.cpp | 13 ++++- src/kernels/activation_swish.cu | 27 ++++++++++ 8 files changed, 111 insertions(+), 2 deletions(-) create mode 100644 include/tkDNN/pluginsRT/ActivationSwishRT.h create mode 100644 src/kernels/activation_swish.cu diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 25c4565..9d6499a 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -19,6 +19,7 @@ enum layerType_t { LAYER_ACTIVATION_CRELU, LAYER_ACTIVATION_LEAKY, LAYER_ACTIVATION_MISH, + LAYER_ACTIVATION_SWISH, LAYER_FLATTEN, LAYER_RESHAPE, LAYER_MULADD, @@ -68,6 +69,7 @@ public: case LAYER_ACTIVATION_CRELU: return "ActivationCReLU"; case LAYER_ACTIVATION_LEAKY: return "ActivationLeaky"; case LAYER_ACTIVATION_MISH: return "ActivationMish"; + case LAYER_ACTIVATION_SWISH: return "ActivationSwish"; case LAYER_FLATTEN: return "Flatten"; case LAYER_RESHAPE: return "Reshape"; case LAYER_MULADD: return "MulAdd"; @@ -212,7 +214,8 @@ public: typedef enum { ACTIVATION_ELU = 100, ACTIVATION_LEAKY = 101, - ACTIVATION_MISH = 102 + ACTIVATION_MISH = 102, + ACTIVATION_SWISH = 103 } tkdnnActivationMode_t; /** @@ -233,6 +236,8 @@ public: return LAYER_ACTIVATION_LEAKY; else if (act_mode == ACTIVATION_MISH) return LAYER_ACTIVATION_MISH; + else if (act_mode == ACTIVATION_SWISH) + return LAYER_ACTIVATION_SWISH; else return LAYER_ACTIVATION; }; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 4c6c816..6efff9b 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -26,6 +26,7 @@ using namespace nvinfer1; #include "pluginsRT/ActivationLeakyRT.h" #include "pluginsRT/ActivationReLUCeilingRT.h" #include "pluginsRT/ActivationMishRT.h" +#include "pluginsRT/ActivationSwishRT.h" #include "pluginsRT/ReorgRT.h" #include "pluginsRT/RegionRT.h" #include "pluginsRT/RouteRT.h" diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index 5d673c8..ebb0dbb 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -9,6 +9,7 @@ void activationReLUCeilingForward(dnnType *srcData, dnnType *dstData, int size, void activationLOGISTICForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationSIGMOIDForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationMishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream= cudaStream_t(0)); +void activationSwishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream= cudaStream_t(0)); void fill(dnnType *data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0)); diff --git a/include/tkDNN/pluginsRT/ActivationSwishRT.h b/include/tkDNN/pluginsRT/ActivationSwishRT.h new file mode 100644 index 0000000..9796e1f --- /dev/null +++ b/include/tkDNN/pluginsRT/ActivationSwishRT.h @@ -0,0 +1,60 @@ +#include +#include "../kernels.h" + +class ActivationSwishRT : public IPlugin { + +public: + ActivationSwishRT() { + + + } + + ~ActivationSwishRT(){ + + } + + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return inputs[0]; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + size = 1; + for(int i=0; i(inputs[0]), + reinterpret_cast(outputs[0]), batchSize*size, stream); + return 0; + } + + + virtual size_t getSerializationSize() override { + return 1*sizeof(int); + } + + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer); + tk::dnn::writeBUF(buf, size); + } + + int size; +}; diff --git a/src/Activation.cpp b/src/Activation.cpp index 28c7624..543e29a 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -52,6 +52,9 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) { else if(act_mode == ACTIVATION_MISH) { activationMishForward(srcData, dstData, dim.tot()); + } + else if(act_mode == ACTIVATION_SWISH) { + activationSwishForward(srcData, dstData, dim.tot()); } else { dnnType alpha = dnnType(1); dnnType beta = dnnType(0); diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index 891c97c..a2441fe 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -192,6 +192,7 @@ namespace tk { namespace dnn { else if(f.activation == "logistic") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_SIGMOID); else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY; else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH; + else if(f.activation == "swish") act = tk::dnn::ACTIVATION_SWISH; else { FatalError("activation not supported: " + f.activation); } netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act); }; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 501ade4..83ddf40 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -226,7 +226,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { return convert_layer(input, (Conv2d*) l); if(type == LAYER_POOLING) return convert_layer(input, (Pooling*) l); - if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH) + if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_SWISH) return convert_layer(input, (Activation*) l); if(type == LAYER_SOFTMAX) return convert_layer(input, (Softmax*) l); @@ -421,6 +421,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) { checkNULL(lRT); return lRT; } + else if(l->act_mode == ACTIVATION_SWISH) { + IPlugin *plugin = new ActivationSwishRT(); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + checkNULL(lRT); + return lRT; + } else { FatalError("this Activation mode is not yet implemented"); return NULL; @@ -653,6 +659,11 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa a->size = readBUF(buf); return a; } + if(name.find("ActivationSwish") == 0) { + ActivationSwishRT *a = new ActivationSwishRT(); + a->size = readBUF(buf); + return a; + } if(name.find("ActivationCReLU") == 0) { ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF(buf)); a->size = readBUF(buf); diff --git a/src/kernels/activation_swish.cu b/src/kernels/activation_swish.cu new file mode 100644 index 0000000..2c7e2e9 --- /dev/null +++ b/src/kernels/activation_swish.cu @@ -0,0 +1,27 @@ +#include "kernels.h" +#include + +// https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu +__device__ float logistic_activate_kernel(float x){return 1.f/(1.f + expf(-x));} + +__global__ +void activation_swish(dnnType *input, dnnType *output, int size) { + int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; + if (i < size) + { + float x_val = input[i]; + float sigmoid = logistic_activate_kernel(x_val); + output[i] = x_val * sigmoid; + } +} + +/** + swish activation function +*/ +void activationSwishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream) +{ + int blocks = (size+255)/256; + int threads = 256; + + activation_swish<<>>(srcData, dstData, size); +} \ No newline at end of file