From 6fd261f628e54d560b063b0e89ff8804f695a73f Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Mon, 22 Jun 2020 18:11:19 +0200 Subject: [PATCH] Shelfnet works on cuDNN. To test everything else Signed-off-by: Micaela Verucchi --- include/tkDNN/Layer.h | 18 ++- include/tkDNN/kernels.h | 2 +- include/tkDNN/pluginsRT/ShortcutRT.h | 12 +- src/LayerWgs.cpp | 3 +- src/NetworkRT.cpp | 4 +- src/Reshape.cpp | 5 + src/Resize.cpp | 38 ++++++ src/Shortcut.cpp | 15 ++- src/kernels/resize.cu | 41 +++---- src/kernels/shortcut.cu | 63 +++++++--- tests/shelfnet/shelfnet.cpp | 175 ++++++++++++++++++++++----- 11 files changed, 290 insertions(+), 86 deletions(-) create mode 100644 src/Resize.cpp diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 9bd8432..adbb5d1 100644 --- a/include/tkDNN/Layer.h +++ b/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; }; /** diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index 5d673c8..3d57132 100644 --- a/include/tkDNN/kernels.h +++ b/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, diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h index 3eadd3f..e4a6a89 100644 --- a/include/tkDNN/pluginsRT/ShortcutRT.h +++ b/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(outputs[0]); checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - for(int b=0; b < batchSize; ++b) - shortcutForward(srcDataBack + b*bc*bh*bw, dstData + b*c*h*w, 1, c, h, w, 1, 1, bc, bh, bw, 1, stream); + shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, mul, 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,12 +63,14 @@ 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); - + } int c, h, w; int bc, bh, bw; + bool mul; }; diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index 4afb7cc..820fb9e 100644 --- a/src/LayerWgs.cpp +++ b/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); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 6e86de1..5b75a46 100644 --- a/src/NetworkRT.cpp +++ b/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(buf); bdim.l = 1; - ShortcutRT *r = new ShortcutRT(bdim); + ShortcutRT *r = new ShortcutRT(bdim, readBUF(buf)); r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); diff --git a/src/Reshape.cpp b/src/Reshape.cpp index c1d814a..f43c4ee 100644 --- a/src/Reshape.cpp +++ b/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() { diff --git a/src/Resize.cpp b/src/Resize.cpp new file mode 100644 index 0000000..434d6b1 --- /dev/null +++ b/src/Resize.cpp @@ -0,0 +1,38 @@ +#include + +#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; +} + +}} \ No newline at end of file diff --git a/src/Shortcut.cpp b/src/Shortcut.cpp index 78a2f23..c9bddf6 100644 --- a/src/Shortcut.cpp +++ b/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 ||*/ - backLayer->output_dim.w != input_dim.w || - backLayer->output_dim.h != input_dim.h ) - FatalError("Shortcut dim missmatch"); + //FIXME + // 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 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; diff --git a/src/kernels/resize.cu b/src/kernels/resize.cu index 169b853..6059ef5 100644 --- a/src/kernels/resize.cu +++ b/src/kernels/resize.cu @@ -1,46 +1,33 @@ #include "kernels.h" #include -#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<<>>(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<<>>(o_size, srcData, i_w, i_h, i_c, o_w, o_h, o_c, n, dstData); } diff --git a/src/kernels/shortcut.cu b/src/kernels/shortcut.cu index 83539a8..c2d99bd 100644 --- a/src/kernels/shortcut.cu +++ b/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<<>>(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<<>>(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<<>>(size, minw, minh, minc, sample, batch, + w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData); + } } diff --git a/tests/shelfnet/shelfnet.cpp b/tests/shelfnet/shelfnet.cpp index 3f8834a..0fb62c1 100644 --- a/tests/shelfnet/shelfnet.cpp +++ b/tests/shelfnet/shelfnet.cpp @@ -1,5 +1,9 @@ #include +#include +#include + #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_out-conv-conv.bin", - "shelfnet/layers/conv_out-conv_out.bin"}; + "shelfnet/layers/conv_out32-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-1-conv-conv.bin", + "shelfnet/layers/decoder-up_conv_list-0-conv_atten.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-1-conv-conv.bin", + "shelfnet/layers/ladder-up_conv_list-0-conv_atten.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[] = { @@ -71,11 +89,11 @@ int main() int classes = 19; - // Network layout + // Network layout 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); @@ -105,24 +123,122 @@ int main() new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, backbone[bi++], true); 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 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<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 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/backbone-layer4-1-bn2.bin"; + const char *output_bin = "shelfnet/debug/fofmaf.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); }