From 000af24d29e1d406f55f3d5f24aa4b01014bfeb4 Mon Sep 17 00:00:00 2001 From: "Cuellar, Adam T" Date: Tue, 27 Jul 2021 13:09:43 +0000 Subject: [PATCH] Add swish --- include/tkDNN/Layer.h | 7 ++- include/tkDNN/NetworkRT.h | 1 + include/tkDNN/kernels.h | 1 + include/tkDNN/pluginsRT/ActivationSwishRT.h | 61 +++++++++++++++++++++ src/Activation.cpp | 4 ++ src/DarknetParser.cpp | 1 + src/NetworkRT.cpp | 14 ++++- src/kernels/activation_swish.cu | 23 ++++++++ 8 files changed, 110 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 d1234a5..a8bcecc 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_ACTIVATION_LOGISTIC, LAYER_FLATTEN, LAYER_RESHAPE, @@ -74,6 +75,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_ACTIVATION_LOGISTIC: return "ActivationLogistic"; case LAYER_FLATTEN: return "Flatten"; case LAYER_RESHAPE: return "Reshape"; @@ -221,7 +223,8 @@ typedef enum { ACTIVATION_ELU = 100, ACTIVATION_LEAKY = 101, ACTIVATION_MISH = 102, - ACTIVATION_LOGISTIC = 103 + ACTIVATION_LOGISTIC = 103, + ACTIVATION_SWISH = 104 } tkdnnActivationMode_t; /** @@ -243,6 +246,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 if (act_mode == ACTIVATION_LOGISTIC) return LAYER_ACTIVATION_LOGISTIC; else diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 9892a24..e5bdd39 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -28,6 +28,7 @@ using namespace nvinfer1; #include "pluginsRT/ActivationLogisticRT.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 d809129..52f8bb3 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..123ff1f --- /dev/null +++ b/include/tkDNN/pluginsRT/ActivationSwishRT.h @@ -0,0 +1,61 @@ +#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),*a=buf; + tk::dnn::writeBUF(buf, size); + assert(buf == a + getSerializationSize()); + } + + int size; +}; diff --git a/src/Activation.cpp b/src/Activation.cpp index 0b113a7..62c26bd 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -52,6 +52,10 @@ 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 if(act_mode == ACTIVATION_LOGISTIC) { activationLOGISTICForward(srcData, dstData, dim.tot()); diff --git a/src/DarknetParser.cpp b/src/DarknetParser.cpp index 69b6b29..247289c 100644 --- a/src/DarknetParser.cpp +++ b/src/DarknetParser.cpp @@ -187,6 +187,7 @@ namespace tk { namespace dnn { if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU); 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 if(f.activation == "logistic") act = tk::dnn::ACTIVATION_LOGISTIC; 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 6ac7235..810c117 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -227,7 +227,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 || type == LAYER_ACTIVATION_LOGISTIC) + if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH || type == LAYER_ACTIVATION_SWISH || type == LAYER_ACTIVATION_LOGISTIC) return convert_layer(input, (Activation*) l); if(type == LAYER_SOFTMAX) return convert_layer(input, (Softmax*) l); @@ -424,6 +424,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 if(l->act_mode == ACTIVATION_LOGISTIC) { IPlugin *plugin = new ActivationLogisticRT(); IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); @@ -674,6 +680,12 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa assert(buf == bufCheck + serialLength); return a; } + if(name.find("ActivationSwish") == 0) { + ActivationSwishRT *a = new ActivationSwishRT(); + a->size = readBUF(buf); + assert(buf == bufCheck + serialLength); + return a; + } if(name.find("ActivationLogistic") == 0) { ActivationLogisticRT *a = new ActivationLogisticRT(); a->size = readBUF(buf); diff --git a/src/kernels/activation_swish.cu b/src/kernels/activation_swish.cu new file mode 100644 index 0000000..7018db6 --- /dev/null +++ b/src/kernels/activation_swish.cu @@ -0,0 +1,23 @@ +#include "kernels.h" + +__global__ +void activation_swish(dnnType *input, dnnType *output, int size) { + + int i = blockDim.x*blockIdx.x + threadIdx.x; + + if(i>>(srcData, dstData, size); +}