Shelfnet works on cuDNN. To test everything else

Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>

include/tkDNN/Layer.h

@@ -21,6 +21,7 @@ enum layerType_t {
     LAYER_ACTIVATION_MISH,
     LAYER_FLATTEN,
     LAYER_RESHAPE,
+    LAYER_RESIZE,
     LAYER_MULADD,
     LAYER_POOLING,
     LAYER_SOFTMAX,
@@ -70,6 +71,7 @@ public:
             case LAYER_ACTIVATION_MISH:     return "ActivationMish";
             case LAYER_FLATTEN:             return "Flatten";
             case LAYER_RESHAPE:             return "Reshape";
+            case LAYER_RESIZE:              return "Resize";
             case LAYER_MULADD:              return "MulAdd";
             case LAYER_POOLING:             return "Pooling";
             case LAYER_SOFTMAX:             return "Softmax";
@@ -427,6 +429,19 @@ public:
 
 };
 
+/**
+    Resize layer
+*/
+class Resize : public Layer {
+
+public:
+    Resize(Network *net, int scale_c, int scale_h, int scale_w, bool fixed=false);
+    virtual ~Resize();
+    virtual layerType_t getLayerType() { return LAYER_RESIZE; };
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+
+};
 
 /**
     MulAdd layer
@@ -545,7 +560,7 @@ public:
 class Shortcut : public Layer {
 
 public:
-    Shortcut(Network *net, Layer *backLayer); 
+    Shortcut(Network *net, Layer *backLayer, bool mul=false); 
     virtual ~Shortcut();
     virtual layerType_t getLayerType() { return LAYER_SHORTCUT; };
 
@@ -553,6 +568,7 @@ public:
 
 public:
     Layer *backLayer;
+    bool mul = false;
 };
 
 /**

include/tkDNN/kernels.h

@@ -24,7 +24,7 @@ void softmaxForward(float *input, int n, int batch, int batch_offset,
                     int groups, int group_offset, int stride, float temp, float *output, cudaStream_t stream = cudaStream_t(0));
 
 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,
+                     int n2, int c2, int h2, int w2, int s2, bool mul,
                      cudaStream_t stream = cudaStream_t(0));
 
 void upsampleForward(dnnType *srcData, dnnType *dstData,

include/tkDNN/pluginsRT/ShortcutRT.h

@@ -4,10 +4,11 @@
 class ShortcutRT : public IPlugin {
 
 public:
-	ShortcutRT(tk::dnn::dataDim_t bdim) {
+	ShortcutRT(tk::dnn::dataDim_t bdim, bool mul) {
 		this->bc = bdim.c;
 		this->bh = bdim.h;
 		this->bw = bdim.w;
+		this->mul = mul;
 	}
 
 	~ShortcutRT(){
@@ -47,15 +48,14 @@ public:
 		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
 
 		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
-		for(int b=0; b < batchSize; ++b)
+		shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, mul, stream);
-			shortcutForward(srcDataBack + b*bc*bh*bw, dstData + b*c*h*w, 1, c, h, w, 1, 1, bc, bh, bw, 1, stream);
 
 		return 0;
 	}
 
 
 	virtual size_t getSerializationSize() override {
-		return 6*sizeof(int);
+		return 6*sizeof(int) + sizeof(bool);
 	}
 
 	virtual void serialize(void* buffer) override {
@@ -63,6 +63,7 @@ public:
 		tk::dnn::writeBUF(buf, bc);
 		tk::dnn::writeBUF(buf, bh);
 		tk::dnn::writeBUF(buf, bw);
+		tk::dnn::writeBUF(buf, mul);
 		tk::dnn::writeBUF(buf, c);
 		tk::dnn::writeBUF(buf, h);
 		tk::dnn::writeBUF(buf, w);
@@ -71,4 +72,5 @@ public:
 
 	int c, h, w;
 	int bc, bh, bw;
+	bool mul;
 };

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

@@ -512,7 +512,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; 
@@ -682,7 +682,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/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,38 @@
+#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) : Layer(net) {
+
+    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,18 @@
 
 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 ||*/
+    //FIXME
-        backLayer->output_dim.w != input_dim.w ||
+    // if( /*backLayer->output_dim.c != input_dim.c ||*/
-        backLayer->output_dim.h != input_dim.h   )
+    //     backLayer->output_dim.w != input_dim.w ||
-        FatalError("Shortcut dim missmatch");
+    //     backLayer->output_dim.h != input_dim.h   )
+    //     FatalError("Shortcut dim missmatch");
+    
 }
 
 Shortcut::~Shortcut() {
@@ -26,7 +29,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/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;
+    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
-    if(i >= i_N) return;
+    if(id >= size) return;
 
-    int out_index = i;
+    int i = id % o_w;
-    int out_w = i%o_w;
+    id /= o_w;
-    i = i/o_w;
+    int j = id % o_h;
-    int out_h = i%o_h;
+    id /= o_h;
-    i = i/o_h;
+    int k = id % o_c;
-    int out_c = i%o_c;
+    id /= o_c;
-    i = i/o_c;
+    int b = id % batch;
 
-    //copying last column/last row as padding
+    int out_index = i + o_w*(j + o_h*(k + o_c*b));
-    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);
+    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[in_index];
+    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);
-    }
 }

