merge

src/Activation.cpp

@@ -48,6 +48,10 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
     if(act_mode == ACTIVATION_LEAKY) {
         activationLEAKYForward(srcData, dstData, dim.tot());
 
+    }
+    else if(act_mode == ACTIVATION_MISH) {
+        activationMishForward(srcData, dstData, dim.tot());
+
     } else {
         dnnType alpha = dnnType(1);
         dnnType beta  = dnnType(0);

src/MobilenetDetection.cpp

@@ -172,7 +172,12 @@ bool MobilenetDetection::init(const std::string& tensor_path, const int n_classe
         colors[c] = cv::Scalar(int(255.0 * b), int(255.0 * g), int(255.0 * r));
     }
 
-    if(classes == 21){
+    if(classes == 11){ //BDD
+        const char *classes_names_[] = {
+        "person","car","truck","bus","motor","bike","rider","traffic light","traffic sign","train"};
+        classesNames = std::vector<std::string>(classes_names_, std::end(classes_names_));
+    }
+    else if(classes == 21){ //VOC
         const char *classes_names_[] = {
         "aeroplane", "bicycle", "bird", "boat", "bottle", "bus",
         "car", "cat", "chair", "cow", "diningtable", "dog", "horse", "motorbike",
@@ -180,7 +185,7 @@ bool MobilenetDetection::init(const std::string& tensor_path, const int n_classe
         classesNames = std::vector<std::string>(classes_names_, std::end(classes_names_));
 
     }
-    else if (classes == 81){
+    else if (classes == 81){ //COCO
         const char *classes_names_[] = {
         "person" , "bicycle" , "car" , "motorbike" , "aeroplane" , "bus" ,
         "train" , "truck" , "boat" , "traffic light" , "fire hydrant" , "stop sign" , 

src/NetworkRT.cpp

@@ -225,7 +225,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)
+    if(type == LAYER_ACTIVATION || type == LAYER_ACTIVATION_CRELU || type == LAYER_ACTIVATION_LEAKY || type == LAYER_ACTIVATION_MISH)
         return convert_layer(input, (Activation*) l);
     if(type == LAYER_SOFTMAX)
         return convert_layer(input, (Softmax*) l);
@@ -413,7 +413,14 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
         IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
         checkNULL(lRT);
         return lRT;
-    } else {
+    } 
+    else if(l->act_mode == ACTIVATION_MISH) {
+        IPlugin *plugin = new ActivationMishRT();
+        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+        checkNULL(lRT);
+        return lRT;
+    }
+    else {
         FatalError("this Activation mode is not yet implemented");
         return NULL;
     }
@@ -518,7 +525,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Yolo *l) {
     //std::cout<<"convert Yolo\n";
 
     //std::cout<<"New plugin YOLO\n";
-    IPlugin *plugin = new YoloRT(l->classes, l->num, l);
+    IPlugin *plugin = new YoloRT(l->classes, l->num, l, l->n_masks, l->scaleXY);
     IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
     checkNULL(lRT);
     return lRT;
@@ -637,6 +644,11 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         a->size = readBUF<int>(buf);
         return a;
     }
+    if(name.find("ActivationMish") == 0) {
+        ActivationMishRT *a = new ActivationMishRT();
+        a->size = readBUF<int>(buf);
+        return a;
+    }
     if(name.find("ActivationCReLU") == 0) {
         ActivationReLUCeiling *a = new ActivationReLUCeiling(readBUF<float>(buf));
         a->size = readBUF<int>(buf);
@@ -728,6 +740,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         r->c = readBUF<int>(buf);
         r->h = readBUF<int>(buf);
         r->w = readBUF<int>(buf);
+        r->scaleXY = readBUF<float>(buf);
         for(int i=0; i<r->n_masks; i++)
             r->mask[i] = readBUF<dnnType>(buf);
         for(int i=0; i<r->n_masks*2*r->num; i++)

src/Yolo.cpp

@@ -11,12 +11,13 @@
 
 namespace tk { namespace dnn {
 
-Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks) : 
+Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks, float scale_xy) : 
     Layer(net) {
     
     this->classes = classes;
     this->num = num;
     this->n_masks = n_masks;
+    this->scaleXY = scale_xy;
 
     // load anchors
     if(fname_weights != "") {
@@ -74,6 +75,8 @@ dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) {
         for(int n = 0; n < n_masks; ++n){
             int index = entry_index(b, n*dim.w*dim.h, 0, classes, input_dim, output_dim);
             activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
+
+            if (this->scaleXY != 1) scalAdd(dstData + index, 2 * dim.w*dim.h, this->scaleXY, -0.5*(this->scaleXY - 1), 1);
             
             index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim);
             activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h);

src/kernels/activation.mish.cu

@@ -0,0 +1,31 @@
+#include "kernels.h"
+#include <math.h>
+
+#define MISH_THRESHOLD 20
+
+__device__ float tanh_activate_kernel(float x){return (2/(1 + expf(-2*x)) - 1);}
+__device__ float softplus_kernel(float x, float threshold = 20) {
+    if (x > threshold) return x;                // too large
+    else if (x < -threshold) return expf(x);    // too small
+    return logf(expf(x) + 1);
+}
+
+// https://github.com/digantamisra98/Mish
+// https://github.com/AlexeyAB/darknet/blob/master/src/activation_kernels.cu
+__global__
+void activation_mish(dnnType *input, dnnType *output, int size) {
+    int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if (i < size) 
+        output[i] = input[i] * tanh_activate_kernel( softplus_kernel(input[i], MISH_THRESHOLD));    
+}
+
+/**
+    Mish activation function
+*/
+void activationMishForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    activation_mish<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+}

src/kernels/scaladd.cu

@@ -0,0 +1,16 @@
+#include "kernels.h"
+#include <math.h>
+
+__global__ void scal_add_kernel(dnnType* dstData, int size, float alpha, float beta, int inc)
+{
+    int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if (i < size) dstData[i*inc] = dstData[i*inc] * alpha + beta;
+}
+
+void scalAdd(dnnType* dstData, int size, float alpha, float beta, int inc, cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    scal_add_kernel<<<blocks, threads, 0, stream>>>(dstData, size, alpha, beta, inc);
+}