Merge branch 'ceccocats:master' into master

src/Activation.cpp

@@ -5,11 +5,12 @@
 
 namespace tk { namespace dnn {
 
-Activation::Activation(Network *net, int act_mode, const float ceiling) : 
+Activation::Activation(Network *net, int act_mode, const float ceiling, const float slope) : 
     Layer(net) {
 
-    this->act_mode = act_mode;
-    this->ceiling = ceiling;
+    this->act_mode  = act_mode;
+    this->ceiling   = ceiling;
+    this->slope     = slope;
     checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
 
     if(int(act_mode) < 100) {
@@ -46,8 +47,7 @@ Activation::~Activation() {
 
 dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
     if(act_mode == ACTIVATION_LEAKY) {
-        activationLEAKYForward(srcData, dstData, dim.tot());
-
+        activationLEAKYForward(srcData, dstData, dim.tot(), this->slope);
     }
     else if(act_mode == ACTIVATION_MISH) {
         activationMishForward(srcData, dstData, dim.tot());

src/LayerWgs.cpp

@@ -26,10 +26,11 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
     }
     
     readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek);
+    seek += outputs;
 
     this->batchnorm = batchnorm;
     if(batchnorm) {
-        seek += outputs;
+        
         readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek);
         seek += outputs;
         readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek);

src/NetworkRT.cpp

@@ -237,6 +237,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
         return convert_layer(input, (Flatten*) l);
     if(type == LAYER_RESHAPE)
         return convert_layer(input, (Reshape*) l);
+    if(type == LAYER_RESIZE)
+        return convert_layer(input, (Resize*) l);
     if(type == LAYER_REORG)
         return convert_layer(input, (Reorg*) l);
     if(type == LAYER_REGION)
@@ -390,13 +392,13 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
         
 #if NV_TENSORRT_MAJOR < 6                
         // plugin version
-        IPlugin *plugin = new ActivationLeakyRT();
+        IPlugin *plugin = new ActivationLeakyRT(l->slope);
         IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
         checkNULL(lRT);
         return lRT;
 #else 
         IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU);
-        lRT->setAlpha(0.1);
+        lRT->setAlpha(l->slope);
         checkNULL(lRT);
         return lRT;
 #endif
@@ -479,13 +481,23 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Flatten *l) {
 ILayer* NetworkRT::convert_layer(ITensor *input, Reshape *l) {
     // std::cout<<"convert Reshape\n";
 
-    l->output_dim.print();
     IPlugin *plugin = new ReshapeRT(l->output_dim);
     IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
     checkNULL(lRT);
     return lRT;
 }
 
+ILayer* NetworkRT::convert_layer(ITensor *input, Resize *l) {
+    // std::cout<<"convert Resize\n";
+
+    IResizeLayer *lRT = networkRT->addResize(*input); //default is kNEAREST
+    checkNULL(lRT);
+    Dims d{};
+    lRT->setResizeMode(ResizeMode(l->mode));
+    lRT->setOutputDimensions(DimsCHW{l->output_dim.c, l->output_dim.h, l->output_dim.w});
+    return lRT;
+}
+
 ILayer* NetworkRT::convert_layer(ITensor *input, Reorg *l) {
     //std::cout<<"convert Reorg\n";
 
@@ -513,7 +525,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
     
     ITensor *back_tens = tensors[l->backLayer];
 
-    if(l->backLayer->output_dim.c == l->output_dim.c)
+    if(l->backLayer->output_dim.c == l->output_dim.c && !l->mul) 
     {
         IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM);
         checkNULL(lRT);
@@ -522,7 +534,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
     else
     {
         // plugin version
-        IPlugin *plugin = new ShortcutRT(l->backLayer->output_dim);
+        IPlugin *plugin = new ShortcutRT(l->backLayer->output_dim, l->mul);
         ITensor **inputs = new ITensor*[2];
         inputs[0] = input;
         inputs[1] = back_tens; 
@@ -651,7 +663,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
     //std::cout<<name<<std::endl;
 
     if(name.find("ActivationLeaky") == 0) {
-        ActivationLeakyRT *a = new ActivationLeakyRT();
+        ActivationLeakyRT *a = new ActivationLeakyRT(readBUF<float>(buf));
         a->size = readBUF<int>(buf);
         assert(buf == bufCheck + serialLength);
         return a;
@@ -710,7 +722,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
         bdim.w = readBUF<int>(buf);
         bdim.l = 1;
 
-        ShortcutRT *r = new ShortcutRT(bdim);
+        ShortcutRT *r = new ShortcutRT(bdim, readBUF<bool>(buf));
         r->c = readBUF<int>(buf);
         r->h = readBUF<int>(buf);
         r->w = readBUF<int>(buf);

src/NetworkViz.cpp

@@ -6,23 +6,389 @@
 
 namespace tk { namespace dnn {
 
-cv::Mat vizFloat2colorMap(cv::Mat map) {
+cv::Mat mapillary_15_map(cv::Mat adjMap){
+
+    // cv::imshow("test", adjMap);
+    // cv::waitKey(0);
+    cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1);
+
+    //animal
+    M3.at<uchar>(0)=165;
+    M2.at<uchar>(0)=42;
+    M1.at<uchar>(0)=45;
+
+    //curb
+    M3.at<uchar>(1)=196;
+    M2.at<uchar>(1)=196;
+    M1.at<uchar>(1)=196;
+
+    //barrier
+    M3.at<uchar>(2)=90;
+    M2.at<uchar>(2)=120;
+    M1.at<uchar>(2)=150;
+
+    //road
+    M3.at<uchar>(3)=128;
+    M2.at<uchar>(3)=64;
+    M1.at<uchar>(3)=128;
+
+    //building
+    M3.at<uchar>(4)=70;
+    M2.at<uchar>(4)=70;
+    M1.at<uchar>(4)=70;
+
+    //person
+    M3.at<uchar>(5)=220;
+    M2.at<uchar>(5)=20;
+    M1.at<uchar>(5)=60;
+
+    //roadmark
+    M3.at<uchar>(6)=255;
+    M2.at<uchar>(6)=255;
+    M1.at<uchar>(6)=255;
+
+    //nature
+    M3.at<uchar>(7)=107;
+    M2.at<uchar>(7)=142;
+    M1.at<uchar>(7)=35;
+
+    //sky
+    M3.at<uchar>(8)=70;
+    M2.at<uchar>(8)=130;
+    M1.at<uchar>(8)=180;
+
+    //billboard
+    M3.at<uchar>(9)=220;
+    M2.at<uchar>(9)=220;
+    M1.at<uchar>(9)=220;
+
+    //pole
+    M3.at<uchar>(10)=153;
+    M2.at<uchar>(10)=153;
+    M1.at<uchar>(10)=153;
+
+    //traffic sign
+    M3.at<uchar>(11)=128;
+    M2.at<uchar>(11)=128;
+    M1.at<uchar>(11)=128;
+
+    //bike
+    M3.at<uchar>(12)=119;
+    M2.at<uchar>(12)=11;
+    M1.at<uchar>(12)=32;
+
+    //vehicle
+    M3.at<uchar>(13)=0;
+    M2.at<uchar>(13)=0;
+    M1.at<uchar>(13)=142;
+
+    //void
+    for(int i=14;i<256;i++)
+    {
+        M1.at<uchar>(i)=0;
+        M2.at<uchar>(i)=0;
+        M3.at<uchar>(i)=0;
+    }
+
+    cv::Mat r1,r2,r3;
+
+    cv::LUT(adjMap,M1,r1);
+    cv::LUT(adjMap,M2,r2);
+    cv::LUT(adjMap,M3,r3);
+
+    std::vector<cv::Mat> planes;
+    planes.push_back(r1);
+    planes.push_back(r2);
+    planes.push_back(r3);
+
+    cv::Mat dst;
+    cv::merge(planes,dst);
+    return dst;
+
+
+}
+
+cv::Mat berkeley_20_map(cv::Mat adjMap){
+
+    cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1);
+
+    //road
+    M3.at<uchar>(0)=128;
+    M2.at<uchar>(0)=64;
+    M1.at<uchar>(0)=128;
+
+    //sidewalk
+    M3.at<uchar>(1)=244;
+    M2.at<uchar>(1)=35;
+    M1.at<uchar>(1)=232;
+
+    //building
+    M3.at<uchar>(2)=70;
+    M2.at<uchar>(2)=70;
+    M1.at<uchar>(2)=70;
+
+    //wall
+    M3.at<uchar>(3)=102;
+    M2.at<uchar>(3)=102;
+    M1.at<uchar>(3)=156;
+
+    //fence
+    M3.at<uchar>(4)=90;
+    M2.at<uchar>(4)=120;
+    M1.at<uchar>(4)=150;
+
+    //pole
+    M3.at<uchar>(5)=153;
+    M2.at<uchar>(5)=153;
+    M1.at<uchar>(5)=153;
+
+    //traffic light
+    M3.at<uchar>(6)=250;
+    M2.at<uchar>(6)=170;
+    M1.at<uchar>(6)=30;
+
+    //traffic sign
+    M3.at<uchar>(7)=128;
+    M2.at<uchar>(7)=128;
+    M1.at<uchar>(7)=128;
+
+    //nature
+    M3.at<uchar>(8)=107;
+    M2.at<uchar>(8)=142;
+    M1.at<uchar>(8)=35;
+
+    //ground
+    M3.at<uchar>(9)=0;
+    M2.at<uchar>(9)=192;
+    M1.at<uchar>(9)=0;
+
+    //sky
+    M3.at<uchar>(10)=70;
+    M2.at<uchar>(10)=130;
+    M1.at<uchar>(10)=180;
+
+    //person
+    M3.at<uchar>(11)=220;
+    M2.at<uchar>(11)=20;
+    M1.at<uchar>(11)=60;
+
+    //rider
+    M3.at<uchar>(12)=255;
+    M2.at<uchar>(12)=0;
+    M1.at<uchar>(12)=100;
+
+    //car
+    M3.at<uchar>(13)=0;
+    M2.at<uchar>(13)=0;
+    M1.at<uchar>(13)=142;
+
+    //truck
+    M3.at<uchar>(14)=0;
+    M2.at<uchar>(14)=0;
+    M1.at<uchar>(14)=70;
+
+    //bus
+    M3.at<uchar>(15)=0;
+    M2.at<uchar>(15)=60;
+    M1.at<uchar>(15)=100;
+
+    //train
+    M3.at<uchar>(16)=0;
+    M2.at<uchar>(16)=0;
+    M1.at<uchar>(16)=192;
+
+    //motorbike
+    M3.at<uchar>(17)=0;
+    M2.at<uchar>(17)=0;
+    M1.at<uchar>(17)=230;
+
+    //bike
+    M3.at<uchar>(18)=119;
+    M2.at<uchar>(18)=11;
+    M1.at<uchar>(18)=32;
+
+    //void
+    for(int i=19;i<256;i++)
+    {
+        M1.at<uchar>(i)=0;
+        M2.at<uchar>(i)=0;
+        M3.at<uchar>(i)=0;
+    }
+
+    cv::Mat r1,r2,r3;
+
+    cv::LUT(adjMap,M1,r1);
+    cv::LUT(adjMap,M2,r2);
+    cv::LUT(adjMap,M3,r3);
+
+    std::vector<cv::Mat> planes;
+    planes.push_back(r1);
+    planes.push_back(r2);
+    planes.push_back(r3);
+
+    cv::Mat dst;
+    cv::merge(planes,dst);
+    return dst;
+
+}
+
+cv::Mat cityscapes_19_map(cv::Mat adjMap){
+
+    cv::Mat M1(1, 256, CV_8UC1), M2(1, 256, CV_8UC1), M3(1, 256, CV_8UC1);
+
+    //road
+    M3.at<uchar>(0)=128;
+    M2.at<uchar>(0)=64;
+    M1.at<uchar>(0)=128;
+
+    //sidewalk
+    M3.at<uchar>(1)=244;
+    M2.at<uchar>(1)=35;
+    M1.at<uchar>(1)=232;
+
+    //building
+    M3.at<uchar>(2)=70;
+    M2.at<uchar>(2)=70;
+    M1.at<uchar>(2)=70;
+
+    //wall
+    M3.at<uchar>(3)=102;
+    M2.at<uchar>(3)=102;
+    M1.at<uchar>(3)=156;
+
+    //fence
+    M3.at<uchar>(4)=190;
+    M2.at<uchar>(4)=153;
+    M1.at<uchar>(4)=153;
+
+    //pole
+    M3.at<uchar>(5)=153;
+    M2.at<uchar>(5)=153;
+    M1.at<uchar>(5)=153;
+
+    //traffic light
+    M3.at<uchar>(6)=250;
+    M2.at<uchar>(6)=170;
+    M1.at<uchar>(6)=30;
+
+    //traffic sign
+    M3.at<uchar>(7)=220;
+    M2.at<uchar>(7)=220;
+    M1.at<uchar>(7)=0;
+
+    //vegetation
+    M3.at<uchar>(8)=107;
+    M2.at<uchar>(8)=142;
+    M1.at<uchar>(8)=35;
+
+    //terrain
+    M3.at<uchar>(9)=152;
+    M2.at<uchar>(9)=251;
+    M1.at<uchar>(9)=152;
+
+    //sky
+    M3.at<uchar>(10)=70;
+    M2.at<uchar>(10)=130;
+    M1.at<uchar>(10)=180;
+
+    //person
+    M3.at<uchar>(11)=220;
+    M2.at<uchar>(11)=20;
+    M1.at<uchar>(11)=60;
+
+    //rider
+    M3.at<uchar>(12)=255;
+    M2.at<uchar>(12)=0;
+    M1.at<uchar>(12)=0;
+
+    //car
+    M3.at<uchar>(13)=0;
+    M2.at<uchar>(13)=0;
+    M1.at<uchar>(13)=142;
+
+    //truck
+    M3.at<uchar>(14)=0;
+    M2.at<uchar>(14)=0;
+    M1.at<uchar>(14)=70;
+
+    //bus
+    M3.at<uchar>(15)=0;
+    M2.at<uchar>(15)=60;
+    M1.at<uchar>(15)=100;
+
+    //train
+    M3.at<uchar>(16)=0;
+    M2.at<uchar>(16)=80;
+    M1.at<uchar>(16)=100;
+
+    //motorcycle
+    M3.at<uchar>(17)=0;
+    M2.at<uchar>(17)=0;
+    M1.at<uchar>(17)=230;
+
+    //bicycle
+    M3.at<uchar>(18)=119;
+    M2.at<uchar>(18)=11;
+    M1.at<uchar>(18)=32;
+
+    //void
+    for(int i=19;i<256;i++)
+    {
+        M1.at<uchar>(i)=0;
+        M2.at<uchar>(i)=0;
+        M3.at<uchar>(i)=0;
+    }
+
+    cv::Mat r1,r2,r3;
+
+    cv::LUT(adjMap,M1,r1);
+    cv::LUT(adjMap,M2,r2);
+    cv::LUT(adjMap,M3,r3);
+
+    std::vector<cv::Mat> planes;
+    planes.push_back(r1);
+    planes.push_back(r2);
+    planes.push_back(r3);
+
+    cv::Mat dst;
+    cv::merge(planes,dst);
+    return dst;
+
+}
+
+
+cv::Mat vizFloat2colorMap(cv::Mat map,double min, double max, int classes) {
+
+    if(min == 0 && max == 0)
+        cv::minMaxIdx(map, &min, &max);
 
-    double min;
-    double max;
-    cv::minMaxIdx(map, &min, &max);
     cv::Mat adjMap;
-    // expand your range to 0..255. Similar to histEq();
-    map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min); 
-    //return adjMap;
-
-    
     cv::Mat falseColorsMap;
-    applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_HOT);
+    
+    switch (classes)
+    {
+    case 15:
+        map.convertTo(adjMap,CV_8UC1);
+        falseColorsMap = mapillary_15_map(adjMap);
+        break;
+    case 20:
+        map.convertTo(adjMap,CV_8UC1);
+        falseColorsMap = berkeley_20_map(adjMap);
+        break;
+    case 19:
+        map.convertTo(adjMap,CV_8UC1);
+        falseColorsMap = cityscapes_19_map(adjMap);
+        break;
+    
+    default:
+        // expand your range to 0..255. Similar to histEq();
+        map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min); 
+        applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_JET);
+    }
     return falseColorsMap;
 }
 
-cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) {
+cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int img_h, int img_w, double min, double max, int classes) {
     dnnType *data = nullptr;
 
     // copy to CPU
@@ -38,14 +404,13 @@ cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) {
     cv::Mat grid = cv::Mat(gridSize, CV_8UC3, cv::Scalar(0));    
 
     for(int i=0; i<dim.c;i++) {
-        cv::Mat raw = vizFloat2colorMap(cv::Mat(cv::Size(dim.w, dim.h),CV_32FC1, data + dim.w*dim.h*i));
+        cv::Mat raw = vizFloat2colorMap(cv::Mat(cv::Size(dim.w, dim.h),CV_32FC1, data + dim.w*dim.h*i), min, max, classes);
         int r = i / gridDim;
         int c = i - r * gridDim;
         raw.copyTo(grid.rowRange(r*dim.h, r*dim.h + dim.h).colRange(c*dim.w, c*dim.w + dim.w));
     }
 
-    float ar = float(dim.w)/dim.h;
-    cv::Size vdim(ar*imgdim, imgdim);
+    cv::Size vdim(img_w, img_h);
     cv::Mat viz;
     cv::resize(grid, viz, vdim, 0, 0, 0);
     
@@ -59,7 +424,7 @@ cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) {
 cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim) {
     if(layer >= net->num_layers)
         FatalError("Could not viz layer\n");
-    return vizData2Mat(net->layers[layer]->dstData, net->layers[layer]->output_dim, imgdim);
+    return vizData2Mat(net->layers[layer]->dstData, net->layers[layer]->output_dim, imgdim, imgdim);
 
     //cv::imwrite("viz/layer" + std::to_string(layer) + ".png", viz);
     //cv::imshow("layer", viz);

src/Reshape.cpp

@@ -15,6 +15,11 @@ Reshape::Reshape(Network *net, dataDim_t new_dim) : Layer(net) {
     output_dim.w = new_dim.w;
     output_dim.l = new_dim.l;
 
+    output_dim = new_dim;
+
+    if(input_dim.tot() != output_dim.tot())
+        FatalError("Reshape dimension mismatch");
+
 }
 
 Reshape::~Reshape() {

src/Resize.cpp

@@ -0,0 +1,39 @@
+#include <iostream>
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tk { namespace dnn {
+
+Resize::Resize(Network *net, int scale_c, int scale_h, int scale_w, bool fixed, ResizeMode_t mode) : Layer(net) {
+
+    this->mode = mode;
+    if(fixed){
+        output_dim.c = scale_c;
+        output_dim.h = scale_h;
+        output_dim.w = scale_w;
+    }
+    else{
+        output_dim.c *= scale_c;
+        output_dim.h *= scale_h;
+        output_dim.w *= scale_w;
+    }
+
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
+}
+
+Resize::~Resize() {
+
+    checkCuda( cudaFree(dstData) );
+}
+
+dnnType* Resize::infer(dataDim_t &dim, dnnType* srcData) {
+
+    resizeForward(srcData, dstData, dim.n, dim.c, dim.h, dim.w,
+                   output_dim.c, output_dim.h, output_dim.w);
+    dim = output_dim;
+
+    return dstData;
+}
+
+}}

src/Shortcut.cpp

@@ -5,15 +5,16 @@
 
 namespace tk { namespace dnn {
 
-Shortcut::Shortcut(Network *net, Layer *backLayer) : Layer(net) {
+Shortcut::Shortcut(Network *net, Layer *backLayer, bool mul) : Layer(net) {
 
     this->backLayer = backLayer;
+    this->mul = mul;
     checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
 
-    if( /*backLayer->output_dim.c != input_dim.c ||*/
-        backLayer->output_dim.w != input_dim.w ||
-        backLayer->output_dim.h != input_dim.h   )
-        FatalError("Shortcut dim mismatch");
+    if( ( backLayer->output_dim.c != input_dim.c && mul ) ||
+        (( backLayer->output_dim.w != input_dim.w ||  backLayer->output_dim.h != input_dim.h ) && !mul ) )  
+        FatalError("Shortcut dim missmatch");
+    
 }
 
 Shortcut::~Shortcut() {
@@ -26,7 +27,7 @@ dnnType* Shortcut::infer(dataDim_t &dim, dnnType* srcData) {
     dataDim_t bdim = this->backLayer->output_dim;
 
     checkCuda(cudaMemcpy(dstData, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
-    shortcutForward(this->backLayer->dstData, dstData, dim.n, dim.c, dim.h, dim.w, 1, bdim.n, bdim.c, bdim.h, bdim.w, 1);
+    shortcutForward(this->backLayer->dstData, dstData, dim.n, dim.c, dim.h, dim.w, 1, bdim.n, bdim.c, bdim.h, bdim.w, 1, mul);
 
     //update data dimensions    
     dim = output_dim;

src/kernels/activation_leaky.cu

@@ -1,7 +1,7 @@
 #include "kernels.h"
 
 __global__
-void activation_leaky(dnnType *input, dnnType *output, int size) {
+void activation_leaky(dnnType *input, dnnType *output, int size, float slope) {
 
     int i = blockDim.x*blockIdx.x + threadIdx.x;
 
@@ -9,7 +9,7 @@ void activation_leaky(dnnType *input, dnnType *output, int size) {
         if (input[i]>0)
             output[i] = input[i];
         else
-            output[i] = 0.1f*input[i];
+            output[i] = slope*input[i];
     }
  }
 
@@ -17,12 +17,12 @@ void activation_leaky(dnnType *input, dnnType *output, int size) {
 /**
     ELU activation function
 */
-void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream)
+void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, float slope, cudaStream_t stream)
 {
     int blocks = (size+255)/256;
     int threads = 256;
     
-    activation_leaky<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+    activation_leaky<<<blocks, threads, 0, stream>>>(srcData, dstData, size, slope);
 }
 
 

src/kernels/postprocessing.cu

@@ -34,6 +34,25 @@ void sortAndTopKonDevice(dnnType *src_begin, int *idsrc, float *topk_scores, int
     sortAndTopK_kernel<<<blocks, threads, 0>>>(src_begin, idsrc, topk_scores, topk_inds, topk_ys, topk_xs, size, K);   
 }
 
+__global__ 
+void maxElem_kernel(float *src_begin, float *dst_begin, const int n_classes, const int size){
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+    if (i > size)
+        return;
+
+    thrust::device_ptr<float> dPbeg ( &src_begin[i*n_classes] ) ;
+    thrust::device_ptr<float> dPend = dPbeg + n_classes;
+    thrust::device_ptr<float> result = thrust::max_element(thrust::device,dPbeg, dPend);
+
+    dst_begin[i] = result - dPbeg;
+}
+
+void maxElem(dnnType *src_begin, dnnType *dst_begin, const int c, const int h, const int w){
+    int blocks = (h*w)/32+1;
+    int threads = 32;
+    maxElem_kernel<<<blocks, threads, 0>>>(src_begin, dst_begin, c, h*w);   
+}
+
 void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, const int wh, int *clses, int *xs, int *ys){    
     thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(wh), clses, thrust::divides<int>());
     thrust::transform(thrust::device, ids_begin, ids_end, thrust::make_constant_iterator(wh), ids_begin, thrust::modulus<int>());

src/kernels/resize.cu

@@ -1,46 +1,33 @@
 #include "kernels.h"
 #include <stdio.h>
-#define MIN(a,b) (((a)<(b))?(a):(b))
-#define MAX(a,b) (((a)>(b))?(a):(b))
 
-__global__ void resize_kernel(  int i_N,float *x, int i_w, int i_h, int i_c,  
+__global__ void resize_kernel(  int size,float *x, int i_w, int i_h, int i_c,  
                                 int o_w, int o_h, int o_c, int batch, float *out)
 {
-    int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
-    if(i >= i_N) return;
+    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if(id >= size) return;
 
-    int out_index = i;
-    int out_w = i%o_w;
-    i = i/o_w;
-    int out_h = i%o_h;
-    i = i/o_h;
-    int out_c = i%o_c;
-    i = i/o_c;
+    int i = id % o_w;
+    id /= o_w;
+    int j = id % o_h;
+    id /= o_h;
+    int k = id % o_c;
+    id /= o_c;
+    int b = id % batch;
 
-    //copying last column/last row as padding
-    int in_index = ((i*i_c + MIN(out_c,i_c-1))*i_h +  MIN(out_h,i_h-1))*i_w + MIN(out_w, i_w-1);
-    out[out_index] = x[in_index];
+    int out_index = i + o_w*(j + o_h*(k + o_c*b));
+    int add_index = i/(o_w/i_w) + i_w*(j/(o_h/i_h) + i_h*(k + i_c*b));
+    out[out_index] = x[add_index];
 }
 
 
 void resizeForward(  dnnType* srcData, dnnType* dstData, int n, int i_c, int i_h, int i_w,
      int o_c, int o_h, int o_w, cudaStream_t stream )
 {
-    int i_size = n*i_c*i_h*i_w;
     int o_size = n*o_c*o_h*o_w;
 
     int blocks = (o_size+255)/256;
     int threads = 256;
 
-    if(i_c == o_c && i_h == o_h && i_w == o_w )
-    {
-        checkCuda(cudaMemcpy(dstData, srcData, i_size*sizeof(dnnType), cudaMemcpyDeviceToDevice));
-    }
-    else
-    {
-        checkCuda(cudaMemset(dstData, 0, o_size*sizeof(dnnType)));
-        resize_kernel<<<blocks, threads, 0, stream>>>(o_size, srcData, i_w, i_h, i_c, o_w, o_h, o_c, n, dstData);
-        // printDeviceVector(i_size, srcData);
-        // printDeviceVector(o_size, dstData);
-    }
+    resize_kernel<<<blocks, threads, 0, stream>>>(o_size, srcData, i_w, i_h, i_c, o_w, o_h, o_c, n, dstData);
 }

src/kernels/shortcut.cu

@@ -21,27 +21,60 @@ __global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stri
     //out[out_index] += add[add_index];
 }
 
+__global__ void shortcut_mul_kernel(int size, int minw, int minh, int minc, int sample, int batch, 
+                                int w1, int h1, int c1, dnnType *mul, 
+                                int w2, int h2, int c2, float s1, float s2, dnnType *out)
+{
+    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if (id >= size) return;
+    int i = id % minw;
+    id /= minw;
+    int j = id % minh;
+    id /= minh;
+    int k = id % minc;
+    id /= minc;
+    int b = id % batch;
+
+    int out_index = i*sample + w1*(j*sample + h1*(k + c1*b));
+    out[out_index] = out[out_index] * mul[k + c2*b];
+}
+
 void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int w1, int s1,
                                                          int n2, int c2, int h2, int w2, int s2, 
-                     cudaStream_t stream)
+                                                         bool mul, cudaStream_t stream)
 {
     assert(n1 == n2);
     int batch = n1;
 
-    int minw = (w1 < w2) ? w1 : w2;
-    int minh = (h1 < h2) ? h1 : h2;
-    int minc = (c1 < c2) ? c1 : c2;
+    if(!mul){
+        int minw = (w1 < w2) ? w1 : w2;
+        int minh = (h1 < h2) ? h1 : h2;
+        int minc = (c1 < c2) ? c1 : c2;
+        int stride = w1/w2;
+        int sample = w2/w1;
+        assert(stride == h1/h2);
+        assert(sample == h2/h1);
+        if(stride < 1) stride = 1;
+        if(sample < 1) sample = 1;
 
-    int stride = w1/w2;
-    int sample = w2/w1;
-    assert(stride == h1/h2);
-    assert(sample == h2/h1);
-    if(stride < 1) stride = 1;
-    if(sample < 1) sample = 1;
+        int size = batch * minw * minh * minc;
+        int blocks = (size+255)/256;
+        int threads = 256;
+        
+        shortcut_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, stride, sample, batch, 
+            w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);  
+    } 
+    else{
+        int minw = w1;
+        int minh = h1;
+        int minc = c1;
+        int sample = 1;
 
-    int size = batch * minw * minh * minc;
-    int blocks = (size+255)/256;
-    int threads = 256;
-    shortcut_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, stride, sample, batch, 
-        w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);
+        int size = batch * minw * minh * minc;
+        int blocks = (size+255)/256;
+        int threads = 256;
+
+        shortcut_mul_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, sample, batch, 
+            w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);  
+    }
 }

src/utils.cpp

@@ -111,6 +111,7 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device, int
     }
     int diffs = 0;
     for(int i=0; i<size; i++) {
+        // data_h[i] = data_h[i]*1e-2;
         if(data_h[i] != data_h[i] || correct_h[i] != correct_h[i] || //nan control
            fabs(data_h[i] - correct_h[i]) > eps) {
             diffs += 1;