src/kernels/shortcut.cu

@@ -21,17 +21,35 @@ __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;
 
+    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);
@@ -42,6 +60,21 @@ void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1,
         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_mul_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, sample, batch, 
+            w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);  
+    }
+}

tests/shelfnet/shelfnet.cpp

@@ -1,5 +1,9 @@
 #include <iostream>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/imgproc/imgproc.hpp>
+
 #include "tkdnn.h"
+#include "NetworkViz.h"
 
 
 const char *output_bin1 = "shelfnet/debug/classification_headers-5.bin";
@@ -29,35 +33,49 @@ const char *backbone[] = {
     "shelfnet/layers/backbone-layer4-1-conv2.bin"};
 
 const char *conv_out[] = {
+    "shelfnet/layers/conv_out-conv-conv.bin",
+    "shelfnet/layers/conv_out-conv_out.bin",
     "shelfnet/layers/conv_out16-conv-conv.bin",
     "shelfnet/layers/conv_out16-conv_out.bin",
     "shelfnet/layers/conv_out32-conv-conv.bin",
-    "shelfnet/layers/conv_out32-conv_out.bin",
+    "shelfnet/layers/conv_out32-conv_out.bin"
-    "shelfnet/layers/conv_out-conv-conv.bin",
+    };
-    "shelfnet/layers/conv_out-conv_out.bin"};
 
 const char *decoder[] = {
     "shelfnet/layers/decoder-bottom-conv1.bin",
-    "shelfnet/layers/decoder-up_conv_list-0-conv_atten.bin",
+    "shelfnet/layers/decoder-bottom-conv12.bin",
     "shelfnet/layers/decoder-up_conv_list-0-conv-conv.bin",
-    "shelfnet/layers/decoder-up_conv_list-1-conv_atten.bin",
+    "shelfnet/layers/decoder-up_conv_list-0-conv_atten.bin",
-    "shelfnet/layers/decoder-up_conv_list-1-conv-conv.bin",
     "shelfnet/layers/decoder-up_dense_list-0-conv.bin",
-    "shelfnet/layers/decoder-up_dense_list-1-conv.bin"};
+    "shelfnet/layers/decoder-up_conv_list-1-conv-conv.bin",
+    "shelfnet/layers/decoder-up_conv_list-1-conv_atten.bin",
+    "shelfnet/layers/decoder-up_dense_list-1-conv.bin"
+    };
 
     
 const char *ladder[] = {
-    "shelfnet/layers/ladder-bottom-conv1.bin",
-    "shelfnet/layers/ladder-down_conv_list-0.bin",
-    "shelfnet/layers/ladder-down_conv_list-1.bin",
-    "shelfnet/layers/ladder-down_module_list-0-conv1.bin",
-    "shelfnet/layers/ladder-down_module_list-1-conv1.bin",
     "shelfnet/layers/ladder-inconv-conv1.bin",
-    "shelfnet/layers/ladder-up_conv_list-0-conv_atten.bin",
+    "shelfnet/layers/ladder-inconv-conv12.bin",
+    "shelfnet/layers/ladder-down_module_list-0-conv1.bin",
+    "shelfnet/layers/ladder-down_module_list-0-conv12.bin",
+    "shelfnet/layers/ladder-down_conv_list-0.bin",
+
+    "shelfnet/layers/ladder-down_module_list-1-conv1.bin",
+    "shelfnet/layers/ladder-down_module_list-1-conv12.bin",
+    "shelfnet/layers/ladder-down_conv_list-1.bin",
+
+    "shelfnet/layers/ladder-bottom-conv1.bin",
+    "shelfnet/layers/ladder-bottom-conv12.bin",
+    
+    
+    
     "shelfnet/layers/ladder-up_conv_list-0-conv-conv.bin",
-    "shelfnet/layers/ladder-up_conv_list-1-conv_atten.bin",
+    "shelfnet/layers/ladder-up_conv_list-0-conv_atten.bin",
-    "shelfnet/layers/ladder-up_conv_list-1-conv-conv.bin",
     "shelfnet/layers/ladder-up_dense_list-0-conv.bin",
+
+    
+    "shelfnet/layers/ladder-up_conv_list-1-conv-conv.bin",
+    "shelfnet/layers/ladder-up_conv_list-1-conv_atten.bin",
     "shelfnet/layers/ladder-up_dense_list-1-conv.bin"};
 
 const char *trans[] = {
@@ -75,7 +93,7 @@ int main()
     tk::dnn::dataDim_t dim(1, 3, 1024, 1024, 1);
     tk::dnn::Network net(dim);
 
-    int bi = 0;
+    int bi = 0, di = 0, li = 0, ci = 0;
     new tk::dnn::Conv2d(&net, 64, 7, 7, 2, 2, 3, 3, backbone[bi++], true);
     new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
     tk::dnn::Layer* last = new tk::dnn::Pooling (&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX);
@@ -106,23 +124,121 @@ int main()
         new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
         new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, backbone[bi++], true);
         
