From e99b353d8bc2670e1f0a7aee1e9c6242192a99ae Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Thu, 19 Dec 2019 14:47:57 +0100 Subject: [PATCH 1/2] Fix the inference operation of the deformable convolutional layer. This commit removes the malloc operation in the inference method and adds the sigmoid kernel. Signed-off-by: Davide Sapienza --- include/tkDNN/Layer.h | 5 + include/tkDNN/kernels.h | 1 + src/DeformConv2d.cpp | 156 +++++++++--------------------- src/kernels/activation_sigmoid.cu | 31 ++++++ src/kernels/deformable_conv.cu | 67 +++---------- 5 files changed, 97 insertions(+), 163 deletions(-) create mode 100644 src/kernels/activation_sigmoid.cu diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 8c501f4..188cffb 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -219,7 +219,12 @@ public: int kernelH, kernelW, strideH, strideW, paddingH, paddingW; protected: + dnnType *ones_d1; + dnnType *ones_d2; cudnnTensorDescriptor_t biasTensorDesc; + int chunk_dim; + dnnType *offset, *mask; + dnnType *output_conv; void initCUDNN(); diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index dfff6e3..fa46efb 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -6,6 +6,7 @@ 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 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)); void fill(dnnType* data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0)); diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp index f135b40..712419d 100644 --- a/src/DeformConv2d.cpp +++ b/src/DeformConv2d.cpp @@ -14,9 +14,36 @@ void DeformConv2d::initCUDNN() { net->tensorFormat, net->dataType, 1, output_dim.c, 1, 1) ); - checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, + checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, net->tensorFormat, net->dataType, output_dim.n, output_dim.c, output_dim.h, output_dim.w)); + const int height_ones = (preconv->input_dim.h + 2 * this->paddingH - (1 * (this->kernelH - 1) + 1)) / this->strideH + 1; + const int width_ones = (preconv->input_dim.w + 2 * this->paddingW - (1 * (this->kernelW - 1) + 1)) / this->strideW + 1; + const int dim_ones = preconv->input_dim.c * this->kernelH * this->kernelW * 1 * height_ones * width_ones; + + int dst_dim = preconv->output_dim.tot(); + if (dst_dim % 3 != 0 ) + std::cout<<"take attention\n\n"; + chunk_dim = dst_dim/3; + checkCuda(cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); + checkCuda(cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); + + // kernel ones + + cudaMallocHost(&ones_d1, (height_ones*width_ones)*sizeof(dnnType)); + float aus1[height_ones*width_ones]; + for(int i=0; iinfer(dim, srcData); - dim = preconv->output_dim; - - //split to chank - dnnType *offset, *mask; - int dst_dim = dim.tot(); - if (dst_dim % 3 != 0 ) - std::cout<<"take attention\n\n"; - int chunk_dim = dst_dim/3; - checkCuda(cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); - checkCuda(cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); - cudaDeviceSynchronize(); - - Conv2dToChunk(chunk_dim, srcData, offset, mask); + // conv2d + output_conv = preconv->infer(dim, srcData); + // split conv2d outputs into offset to mask + checkCuda(cudaMemcpy(offset, output_conv, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); // kernel sigmoide - dnnType *vec; - vec = new dnnType[chunk_dim]; - cudaDeviceSynchronize(); - cudaMemcpy(vec, mask, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToHost); - cudaDeviceSynchronize(); - for(int i=0; iinput_dim.h + 2 * this->paddingH - (1 * (this->kernelH - 1) + 1)) / this->strideH + 1; - const int width_ones = (preconv->input_dim.w + 2 * this->paddingW - (1 * (this->kernelW - 1) + 1)) / this->strideW + 1; - const int dim_ones = preconv->input_dim.c * this->kernelH * this->kernelW * 1 * height_ones * width_ones; - - // kernel ones - dnnType *ones_d1; - cudaMallocHost(&ones_d1, (height_ones*width_ones)*sizeof(dnnType)); - float aus1[height_ones*width_ones]; - for(int i=0; idata_d, + // deformable convolution + dcn_v2_cuda_forward(srcData, this->data_d, this->bias2_d, ones_d1, offset, mask, dstData, ones_d2, @@ -148,44 +107,18 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { this->deformableGroup, preconv->input_dim.n, preconv->input_dim.c, preconv->input_dim.h, preconv->input_dim.w, this->output_dim.n, this->output_dim.c, this->output_dim.h, this->output_dim.w, - dst_dim); - - cudaFree(offset); - cudaFree(mask); - cudaFree(input); - cudaFreeHost(ones_d1); - cudaFreeHost(ones_d2); - - - // dnnType *aus3; - // cudaMallocHost(&aus3, 256*7*7*sizeof(dnnType)); - // cudaMemcpy(aus3, dstData, (256*7*7)*sizeof(dnnType), cudaMemcpyDeviceToHost); - // checkCuda(cudaDeviceSynchronize()); - // std::cout<<"OutDim:\n"; - // this->output_dim.print(); - // std::cout<<"\n\n\nprint dstData: \n"; - // for (int i = 0 ; i < 256*7*7; i++){ - // if(i==294) - // std::cout<<"\n\n\n"; - // std::cout<cudnnHandle, &alpha, biasTensorDesc, bias_d, &beta, dstTensorDesc, dstData) ); } else { - std::cout<<"LOL\n"; alpha = dnnType(1); beta = dnnType(0); checkCUDNN( cudnnBatchNormalizationForwardInference(net->cudnnHandle, @@ -195,9 +128,8 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { scales_d, bias_d, mean_d, variance_d, CUDNN_BN_MIN_EPSILON) ); } + //update data dimensions - std::cout<<"dstData BN:\n"; - printDeviceVector(64, dstData); dim = output_dim; return dstData; } diff --git a/src/kernels/activation_sigmoid.cu b/src/kernels/activation_sigmoid.cu new file mode 100644 index 0000000..ba6c997 --- /dev/null +++ b/src/kernels/activation_sigmoid.cu @@ -0,0 +1,31 @@ +#include "kernels.h" + +__device__ +__forceinline__ +double sigmoid (double a) +{ + return 1.0 / (1.0 + exp (-a)); +} + + +__global__ +void activation_sigmoid(dnnType *input, dnnType *output, int size) { + + int stride = gridDim.x * blockDim.x; + int tid = blockDim.x * blockIdx.x + threadIdx.x; + for (int i = tid; i < size; i += stride) { + output[i] = sigmoid (input[i]); + } + } + + +/** + ELU activation function +*/ +void activationSIGMOIDForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream) +{ + int blocks = (size+255)/256; + int threads = 256; + + activation_sigmoid<<>>(srcData, dstData, size); +} \ No newline at end of file diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index 71efb0d..cb9ded0 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -149,10 +149,8 @@ void dcn_v2_cuda_forward(float *input, float *weight, const int deformable_group, const int in_n, const int in_c, const int in_h, const int in_w, const int out_n, const int out_c, const int out_h, const int out_w, - const int dst_dim, cudaStream_t stream) -{ - checkCuda(cudaDeviceSynchronize()); - cudaError_t cudaStat; + const int chunk_dim, cudaStream_t stream) +{ cublasStatus_t stat; cublasHandle_t handle; stat = cublasCreate(&handle); @@ -160,83 +158,50 @@ void dcn_v2_cuda_forward(float *input, float *weight, printf ("CUBLAS initialization failed\n"); return; } - checkCuda(cudaDeviceSynchronize()); - - const int batch = in_n; + const int channels = in_c; const int height = in_h; const int width = in_w; + const int channels_out = out_c; - const int channels_kernel = in_c; - const int kernel_h_ = kernel_h; - const int kernel_w_ = kernel_w; const int height_out = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1; const int width_out = (width + 2 * pad_w - (dilation_w * (kernel_w - 1) + 1)) / stride_w + 1; - - float *input_n; - cudaMalloc(&input_n, (in_n*in_c*in_h*in_w)*sizeof(float)); - cudaMemcpy(input_n, input, (in_n*in_c*in_h*in_w)*sizeof(float), cudaMemcpyDeviceToDevice); - checkCuda(cudaDeviceSynchronize()); - - float *offset_n; - cudaMalloc(&offset_n, ((dst_dim/3)*2)*sizeof(float)); - cudaMemcpy(offset_n, offset, ((dst_dim/3)*2)*sizeof(float), cudaMemcpyDeviceToDevice); - checkCuda(cudaDeviceSynchronize()); - - float *mask_n; - cudaMalloc(&mask_n, (dst_dim/3)*sizeof(float)); - cudaMemcpy(mask_n, mask, (dst_dim/3)*sizeof(float), cudaMemcpyDeviceToDevice); - checkCuda(cudaDeviceSynchronize()); - - float *output_n; - checkCuda(cudaMalloc(&output_n, (channels_out*height_out*width_out)*sizeof(float))); - checkCuda(cudaDeviceSynchronize()); - long m_ = channels_out; - long n_ = height_out * width_out; - long k_ = 1; + long m = channels_out; + long n = height_out * width_out; + long k = 1; float alpha = 1.0; float beta = 0.0; - checkCuda(cudaDeviceSynchronize()); + stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, - n_, m_, k_, &alpha, - ones, k_, bias, k_, - &beta, output_n, n_); + n, m, k, &alpha, + ones, k, bias, k, + &beta, output, n); if (stat != CUBLAS_STATUS_SUCCESS) { printf ("CUBLAS initialization failed\n"); return ; } - checkCuda(cudaDeviceSynchronize()); - modulated_deformable_im2col_cuda(stream, - input_n, offset_n, - mask_n, + input, offset, + mask, 1, channels, height, width, height_out, width_out, kernel_h, kernel_w, pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w, deformable_group, columns); - checkCuda(cudaDeviceSynchronize()); - + //(k * m) x (m * n) // Y = WC - long m = channels_out; - long n = height_out * width_out; - long k = channels * kernel_h * kernel_w; - - alpha = 1.0; + k = channels * kernel_h * kernel_w; beta = 1.0; stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, m, k, &alpha, columns, n, weight, k, - &beta, output_n, n); + &beta, output, n); - cudaMemcpy(output, output_n, (n*m)*sizeof(float), cudaMemcpyDeviceToDevice); - checkCuda(cudaDeviceSynchronize()); - if (stat != CUBLAS_STATUS_SUCCESS) { printf ("CUBLAS initialization failed\n"); return ; From d889ed385d55ec433468a2e49bcda0a6e67fdcf3 Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Fri, 20 Dec 2019 11:05:03 +0100 Subject: [PATCH 2/2] CenterNet TensorRT works. TensorRT serialization not yet implemented Signed-oof-by: Davide Sapienza --- include/tkDNN/Layer.h | 13 ++-- include/tkDNN/NetworkRT.h | 2 + include/tkDNN/pluginsRT/DeformableConvRT.h | 80 ++++++++++++++++++++++ src/Conv2d.cpp | 4 +- src/DeformConv2d.cpp | 16 ++--- src/Layer.cpp | 4 +- src/LayerWgs.cpp | 4 +- src/NetworkRT.cpp | 53 +++++++++++++- tests/resnet101_cnet/resnet101_cnet.cpp | 26 +++---- 9 files changed, 167 insertions(+), 35 deletions(-) create mode 100644 include/tkDNN/pluginsRT/DeformableConvRT.h diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 188cffb..a6c639b 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -32,7 +32,7 @@ enum layerType_t { class Layer { public: - Layer(Network *net); + Layer(Network *net, bool final = false); virtual ~Layer(); virtual layerType_t getLayerType() = 0; @@ -43,6 +43,7 @@ public: dataDim_t input_dim, output_dim; dnnType *dstData; //where results will be putted + bool final; //if the layer is the final one std::string getLayerName() { layerType_t type = getLayerType(); @@ -80,7 +81,7 @@ class LayerWgs : public Layer { public: LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kt, - std::string fname_weights, bool batchnorm = false, bool additional_bias = false); + std::string fname_weights, bool batchnorm = false, bool additional_bias = false, bool final = false); virtual ~LayerWgs(); int inputs, outputs; @@ -160,7 +161,7 @@ class Conv2d : public LayerWgs { public: Conv2d( Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, - std::string fname_weights, bool batchnorm = false, bool deConv = false); + std::string fname_weights, bool batchnorm = false, bool deConv = false, bool final = false); virtual ~Conv2d(); virtual layerType_t getLayerType() { return LAYER_CONV2D; }; @@ -217,15 +218,15 @@ public: int out_ch; int deformableGroup; int kernelH, kernelW, strideH, strideW, paddingH, paddingW; -protected: - dnnType *ones_d1; dnnType *ones_d2; - cudnnTensorDescriptor_t biasTensorDesc; int chunk_dim; dnnType *offset, *mask; dnnType *output_conv; +protected: + + cudnnTensorDescriptor_t biasTensorDesc; void initCUDNN(); }; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index d30da03..590b323 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -31,6 +31,7 @@ using namespace nvinfer1; #include "pluginsRT/YoloRT.h" #include "pluginsRT/UpsampleRT.h" //#include "pluginsRT/Int8Calibrator.h" +#include "pluginsRT/DeformableConvRT.h" class PluginFactory : IPluginFactory { @@ -85,6 +86,7 @@ public: nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Yolo *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Upsample *l); + nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, DeformConv2d *l); bool serialize(const char *filename); bool deserialize(const char *filename); diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h new file mode 100644 index 0000000..6a85538 --- /dev/null +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -0,0 +1,80 @@ +#include +#include "../kernels.h" + + +class DeformableConvRT : public IPlugin { + + + +public: + DeformableConvRT(tk::dnn::DeformConv2d *deformable) { + this->defRT = deformable; + } + + ~DeformableConvRT(){ + + } + + int getNbOutputs() const override { + return 1; + } + + Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override { + return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; + } + + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { + } + + int initialize() override { + + return 0; + } + + virtual void terminate() override { + } + + virtual size_t getWorkspaceSize(int maxBatchSize) const override { + return 0; + } + + virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { + + dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); + dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); + + // split conv2d outputs into offset to mask + checkCuda(cudaMemcpy(defRT->offset, defRT->output_conv, 2*defRT->chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + checkCuda(cudaMemcpy(defRT->mask, defRT->output_conv + 2*defRT->chunk_dim, defRT->chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + // kernel sigmoide + activationSIGMOIDForward(defRT->mask, defRT->mask, defRT->chunk_dim); + + // deformable convolution + dcn_v2_cuda_forward(srcData, defRT->data_d, + defRT->bias2_d, defRT->ones_d1, + defRT->offset, defRT->mask, + reinterpret_cast(outputs[0]), defRT->ones_d2, + defRT->kernelH, defRT->kernelW, + defRT->strideH, defRT->strideW, + defRT->paddingH, defRT->paddingW, + 1, 1, + defRT->deformableGroup, + defRT->preconv->input_dim.n, defRT->preconv->input_dim.c, defRT->preconv->input_dim.h, defRT->preconv->input_dim.w, + defRT->output_dim.n, defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w, + defRT->chunk_dim); + + return 0; + } + + + virtual size_t getSerializationSize() override { + return 0; + } + + virtual void serialize(void* buffer) override { + char *buf = reinterpret_cast(buffer); + } + + int size; + tk::dnn::DeformConv2d *defRT; +}; diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 05fec83..b8296ac 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -119,10 +119,10 @@ void Conv2d::inferCUDNN(dnnType* srcData, bool back) { Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, - std::string fname_weights, bool batchnorm, bool deConv) : + std::string fname_weights, bool batchnorm, bool deConv, bool final) : LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, - fname_weights, batchnorm) { + fname_weights, batchnorm, false, final) { this->kernelH = kernelH; this->kernelW = kernelW; diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp index 712419d..3d1fad2 100644 --- a/src/DeformConv2d.cpp +++ b/src/DeformConv2d.cpp @@ -25,25 +25,23 @@ void DeformConv2d::initCUDNN() { if (dst_dim % 3 != 0 ) std::cout<<"take attention\n\n"; chunk_dim = dst_dim/3; - checkCuda(cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); - checkCuda(cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); + checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); + checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); // kernel ones - cudaMallocHost(&ones_d1, (height_ones*width_ones)*sizeof(dnnType)); + checkCuda( cudaMalloc(&ones_d1, (height_ones*width_ones)*sizeof(dnnType)) ); float aus1[height_ones*width_ones]; for(int i=0; inet = net; - + this->final = final; if(net != nullptr) { this->input_dim = net->getOutputDim(); this->output_dim = input_dim; diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index 0f623c2..2f7c7fc 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -8,12 +8,12 @@ namespace tk { namespace dnn { LayerWgs::LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kl, - std::string fname_weights, bool batchnorm, bool additional_bias) : Layer(net) { + std::string fname_weights, bool batchnorm, bool additional_bias, bool final) : Layer(net, final) { this->inputs = inputs; this->outputs = outputs; this->weights_path = std::string(fname_weights); - + std::cout<<"Reading weights: I="<dontLoadWeights); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 4500be5..884a4a3 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -75,7 +75,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { input = Ilay->getOutput(0); input->setName( (l->getLayerName() + std::to_string(i) + "_out").c_str() ); - if(l->getLayerType() == LAYER_YOLO) + if(l->getLayerType() == LAYER_YOLO || l->final) networkRT->markOutput(*input); tensors[l] = input; } @@ -182,6 +182,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { return convert_layer(input, (Yolo*) l); if(type == LAYER_UPSAMPLE) return convert_layer(input, (Upsample*) l); + if(type == LAYER_DEFORMCONV2D) + return convert_layer(input, (DeformConv2d*) l); std::cout<getLayerName()<<"\n"; FatalError("Layer not implemented in tensorRT"); @@ -254,6 +256,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); lRT = (ILayer*) lRTconv; + Dims d = lRTconv->getOutput(0)->getDimensions(); + std::cout<<"DECONV: "<preconv); + + ITensor **inputs = new ITensor*[2]; + inputs[0] = input; + inputs[1] = preconv->getOutput(0); + + //std::cout<<"New plugin DEFORMABLE\n"; + IPlugin *plugin = new DeformableConvRT(l); + IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin); + checkNULL(lRT); + + // batchnorm + void *bias_b, *power_b, *mean_b, *variance_b, *scales_b; + if(dtRT == DataType::kHALF) { + bias_b = l->bias16_h; + power_b = l->power16_h; + mean_b = l->mean16_h; + variance_b = l->variance16_h; + scales_b = l->scales16_h; + } else { + bias_b = l->bias_h; + power_b = l->power_h; + mean_b = l->mean_h; + variance_b = l->variance_h; + scales_b = l->scales_h; + } + + Weights power{dtRT, power_b, l->outputs}; + Weights shift{dtRT, mean_b, l->outputs}; + Weights scale{dtRT, variance_b, l->outputs}; + std::cout<getNbOutputs()<addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, + shift, scale, power); + + checkNULL(lRT2); + + Weights shift2{dtRT, bias_b, l->outputs}; + Weights scale2{dtRT, scales_b, l->outputs}; + IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, + shift2, scale2, power); + checkNULL(lRT3); + + return lRT3; +} + bool NetworkRT::serialize(const char *filename) { std::ofstream p(filename); diff --git a/tests/resnet101_cnet/resnet101_cnet.cpp b/tests/resnet101_cnet/resnet101_cnet.cpp index 33d60b8..c89591a 100644 --- a/tests/resnet101_cnet/resnet101_cnet.cpp +++ b/tests/resnet101_cnet/resnet101_cnet.cpp @@ -172,9 +172,9 @@ const char *reg_conv2_bin = "../tests/resnet101_cnet/layers/reg-2.bin"; const char *fc_bin = "../tests/resnet101_cnet/layers/fc.bin"; const char *output_bin[]={ -"../tests/resnet101_cnet/debug/hm.bin", -"../tests/resnet101_cnet/debug/wh.bin", -"../tests/resnet101_cnet/debug/reg.bin"}; + "../tests/resnet101_cnet/debug/hm.bin", + "../tests/resnet101_cnet/debug/wh.bin", + "../tests/resnet101_cnet/debug/reg.bin"}; int main() { @@ -317,17 +317,17 @@ int main() tk::dnn::Layer *route_1_0_layers[1] = { layer2_deconv1_relu }; tk::dnn::Conv2d *hm_conv1 = new tk::dnn::Conv2d(&net, 64, 3, 3, 1, 1, 1, 1, hm_conv1_bin, false); tk::dnn::Activation *hm_relu1 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); - tk::dnn::Conv2d *hm = new tk::dnn::Conv2d(&net, 80, 1, 1, 1, 1, 0, 0, hm_conv2_bin, false); + tk::dnn::Conv2d *hm = new tk::dnn::Conv2d(&net, 80, 1, 1, 1, 1, 0, 0, hm_conv2_bin, false, false, true); tk::dnn::Route *route_1_0 = new tk::dnn::Route(&net, route_1_0_layers, 1); tk::dnn::Conv2d *wh_conv1 = new tk::dnn::Conv2d(&net, 64, 3, 3, 1, 1, 1, 1, wh_conv1_bin, false); tk::dnn::Activation *wh_relu1 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); - tk::dnn::Conv2d *wh = new tk::dnn::Conv2d(&net, 2, 1, 1, 1, 1, 0, 0, wh_conv2_bin, false); + tk::dnn::Conv2d *wh = new tk::dnn::Conv2d(&net, 2, 1, 1, 1, 1, 0, 0, wh_conv2_bin, false, false, true); tk::dnn::Route *route_2_0 = new tk::dnn::Route(&net, route_1_0_layers, 1); tk::dnn::Conv2d *reg_conv1 = new tk::dnn::Conv2d(&net, 64, 3, 3, 1, 1, 1, 1, reg_conv1_bin, false); tk::dnn::Activation *reg_relu1 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); - tk::dnn::Conv2d *reg = new tk::dnn::Conv2d(&net, 2, 1, 1, 1, 1, 0, 0, reg_conv2_bin, false); + tk::dnn::Conv2d *reg = new tk::dnn::Conv2d(&net, 2, 1, 1, 1, 1, 0, 0, reg_conv2_bin, false, false, true); // Load input dnnType *data; @@ -339,7 +339,7 @@ int main() net.print(); //convert network to tensorRT -// tk::dnn::NetworkRT netRT(&net, "resnet101_cnet.rt"); + tk::dnn::NetworkRT netRT(&net, "resnet101_cnet.rt"); tk::dnn::dataDim_t dim1 = dim; //input dim @@ -354,7 +354,7 @@ int main() // printDeviceVector(64, cudnn_out, true); -/* tk::dnn::dataDim_t dim2 = dim; + tk::dnn::dataDim_t dim2 = dim; printCenteredTitle(" TENSORRT inference ", '=', 30); { dim2.print(); @@ -363,10 +363,9 @@ int main() TIMER_STOP dim2.print(); } - rt_out = (dnnType *)netRT.buffersRT[1]; -*/ - tk::dnn::Conv2d *outs[3] = { hm, wh, reg }; + tk::dnn::Layer *outs[3] = { hm, wh, reg }; + for(int i=0; i<3; i++) { printCenteredTitle((std::string(" RESNET CHECK RESULTS ") + std::to_string(i) + " ").c_str(), '=', 30); @@ -382,14 +381,15 @@ int main() dnnType *cudnn_out, *rt_out; cudnn_out = outs[i]->dstData; + rt_out = (dnnType *)netRT.buffersRT[i+1]; std::cout << "CUDNN vs correct"; checkResult(odim, cudnn_out, out); - /* std::cout << "TRT vs correct"; + std::cout << "TRT vs correct"; checkResult(odim, rt_out, out); std::cout << "CUDNN vs TRT "; - checkResult(odim, cudnn_out, rt_out);*/ + checkResult(odim, cudnn_out, rt_out); } return 0; }