From 94e558003d1b87b751060d606d3137456a1eb7b1 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Tue, 23 Jun 2020 12:50:24 +0200 Subject: [PATCH] Shelfnet works on tensorRT (shortcut need to be fixed) Signed-off-by: Micaela Verucchi --- include/tkDNN/Layer.h | 3 +- include/tkDNN/NetworkRT.h | 1 + include/tkDNN/kernels.h | 2 +- include/tkDNN/pluginsRT/ActivationLeakyRT.h | 11 +- src/Activation.cpp | 10 +- src/NetworkRT.cpp | 21 +++- src/Shortcut.cpp | 8 +- src/kernels/activation_leaky.cu | 8 +- tests/shelfnet/shelfnet.cpp | 111 +++++++------------- 9 files changed, 75 insertions(+), 100 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index adbb5d1..7ea5e6c 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -225,8 +225,9 @@ class Activation : public Layer { public: int act_mode; float ceiling; + float slope; - Activation(Network *net, int act_mode, const float ceiling=0.0); + Activation(Network *net, int act_mode, const float ceiling=0.0, const float slope=0.1); virtual ~Activation(); virtual layerType_t getLayerType() { if(act_mode == CUDNN_ACTIVATION_CLIPPED_RELU) diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index ee1f728..3bd53d8 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -105,6 +105,7 @@ public: nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Route *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Flatten *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Reshape *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Resize *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Reorg *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Region *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l); diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index 3d57132..d809129 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -4,7 +4,7 @@ #include "utils.h" void activationELUForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); -void activationLEAKYForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); +void activationLEAKYForward(dnnType *srcData, dnnType *dstData, int size, float slope, cudaStream_t stream = cudaStream_t(0)); void activationReLUCeilingForward(dnnType *srcData, dnnType *dstData, int size, const float ceiling, cudaStream_t stream = cudaStream_t(0)); void activationLOGISTICForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); void activationSIGMOIDForward(dnnType *srcData, dnnType *dstData, int size, cudaStream_t stream = cudaStream_t(0)); diff --git a/include/tkDNN/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h index d3f66fb..30bed7e 100644 --- a/include/tkDNN/pluginsRT/ActivationLeakyRT.h +++ b/include/tkDNN/pluginsRT/ActivationLeakyRT.h @@ -4,9 +4,8 @@ class ActivationLeakyRT : public IPlugin { public: - ActivationLeakyRT() { - - + ActivationLeakyRT(float s) { + slope = s; } ~ActivationLeakyRT(){ @@ -42,19 +41,21 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { activationLEAKYForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), batchSize*size, stream); + reinterpret_cast(outputs[0]), batchSize*size, slope, stream); return 0; } virtual size_t getSerializationSize() override { - return 1*sizeof(int); + return 1*sizeof(int) + 1*sizeof(float); } virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer); + tk::dnn::writeBUF(buf, slope); tk::dnn::writeBUF(buf, size); } int size; + float slope; }; diff --git a/src/Activation.cpp b/src/Activation.cpp index 28c7624..c97a717 100644 --- a/src/Activation.cpp +++ b/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()); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 5b75a46..71662d3 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -236,6 +236,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) @@ -389,13 +391,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 @@ -469,13 +471,22 @@ 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->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"; @@ -503,7 +514,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(false) //l->backLayer->output_dim.c == l->output_dim.c && !l->mul) FIXME { IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM); checkNULL(lRT); @@ -641,7 +652,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa //std::cout<(buf)); a->size = readBUF(buf); return a; } diff --git a/src/Shortcut.cpp b/src/Shortcut.cpp index c9bddf6..b1053c8 100644 --- a/src/Shortcut.cpp +++ b/src/Shortcut.cpp @@ -11,11 +11,9 @@ Shortcut::Shortcut(Network *net, Layer *backLayer, bool mul) : Layer(net) { this->mul = mul; checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); - //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"); + 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"); } diff --git a/src/kernels/activation_leaky.cu b/src/kernels/activation_leaky.cu index 47dc18f..2ecb3af 100644 --- a/src/kernels/activation_leaky.cu +++ b/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.01f*input[i]; //FIME!! + 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<<>>(srcData, dstData, size); + activation_leaky<<>>(srcData, dstData, size, slope); } diff --git a/tests/shelfnet/shelfnet.cpp b/tests/shelfnet/shelfnet.cpp index 0fb62c1..ab28d90 100644 --- a/tests/shelfnet/shelfnet.cpp +++ b/tests/shelfnet/shelfnet.cpp @@ -6,8 +6,6 @@ #include "NetworkViz.h" -const char *output_bin1 = "shelfnet/debug/classification_headers-5.bin"; -const char *output_bin2 = "shelfnet/debug/regression_headers-5.bin"; const char *input_bin = "shelfnet/debug/input.bin"; const char *backbone[] = { @@ -95,14 +93,14 @@ int main() 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); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); tk::dnn::Layer* last = new tk::dnn::Pooling (&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX); for(int i=0; i<2; ++i){ new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true); - new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY); + new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true); new tk::dnn::Shortcut(&net, last); last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU); @@ -113,7 +111,7 @@ int main() int out_channel = pow(2,7+i); std::cout< 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::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); 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); @@ -153,7 +151,7 @@ int main() //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); @@ -168,7 +166,7 @@ int main() //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); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); up_out.push_back(last); } @@ -176,7 +174,7 @@ int main() 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::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); 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); @@ -186,7 +184,7 @@ int main() 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::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); 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); @@ -197,7 +195,7 @@ int main() } 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::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); 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); @@ -208,7 +206,7 @@ int main() 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); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); 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); @@ -223,7 +221,7 @@ int main() // //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); + last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); up_out.push_back(last); } @@ -231,29 +229,26 @@ int main() // 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::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01); 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); + // /*up_out[i] =*/ new tk::dnn::Resize(&net, 19, net.input_dim.h, net.input_dim.w, true); // } - new tk::dnn::Softmax(&net); + // new tk::dnn::Softmax(&net); - const char *output_bin = "shelfnet/debug/fofmaf.bin"; - + const char *output_bin = "shelfnet/debug/conv_out-conv_out.bin"; - // Load input dnnType *data; dnnType *input_h; readBinaryFile(input_bin, dim.tot(), &input_h, &data); std::cout<<"Input:"<output_dim; - // dnnType *cudnn_out2 = loc5[0]->dstData; - // tk::dnn::dataDim_t out_dim2 = loc5[0]->output_dim; + tk::dnn::dataDim_t dim2 = dim; + printCenteredTitle(" TENSORRT inference ", '=', 30); + { + dim2.print(); + TKDNN_TSTART + netRT.infer(dim2, data); + TKDNN_TSTOP + dim2.print(); + } - // tk::dnn::dataDim_t dim2 = dim; - // printCenteredTitle(" TENSORRT inference ", '=', 30); - // { - // dim2.print(); - // TKDNN_TSTART - // netRT.infer(dim2, data); - // TKDNN_TSTOP - // dim2.print(); - // } + dnnType *rt_out1 = (dnnType *)netRT.buffersRT[1]; - // dnnType *rt_out1 = (dnnType *)netRT.buffersRT[1]; - // dnnType *rt_out2 = (dnnType *)netRT.buffersRT[2]; - // dnnType *rt_out3 = (dnnType *)netRT.buffersRT[3]; - // dnnType *rt_out4 = (dnnType *)netRT.buffersRT[4]; - - printCenteredTitle(std::string(" RESNET CHECK RESULTS ").c_str(), '=', 30); + printCenteredTitle(std::string(" CHECK RESULTS ").c_str(), '=', 30); dnnType *out1, *out1_h; int odim1 = dim1.tot(); readBinaryFile(output_bin, odim1, &out1_h, &out1); - printDeviceVector(64, out1); - - // dnnType *out2, *out2_h; - // int odim2 = out_dim2.tot(); - // readBinaryFile(output_bin2, odim2, &out2_h, &out2); - // int ret_cudnn = 0, ret_tensorrt = 0, ret_cudnn_tensorrt = 0; - + int ret_cudnn = 0, ret_tensorrt = 0, ret_cudnn_tensorrt = 0; std::cout << "CUDNN vs correct" << std::endl; - checkResult(odim1, cudnn_out, out1, true, 20) == 0 ? 0 : ERROR_CUDNN; + ret_cudnn |= checkResult(odim1, cudnn_out, out1, true, 20) == 0 ? 0 : ERROR_CUDNN; - // std::cout << "TRT vs correct" << std::endl; - // checkResult(odim1, rt_out1, out1) == 0 ? 0 : ERROR_TENSORRT; - // ret_tensorrt |= checkResult(odim2, rt_out2, out2) == 0 ? 0 : ERROR_TENSORRT; + std::cout << "TRT vs correct" << std::endl; + ret_tensorrt |=checkResult(odim1, rt_out1, out1) == 0 ? 0 : ERROR_TENSORRT; - // std::cout << "CUDNN vs TRT " << std::endl; - // ret_cudnn_tensorrt |= checkResult(odim1, cudnn_out1, rt_out1) == 0 ? 0 : ERROR_CUDNNvsTENSORRT; - // ret_cudnn_tensorrt |= checkResult(odim2, cudnn_out2, rt_out2) == 0 ? 0 : ERROR_CUDNNvsTENSORRT; - - // std::cout << "---------------------------------------------------" << std::endl; - // std::cout << "Confidence CUDNN" << std::endl; - // printDeviceVector(64, conf->dstData, true); - // std::cout << "Locations CUDNN" << std::endl; - // printDeviceVector(64, loc->dstData, true); - // std::cout << "---------------------------------------------------" << std::endl; - - // std::cout << "Confidence tensorRT" << std::endl; - // printDeviceVector(64, rt_out3, true); - // std::cout << "Locations tensorRT" << std::endl; - // printDeviceVector(64, rt_out4, true); - // std::cout << "---------------------------------------------------" << std::endl; - - // std::cout << "CUDNN vs TRT " << std::endl; - // ret_cudnn_tensorrt |= checkResult(conf->output_dim.tot(), conf->dstData, rt_out3) == 0 ? 0 : ERROR_CUDNNvsTENSORRT; - // ret_cudnn_tensorrt |= checkResult(loc->output_dim.tot(), loc->dstData, rt_out4) == 0 ? 0 : ERROR_CUDNNvsTENSORRT; - - // return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt; + std::cout << "CUDNN vs TRT " << std::endl; + ret_cudnn_tensorrt |= checkResult(odim1, cudnn_out, rt_out1) == 0 ? 0 : ERROR_CUDNNvsTENSORRT; cv::Mat viz = vizLayer2Mat(&net, net.num_layers-1); cv::imwrite("test.png", viz); + + return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt; }