-        if(i != 2)
+        new tk::dnn::Shortcut(&net, last);
-        
-        {new tk::dnn::Shortcut(&net, last);
         last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
-        features.push_back(last);}
+        features.push_back(last);
     }
 
-    // for(int i=0; i<features.size(); ++i){
+    for(int i=0; i<features.size(); ++i){
-    //     new tk::dnn::Route(&net, &features[i], 1);
+        new tk::dnn::Route(&net, &features[i], 1);
-    //     int out_channel = pow(2,6+i);
+        int out_channel = pow(2,6+i);
-    //     new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, trans[i], true);
+        new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, trans[i], true);
-    //     new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+        features[i] = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+    }
+
+    //DECODER
+    
+    last = features[2];
+    std::vector<tk::dnn::Layer*> up_out;
+    //bottom
+    new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, decoder[di++], true, false, 1, true);
+    new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+    new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, decoder[di++], true, false, 1, true);
+    new tk::dnn::Shortcut(&net, last);
+    last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+    up_out.push_back(last);
+
+    for(int i=0; i<2; ++i){
+        int out_channel = pow(2,7-i);
+        //up-conv
+        std::cout<<out_channel<<std::endl;
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, decoder[di++], true);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+        
+        new tk::dnn::Pooling(&net, last->output_dim.w, last->output_dim.h, last->output_dim.w, last->output_dim.h, 0, 0, tk::dnn::POOLING_AVERAGE);
+        new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, decoder[di++], true);
+        
+        tk::dnn::Layer* act = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_SIGMOID);
+        new tk::dnn::Route(&net, &last, 1);
+        new tk::dnn::Shortcut(&net, act, true);
+
+        //interpolate
+        new tk::dnn::Resize(&net, 1,2,2);
+        new tk::dnn::Shortcut(&net, features[1-i]);
+
+        //up-dense
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, decoder[di++], true);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+        up_out.push_back(last);
+    }
+
+    //LADDER
+
+    std::vector<tk::dnn::Layer*> down_out;
+    new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+    new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+    new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+    new tk::dnn::Shortcut(&net, last);
+    new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+    
+    for(int i=0; i<2;++i){
+        int out_channel = pow(2,6+i);
+        tk::dnn::Layer* l_last = new tk::dnn::Shortcut(&net, up_out[2-i]);
+    
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+        new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+        new tk::dnn::Shortcut(&net, l_last);
+        l_last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+        down_out.push_back(l_last);
+
+        new tk::dnn::Conv2d (&net, out_channel*2, 3, 3, 2, 2, 1, 1, ladder[li++], false);
+        last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+    }
+
+    new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+    new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+    new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+    new tk::dnn::Shortcut(&net, last);
+    last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+    up_out.clear();
+    up_out.push_back(last);
+
+    for(int i=0; i<2; ++i){
+        int out_channel = pow(2,7-i);
+        //up-conv
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+        
+        new tk::dnn::Pooling(&net, last->output_dim.w, last->output_dim.h, last->output_dim.w, last->output_dim.h, 0, 0, tk::dnn::POOLING_AVERAGE);
+        new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, ladder[li++], true);
+        
+        tk::dnn::Layer* act = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_SIGMOID);
+        new tk::dnn::Route(&net, &last, 1);
+        new tk::dnn::Shortcut(&net, act, true);
+
+        //interpolate
+        new tk::dnn::Resize(&net, 1,2,2);
+        new tk::dnn::Shortcut(&net, down_out[1-i]);
+
+        // //up-dense
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+        up_out.push_back(last);
+    }
+
+
+    // for(int i=2;i>=0;--i){
+        // new tk::dnn::Route(&net, &up_out[i], 1);
+        new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, conv_out[ci++], true);
+        new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY);
+        new tk::dnn::Conv2d (&net, 19, 3, 3, 1, 1, 1, 1, conv_out[ci++], false);
+        /*up_out[i] =*/ new tk::dnn::Resize(&net, 19, net.input_dim.h, net.input_dim.w, true);
     // }
 
+    new tk::dnn::Softmax(&net);
     
-
+    const char *output_bin = "shelfnet/debug/fofmaf.bin";
-    const char *output_bin = "shelfnet/debug/backbone-layer4-1-bn2.bin";
 
     
 
@@ -210,4 +326,7 @@ int main()
     // ret_cudnn_tensorrt |= checkResult(loc->output_dim.tot(), loc->dstData, rt_out4) == 0 ? 0 : ERROR_CUDNNvsTENSORRT;
 
     // return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt;
+    
+    cv::Mat viz = vizLayer2Mat(&net, net.num_layers-1);
+    cv::imwrite("test.png", viz);
 }