From df888a345731ceb50560f6e3f35a218b00d09c8e Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Wed, 18 Dec 2019 18:59:54 +0100 Subject: [PATCH] Add CenterNet based on Resnet101, TensorRT not implemented. Signed-off-by: Davide Sapienza --- CMakeLists.txt | 3 + include/tkDNN/Layer.h | 33 +- include/tkDNN/Network.h | 2 +- include/tkDNN/kernels.h | 13 + src/Activation.cpp | 1 - src/Conv2d.cpp | 11 +- src/DeformConv2d.cpp | 206 ++++++++++ src/LayerWgs.cpp | 8 +- src/kernels/deformable_conv.cu | 108 +++++ src/utils.cpp | 1 - tests/resnet101/resnet101.cpp | 1 - tests/resnet101/resnet101_weightsexporter.py | 2 + tests/resnet101_cnet/resnet101_cnet.cpp | 395 +++++++++++++++++++ 13 files changed, 773 insertions(+), 11 deletions(-) create mode 100644 src/DeformConv2d.cpp create mode 100644 tests/resnet101_cnet/resnet101_cnet.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index f50b96a..8043e4e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -102,6 +102,9 @@ target_link_libraries(test_yolo3_flir tkDNN) add_executable(test_resnet101 tests/resnet101/resnet101.cpp) target_link_libraries(test_resnet101 tkDNN) +add_executable(test_resnet101_cnet tests/resnet101_cnet/resnet101_cnet.cpp) +target_link_libraries(test_resnet101_cnet tkDNN) + ################################################################################ diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index d001497..8c501f4 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -12,6 +12,7 @@ enum layerType_t { LAYER_DENSE, LAYER_CONV2D, LAYER_DECONV2D, + LAYER_DEFORMCONV2D, LAYER_ACTIVATION, LAYER_FLATTEN, LAYER_MULADD, @@ -49,6 +50,7 @@ public: case LAYER_DENSE: return "Dense"; case LAYER_CONV2D: return "Conv2d"; case LAYER_DECONV2D: return "DeConv2d"; + case LAYER_DEFORMCONV2D:return "DeformConv2d"; case LAYER_ACTIVATION: return "Activation"; case LAYER_FLATTEN: return "Flatten"; case LAYER_MULADD: return "MulAdd"; @@ -78,7 +80,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); + std::string fname_weights, bool batchnorm = false, bool additional_bias = false); virtual ~LayerWgs(); int inputs, outputs; @@ -87,6 +89,10 @@ public: dnnType *data_h, *data_d; dnnType *bias_h, *bias_d; + // additional bias for DCN + bool additional_bias; + dnnType *bias2_h, *bias2_d; + //batchnorm bool batchnorm; dnnType *power_h; @@ -194,6 +200,31 @@ public: }; +/** + Deformable Convolutionl 2d layer +*/ +class DeformConv2d : public LayerWgs { + +public: + DeformConv2d( Network *net, int out_ch, int deformable_group, int kernelH, int kernelW, + int strideH, int strideW, int paddingH, int paddingW, + std::string d_fname_weights, std::string fname_weights, bool batchnorm); + virtual ~DeformConv2d(); + virtual layerType_t getLayerType() { return LAYER_DEFORMCONV2D; }; + + virtual dnnType* infer(dataDim_t &dim, dnnType* srcData); + tk::dnn::Conv2d *preconv; + int out_ch; + int deformableGroup; + int kernelH, kernelW, strideH, strideW, paddingH, paddingW; +protected: + + cudnnTensorDescriptor_t biasTensorDesc; + + void initCUDNN(); + +}; + /** Flatten layer is actually a matrix transposition diff --git a/include/tkDNN/Network.h b/include/tkDNN/Network.h index c56bd2c..b234f71 100644 --- a/include/tkDNN/Network.h +++ b/include/tkDNN/Network.h @@ -32,7 +32,7 @@ struct dataDim_t { }; class Layer; -const int MAX_LAYERS = 256; +const int MAX_LAYERS = 512; class Network { diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index ff4a334..dfff6e3 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -33,4 +33,17 @@ void modulated_deformable_im2col_cuda(cudaStream_t stream, const int pad_h, const int pad_w, const int stride_h, const int stride_w, const int dilation_h, const int dilation_w, const int deformable_group, float *data_col); + +void dcn_v2_cuda_forward(float *input, float *weight, + float *bias, float *ones, + float *offset, float *mask, + float *output, float *columns, + int kernel_h, int kernel_w, + const int stride_h, const int stride_w, + const int pad_h, const int pad_w, + const int dilation_h, const int dilation_w, + 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 = cudaStream_t(0)); #endif //KERNELS_H diff --git a/src/Activation.cpp b/src/Activation.cpp index 3c43e68..1f6f0e7 100644 --- a/src/Activation.cpp +++ b/src/Activation.cpp @@ -44,7 +44,6 @@ Activation::~Activation() { } dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) { - if(act_mode == ACTIVATION_LEAKY) { activationLEAKYForward(srcData, dstData, dim.tot()); diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index cc4fbd5..05fec83 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -44,7 +44,8 @@ void Conv2d::initCUDNN(bool back) { convDesc, srcTensor, filterDesc, &tmpdim.n, &tmpdim.c, &tmpdim.h, &tmpdim.w) ); if(odim.n != tmpdim.n || odim.c != tmpdim.c || odim.h != tmpdim.h || odim.w != tmpdim.w) { - std::cout<<"tkdim: "; odim.print(); + std::cout<<"tkdim input: "; idim.print(); + std::cout<<"tkdim output: "; odim.print(); std::cout<<"cudnndim: "; tmpdim.print(); FatalError("Eror conv dimension mismatch"); } @@ -107,12 +108,12 @@ void Conv2d::inferCUDNN(dnnType* srcData, bool back) { } else { alpha = dnnType(1); beta = dnnType(0); - cudnnBatchNormalizationForwardInference(net->cudnnHandle, + checkCUDNN( cudnnBatchNormalizationForwardInference(net->cudnnHandle, CUDNN_BATCHNORM_SPATIAL, &alpha, &beta, dstTensorDesc, dstData, dstTensorDesc, dstData, biasTensorDesc, //same tensor descriptor as bias scales_d, bias_d, mean_d, variance_d, - CUDNN_BN_MIN_EPSILON); + CUDNN_BN_MIN_EPSILON) ); } } @@ -140,8 +141,8 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, } else { output_dim.n = input_dim.n; output_dim.c = out_ch; - output_dim.h = (input_dim.h * strideH) - 2*paddingH + kernelH -1; - output_dim.w = (input_dim.w * strideW) - 2*paddingW + kernelW -1; + output_dim.h = ((input_dim.h-1) * strideH) - 2*paddingH + kernelH; + output_dim.w = ((input_dim.w-1) * strideW) - 2*paddingW + kernelW; output_dim.l = 1; } initCUDNN(deConv); diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp new file mode 100644 index 0000000..f135b40 --- /dev/null +++ b/src/DeformConv2d.cpp @@ -0,0 +1,206 @@ +#include + +#include "Layer.h" +#include "kernels.h" +#include + + +namespace tk { namespace dnn { + +void DeformConv2d::initCUDNN() { + + checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) ); + checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc, + net->tensorFormat, net->dataType, + 1, output_dim.c, 1, 1) ); + + checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, + net->tensorFormat, net->dataType, output_dim.n, output_dim.c, output_dim.h, output_dim.w)); + +} + +DeformConv2d::DeformConv2d( Network *net, int out_ch, int deformable_group, int kernelH, int kernelW, + int strideH, int strideW, int paddingH, int paddingW, + std::string d_fname_weights, std::string fname_weights, bool batchnorm) : + + LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, + d_fname_weights, batchnorm, true){ + + this->out_ch = out_ch; + this->deformableGroup = deformable_group; + this->kernelH = kernelH; + this->kernelW = kernelW; + this->strideH = strideH; + this->strideW = strideW; + this->paddingH = paddingH; + this->paddingW = paddingW; + + preconv = new tk::dnn::Conv2d(net, deformable_group * 3 * kernelH * kernelW, kernelH, kernelW, + strideH, strideW, paddingH, paddingW, fname_weights, false); + net->num_layers--; + + output_dim = preconv->output_dim; + + output_dim.c = out_ch; + initCUDNN(); + //allocate data for infer result + checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); +} + +DeformConv2d::~DeformConv2d() { + + checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) ); + checkCuda( cudaFree(dstData) ); +} + +void Conv2dToChunk(int dim, dnnType* srcData, dnnType* offset, dnnType* mask) +{ + // std::cout<<"9\n"; + // cudaMemcpyFromArray(offset, (const struct cudaArray *)srcData, 0, 2*dim.tot()/3, dim.tot()/3, cudaMemcpyDeviceToHost); + // checkCuda(cudaMemcpyFromArray(offset, (const struct cudaArray *)srcData, 0, 0, 2*dim, cudaMemcpyDeviceToDevice)); + checkCuda(cudaMemcpy(offset, srcData, 2*dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + + // std::cout<<"9a\n"; + cudaDeviceSynchronize(); + // cudaMemcpyFromArray(mask, (const struct cudaArray *)srcData, 2*dim.tot()/3, dim.tot(), dim.tot()/3, cudaMemcpyDeviceToHost); + // checkCuda(cudaMemcpyFromArray(mask, (const struct cudaArray *)srcData, 0, 2*dim, dim, cudaMemcpyDeviceToDevice)); + checkCuda(cudaMemcpy(mask, srcData + 2*dim, dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + // std::cout<<"9b\n"; +} + +dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { + dnnType *input; + checkCuda(cudaMalloc(&input, dim.tot()*sizeof(dnnType))); + checkCuda(cudaMemcpy(input, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + cudaDeviceSynchronize(); + srcData = preconv->infer(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); + + // 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, + this->bias2_d, ones_d1, + offset, mask, + dstData, ones_d2, + this->kernelH, this->kernelW, + this->strideH, this->strideW, + this->paddingH, this->paddingW, + 1, 1, + 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, + CUDNN_BATCHNORM_SPATIAL, &alpha, &beta, + dstTensorDesc, dstData, dstTensorDesc, + dstData, biasTensorDesc, //same tensor descriptor as bias + 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/LayerWgs.cpp b/src/LayerWgs.cpp index 8c739b0..0f623c2 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -8,7 +8,7 @@ namespace tk { namespace dnn { LayerWgs::LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kl, - std::string fname_weights, bool batchnorm) : Layer(net) { + std::string fname_weights, bool batchnorm, bool additional_bias) : Layer(net) { this->inputs = inputs; this->outputs = outputs; @@ -18,6 +18,12 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs, int seek = 0; readBinaryFile(weights_path.c_str(), inputs*outputs*kh*kw*kl, &data_h, &data_d, seek, net->dontLoadWeights); seek += inputs*outputs*kh*kw*kl; + this->additional_bias = additional_bias; + if(additional_bias) { + readBinaryFile(weights_path.c_str(), outputs, &bias2_h, &bias2_d, seek, net->dontLoadWeights); + seek += outputs; + } + readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek, net->dontLoadWeights); this->batchnorm = batchnorm; diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index 22b4dab..71efb0d 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -1,6 +1,8 @@ #include #include #include +#include "kernels.h" +#include #define CUDA_KERNEL_LOOP(i, n) \ for (int i = blockIdx.x * blockDim.x + threadIdx.x; \ @@ -134,3 +136,109 @@ void modulated_deformable_im2col_cuda(cudaStream_t stream, } } + + +void dcn_v2_cuda_forward(float *input, float *weight, + float *bias, float *ones, + float *offset, float *mask, + float *output, float *columns, + int kernel_h, int kernel_w, + const int stride_h, const int stride_w, + const int pad_h, const int pad_w, + const int dilation_h, const int dilation_w, + 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; + cublasStatus_t stat; + cublasHandle_t handle; + stat = cublasCreate(&handle); + if (stat != CUBLAS_STATUS_SUCCESS) { + 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; + 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_); + if (stat != CUBLAS_STATUS_SUCCESS) { + printf ("CUBLAS initialization failed\n"); + return ; + } + + checkCuda(cudaDeviceSynchronize()); + + modulated_deformable_im2col_cuda(stream, + input_n, offset_n, + mask_n, + 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; + beta = 1.0; + + stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, + n, m, k, &alpha, + columns, n, weight, k, + &beta, output_n, n); + + cudaMemcpy(output, output_n, (n*m)*sizeof(float), cudaMemcpyDeviceToDevice); + checkCuda(cudaDeviceSynchronize()); + + if (stat != CUBLAS_STATUS_SUCCESS) { + printf ("CUBLAS initialization failed\n"); + return ; + } +} diff --git a/src/utils.cpp b/src/utils.cpp index 250cadc..deb4def 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -89,7 +89,6 @@ int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) { data_h = data_d; correct_h = correct_d; } - int diffs = 0; for(int i=0; i +#include "tkdnn.h" + +const char *input_bin = "../tests/resnet101_cnet/debug/input.bin"; +const char *conv1_bin = "../tests/resnet101_cnet/layers/conv1.bin"; + +//layer1 +const char *layer1_bin[]={ +"../tests/resnet101_cnet/layers/layer1-0-conv1.bin", +"../tests/resnet101_cnet/layers/layer1-0-conv2.bin", +"../tests/resnet101_cnet/layers/layer1-0-conv3.bin", +"../tests/resnet101_cnet/layers/layer1-0-downsample-0.bin", + +"../tests/resnet101_cnet/layers/layer1-1-conv1.bin", +"../tests/resnet101_cnet/layers/layer1-1-conv2.bin", +"../tests/resnet101_cnet/layers/layer1-1-conv3.bin", + +"../tests/resnet101_cnet/layers/layer1-2-conv1.bin", +"../tests/resnet101_cnet/layers/layer1-2-conv2.bin", +"../tests/resnet101_cnet/layers/layer1-2-conv3.bin"}; + + +//layer2 +const char *layer2_bin[]={ +"../tests/resnet101_cnet/layers/layer2-0-conv1.bin", +"../tests/resnet101_cnet/layers/layer2-0-conv2.bin", +"../tests/resnet101_cnet/layers/layer2-0-conv3.bin", +"../tests/resnet101_cnet/layers/layer2-0-downsample-0.bin", + +"../tests/resnet101_cnet/layers/layer2-1-conv1.bin", +"../tests/resnet101_cnet/layers/layer2-1-conv2.bin", +"../tests/resnet101_cnet/layers/layer2-1-conv3.bin", + +"../tests/resnet101_cnet/layers/layer2-2-conv1.bin", +"../tests/resnet101_cnet/layers/layer2-2-conv2.bin", +"../tests/resnet101_cnet/layers/layer2-2-conv3.bin", + +"../tests/resnet101_cnet/layers/layer2-3-conv1.bin", +"../tests/resnet101_cnet/layers/layer2-3-conv2.bin", +"../tests/resnet101_cnet/layers/layer2-3-conv3.bin" +}; +//layer3 +const char *layer3_bin[]={ +"../tests/resnet101_cnet/layers/layer3-0-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-0-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-0-conv3.bin", +"../tests/resnet101_cnet/layers/layer3-0-downsample-0.bin", + +"../tests/resnet101_cnet/layers/layer3-1-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-1-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-1-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-2-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-2-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-2-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-3-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-3-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-3-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-4-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-4-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-4-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-5-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-5-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-5-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-6-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-6-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-6-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-7-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-7-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-7-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-8-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-8-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-8-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-9-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-9-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-9-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-10-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-10-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-10-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-11-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-11-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-11-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-12-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-12-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-12-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-13-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-13-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-13-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-14-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-14-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-14-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-15-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-15-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-15-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-16-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-16-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-16-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-17-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-17-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-17-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-18-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-18-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-18-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-19-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-19-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-19-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-20-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-20-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-20-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-21-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-21-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-21-conv3.bin", + +"../tests/resnet101_cnet/layers/layer3-22-conv1.bin", +"../tests/resnet101_cnet/layers/layer3-22-conv2.bin", +"../tests/resnet101_cnet/layers/layer3-22-conv3.bin"}; + + +//layer4 +const char *layer4_bin[]={ +"../tests/resnet101_cnet/layers/layer4-0-conv1.bin", +"../tests/resnet101_cnet/layers/layer4-0-conv2.bin", +"../tests/resnet101_cnet/layers/layer4-0-conv3.bin", +"../tests/resnet101_cnet/layers/layer4-0-downsample-0.bin", + +"../tests/resnet101_cnet/layers/layer4-1-conv1.bin", +"../tests/resnet101_cnet/layers/layer4-1-conv2.bin", +"../tests/resnet101_cnet/layers/layer4-1-conv3.bin", + +"../tests/resnet101_cnet/layers/layer4-2-conv1.bin", +"../tests/resnet101_cnet/layers/layer4-2-conv2.bin", +"../tests/resnet101_cnet/layers/layer4-2-conv3.bin"}; + +const char *d_conv1_bin = "../tests/resnet101_cnet/layers/deconv_layers-0-conv_offset_mask.bin"; +const char *deform1_bin = "../tests/resnet101_cnet/layers/deconv_layers-0.bin"; +const char *deconv1_bin = "../tests/resnet101_cnet/layers/deconv_layers-3.bin"; + +const char *d_conv2_bin = "../tests/resnet101_cnet/layers/deconv_layers-6-conv_offset_mask.bin"; +const char *deform2_bin = "../tests/resnet101_cnet/layers/deconv_layers-6.bin"; +const char *deconv2_bin = "../tests/resnet101_cnet/layers/deconv_layers-9.bin"; + +const char *d_conv3_bin = "../tests/resnet101_cnet/layers/deconv_layers-12-conv_offset_mask.bin"; +const char *deform3_bin = "../tests/resnet101_cnet/layers/deconv_layers-12.bin"; +const char *deconv3_bin = "../tests/resnet101_cnet/layers/deconv_layers-15.bin"; + +const char *hm_conv1_bin = "../tests/resnet101_cnet/layers/hm-0.bin"; +const char *hm_conv2_bin = "../tests/resnet101_cnet/layers/hm-2.bin"; +const char *wh_conv1_bin = "../tests/resnet101_cnet/layers/wh-0.bin"; +const char *wh_conv2_bin = "../tests/resnet101_cnet/layers/wh-2.bin"; +const char *reg_conv1_bin = "../tests/resnet101_cnet/layers/reg-0.bin"; +const char *reg_conv2_bin = "../tests/resnet101_cnet/layers/reg-2.bin"; +//final +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"}; + +int main() +{ + + // Network layout + tk::dnn::dataDim_t dim(1, 3, 224, 224, 1); + tk::dnn::Network net(dim); + + tk::dnn::Conv2d conv1(&net, 64, 7, 7, 2, 2, 3, 3, conv1_bin, true); + tk::dnn::Activation relu3(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Pooling maxpool4(&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX); + + + //layer 1 + int id_layer1_bin = 0; + tk::dnn::Layer *last = &maxpool4; + for(int i=0; i<3;i++) + { + tk::dnn::Conv2d *layer1_0_conv1 = new tk::dnn::Conv2d(&net, 64, 1, 1, 1, 1, 0, 0, layer1_bin[id_layer1_bin++], true); + tk::dnn::Activation *relu1_0_1 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d *layer1_0_conv2 = new tk::dnn::Conv2d(&net, 64, 3, 3, 1, 1, 1, 1, layer1_bin[id_layer1_bin++], true); + tk::dnn::Activation *relu1_0_2 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d *layer1_0_conv3 = new tk::dnn::Conv2d(&net, 256, 1, 1, 1, 1, 0, 0, layer1_bin[id_layer1_bin++], true); + if(i==0) { + tk::dnn::Layer *route_1_0_layers[1] = { last }; + tk::dnn::Route *route_1_0 = new tk::dnn::Route(&net, route_1_0_layers, 1); + tk::dnn::Conv2d *layer1_0_downsample_0 = new tk::dnn::Conv2d(&net, 256, 1, 1, 1, 1, 0, 0, layer1_bin[id_layer1_bin++], true); + tk::dnn::Shortcut *s1_0 = new tk::dnn::Shortcut(&net, layer1_0_conv3); + } else { + tk::dnn::Shortcut *s1_0 = new tk::dnn::Shortcut(&net, last); + } + tk::dnn::Activation *layer1_0_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + last = layer1_0_relu; + } + + // tk::dnn::Activation *last_activation = (tk::dnn::Activation *) net.layers[net.num_layers-1]; + // layer 2 + int id_layer2_bin = 0; + for(int i=0; i<4;i++) + { + tk::dnn::Conv2d *layer1_0_conv1 = new tk::dnn::Conv2d(&net, 128, 1, 1, 1, 1, 0, 0, layer2_bin[id_layer2_bin++], true); + tk::dnn::Activation *relu1_0_1 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d *layer1_0_conv2; + if(i==0) + layer1_0_conv2 = new tk::dnn::Conv2d(&net, 128, 3, 3, 2, 2, 1, 1, layer2_bin[id_layer2_bin++], true); + else + layer1_0_conv2 = new tk::dnn::Conv2d(&net, 128, 3, 3, 1, 1, 1, 1, layer2_bin[id_layer2_bin++], true); + + tk::dnn::Activation *relu1_0_2 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d *layer1_0_conv3 = new tk::dnn::Conv2d(&net, 512, 1, 1, 1, 1, 0, 0, layer2_bin[id_layer2_bin++], true); + if(i==0) + { + tk::dnn::Layer *route_1_0_layers[1] = { last }; + tk::dnn::Route *route_1_0 = new tk::dnn::Route(&net, route_1_0_layers, 1); + tk::dnn::Conv2d *layer1_0_downsample_0 = new tk::dnn::Conv2d(&net, 512, 1, 1, 2, 2, 0, 0, layer2_bin[id_layer2_bin++], true); + tk::dnn::Shortcut *s1_0 = new tk::dnn::Shortcut(&net, layer1_0_conv3); + } + else + { + tk::dnn::Shortcut *s1_0 = new tk::dnn::Shortcut(&net, last); + } + tk::dnn::Activation *layer1_0_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + last = layer1_0_relu; + } + + // layer 3 + int id_layer3_bin = 0; + for(int i=0; i<23;i++) + { + tk::dnn::Conv2d *layer1_0_conv1 = new tk::dnn::Conv2d(&net, 256, 1, 1, 1, 1, 0, 0, layer3_bin[id_layer3_bin++], true); + tk::dnn::Activation *relu1_0_1 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d *layer1_0_conv2; + if(i==0) + layer1_0_conv2 = new tk::dnn::Conv2d(&net, 256, 3, 3, 2, 2, 1, 1, layer3_bin[id_layer3_bin++], true); + else + layer1_0_conv2 = new tk::dnn::Conv2d(&net, 256, 3, 3, 1, 1, 1, 1, layer3_bin[id_layer3_bin++], true); + + tk::dnn::Activation *relu1_0_2 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d *layer1_0_conv3 = new tk::dnn::Conv2d(&net, 1024, 1, 1, 1, 1, 0, 0, layer3_bin[id_layer3_bin++], true); + if(i==0) + { + tk::dnn::Layer *route_1_0_layers[1] = { last }; + tk::dnn::Route *route_1_0 = new tk::dnn::Route(&net, route_1_0_layers, 1); + tk::dnn::Conv2d *layer1_0_downsample_0 = new tk::dnn::Conv2d(&net, 1024, 1, 1, 2, 2, 0, 0, layer3_bin[id_layer3_bin++], true); + tk::dnn::Shortcut *s1_0 = new tk::dnn::Shortcut(&net, layer1_0_conv3); + } + else + { + tk::dnn::Shortcut *s1_0 = new tk::dnn::Shortcut(&net, last); + } + tk::dnn::Activation *layer1_0_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + last = layer1_0_relu; + } + + // layer 4 + int id_layer4_bin = 0; + for(int i=0; i<3;i++) + { + tk::dnn::Conv2d *layer1_0_conv1 = new tk::dnn::Conv2d(&net, 512, 1, 1, 1, 1, 0, 0, layer4_bin[id_layer4_bin++], true); + tk::dnn::Activation *relu1_0_1 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d *layer1_0_conv2; + if(i==0) + layer1_0_conv2 = new tk::dnn::Conv2d(&net, 512, 3, 3, 2, 2, 1, 1, layer4_bin[id_layer4_bin++], true); + else + layer1_0_conv2 = new tk::dnn::Conv2d(&net, 512, 3, 3, 1, 1, 1, 1, layer4_bin[id_layer4_bin++], true); + + tk::dnn::Activation *relu1_0_2 = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::Conv2d *layer1_0_conv3 = new tk::dnn::Conv2d(&net, 2048, 1, 1, 1, 1, 0, 0, layer4_bin[id_layer4_bin++], true); + if(i==0) + { + tk::dnn::Layer *route_1_0_layers[1] = { last }; + tk::dnn::Route *route_1_0 = new tk::dnn::Route(&net, route_1_0_layers, 1); + tk::dnn::Conv2d *layer1_0_downsample_0 = new tk::dnn::Conv2d(&net, 2048, 1, 1, 2, 2, 0, 0, layer4_bin[id_layer4_bin++], true); + tk::dnn::Shortcut *s1_0 = new tk::dnn::Shortcut(&net, layer1_0_conv3); + } + else + { + tk::dnn::Shortcut *s1_0 = new tk::dnn::Shortcut(&net, last); + } + tk::dnn::Activation *layer1_0_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + last = layer1_0_relu; + } + + tk::dnn::DeformConv2d *layer0_deform1 = new tk::dnn::DeformConv2d(&net, 256, 1, 3, 3, 1, 1, 1, 1, deform1_bin, d_conv1_bin, true); + tk::dnn::Activation *layer0_deform1_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d *layer0_deconv1 = new tk::dnn::DeConv2d(&net, 256, 4, 4, 2, 2, 1, 1, deconv1_bin, true); + tk::dnn::Activation *layer0_deconv1_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::DeformConv2d *layer1_deform1 = new tk::dnn::DeformConv2d(&net, 128, 1, 3, 3, 1, 1, 1, 1, deform2_bin, d_conv2_bin, true); + tk::dnn::Activation *layer1_deform1_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d *layer1_deconv1 = new tk::dnn::DeConv2d(&net, 128, 4, 4, 2, 2, 1, 1, deconv2_bin, true); + tk::dnn::Activation *layer1_deconv1_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::DeformConv2d *layer2_deform1 = new tk::dnn::DeformConv2d(&net, 64, 1, 3, 3, 1, 1, 1, 1, deform3_bin, d_conv3_bin, true); + tk::dnn::Activation *layer2_deform1_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d *layer2_deconv1 = new tk::dnn::DeConv2d(&net, 64, 4, 4, 2, 2, 1, 1, deconv3_bin, true); + tk::dnn::Activation *layer2_deconv1_relu = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_RELU); + + 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::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::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); + + // Load input + dnnType *data; + dnnType *input_h; + readBinaryFile(input_bin, dim.tot(), &input_h, &data); + // printDeviceVector(64, data, true); + + //print network model + net.print(); + + //convert network to tensorRT +// tk::dnn::NetworkRT netRT(&net, "resnet101_cnet.rt"); + + + tk::dnn::dataDim_t dim1 = dim; //input dim + printCenteredTitle(" CUDNN inference ", '=', 30); + { + dim1.print(); + TIMER_START + net.infer(dim1, data); + TIMER_STOP + dim1.print(); + } + + // printDeviceVector(64, cudnn_out, true); + +/* tk::dnn::dataDim_t dim2 = dim; + printCenteredTitle(" TENSORRT inference ", '=', 30); + { + dim2.print(); + TIMER_START + netRT.infer(dim2, data); + TIMER_STOP + dim2.print(); + } + rt_out = (dnnType *)netRT.buffersRT[1]; +*/ + + tk::dnn::Conv2d *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); + + outs[i]->output_dim.print(); + + dnnType *out, *out_h; + int odim = outs[i]->output_dim.tot(); + readBinaryFile(output_bin[i], odim, &out_h, &out); + // std::cout<<"OUTPUT BIN:\n"; + // printDeviceVector(odim, cudnn_out, true); + // std::cout<<"FILE BIN:\n"; + // printDeviceVector(odim, out, true); + + dnnType *cudnn_out, *rt_out; + cudnn_out = outs[i]->dstData; + + std::cout << "CUDNN vs correct"; + checkResult(odim, cudnn_out, out); + + /* std::cout << "TRT vs correct"; + checkResult(odim, rt_out, out); + std::cout << "CUDNN vs TRT "; + checkResult(odim, cudnn_out, rt_out);*/ + } + return 0; +}