support swish activation

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);

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);
         };

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<int>(buf);
         return a;
     }
+    if(name.find("ActivationSwish") == 0) {
+        ActivationSwishRT *a = new ActivationSwishRT();
+        a->size = readBUF<int>(buf);
+        return a;
+    }
     if(name.find("ActivationCReLU") == 0) {
         ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF<float>(buf));
         a->size = readBUF<int>(buf);

src/kernels/activation_swish.cu

@@ -0,0 +1,27 @@
+#include "kernels.h"
+#include <math.h>
+
+// 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<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+}