diff --git a/CMakeLists.txt b/CMakeLists.txt index eae11be..27d99ea 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -102,6 +102,9 @@ target_link_libraries(test_resnet101_cnet tkDNN) add_executable(test_dla34 tests/dla34/dla34.cpp) target_link_libraries(test_dla34 tkDNN) +add_executable(test_dla34_cnet tests/dla34_cnet/dla34_cnet.cpp) +target_link_libraries(test_dla34_cnet tkDNN) + ################################################################################ diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 7671c07..fc55922 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -85,7 +85,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, bool final = false); + std::string fname_weights, bool batchnorm = false, bool additional_bias = false, bool final = false, bool deConv = false, int groups = 1); virtual ~LayerWgs(); int inputs, outputs; @@ -165,7 +165,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, bool final = false); + std::string fname_weights, bool batchnorm = false, bool deConv = false, bool final = false, int groups = 1); virtual ~Conv2d(); virtual layerType_t getLayerType() { return LAYER_CONV2D; }; @@ -173,6 +173,7 @@ public: int kernelH, kernelW, strideH, strideW, paddingH, paddingW; bool deConv; + int groups; protected: cudnnFilterDescriptor_t filterDesc; @@ -196,8 +197,8 @@ class DeConv2d : public Conv2d { public: DeConv2d( Network *net, int out_ch, int kernelH, int kernelW, int strideH, int strideW, int paddingH, int paddingW, - std::string fname_weights, bool batchnorm = false) : - Conv2d(net, out_ch, kernelH, kernelW, strideH, strideW, paddingH, paddingW, fname_weights, batchnorm, true) {} + std::string fname_weights, bool batchnorm = false, int groups = 1) : + Conv2d(net, out_ch, kernelH, kernelW, strideH, strideW, paddingH, paddingW, fname_weights, batchnorm, true, false, groups) {} virtual ~DeConv2d() {} virtual layerType_t getLayerType() { return LAYER_DECONV2D; }; diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index f984f67..18e907d 100644 --- a/src/Conv2d.cpp +++ b/src/Conv2d.cpp @@ -17,8 +17,6 @@ void Conv2d::initCUDNN(bool back) { idim = output_dim; odim = input_dim; } - //idim.print(); - //odim.print(); checkCUDNN( cudnnCreateFilterDescriptor(&filterDesc) ); checkCUDNN( cudnnCreateConvolutionDescriptor(&convDesc) ); @@ -29,7 +27,7 @@ void Conv2d::initCUDNN(bool back) { net->tensorFormat, net->dataType, idim.n, idim.c, idim.h, idim.w) ); checkCUDNN( cudnnSetFilter4dDescriptor(filterDesc, - net->dataType, net->tensorFormat, odim.c, idim.c, + net->dataType, net->tensorFormat, odim.c, idim.c/groups, kernelH, kernelW) ); checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc, @@ -38,16 +36,20 @@ void Conv2d::initCUDNN(bool back) { 1,1, // upscale CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT) ); + checkCUDNN( cudnnSetConvolutionGroupCount(convDesc, + groups) ); + // check dimension of convolution output dataDim_t tmpdim; checkCUDNN( cudnnGetConvolution2dForwardOutputDim( 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 input: "; idim.print(); std::cout<<"tkdim output: "; odim.print(); std::cout<<"cudnndim: "; tmpdim.print(); - FatalError("Eror conv dimension mismatch"); + FatalError("Error conv dimension mismatch"); } checkCUDNN( cudnnSetTensor4dDescriptor(dstTensor, @@ -119,11 +121,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, bool final) : + std::string fname_weights, bool batchnorm, bool deConv, bool final, int groups) : LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, - fname_weights, batchnorm, false, final) { - + fname_weights, batchnorm, false, final, deConv, groups) { this->kernelH = kernelH; this->kernelW = kernelW; this->strideH = strideH; @@ -131,6 +132,7 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, this->paddingH = paddingH; this->paddingW = paddingW; this->deConv = deConv; + this->groups = groups; if(!deConv) { output_dim.n = input_dim.n; diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp index 8997593..40b6b9e 100644 --- a/src/DeformConv2d.cpp +++ b/src/DeformConv2d.cpp @@ -76,8 +76,8 @@ DeformConv2d::~DeformConv2d() { checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) ); checkCuda( cudaFree(dstData) ); - checkCuda( cudaFreeHost(ones_d1) ); - checkCuda( cudaFreeHost(ones_d2) ); + checkCuda( cudaFree(ones_d1) ); + checkCuda( cudaFree(ones_d2) ); checkCuda( cudaFree(offset) ); checkCuda( cudaFree(mask) ); checkCuda( cudaFree(output_conv) ); diff --git a/src/LayerWgs.cpp b/src/LayerWgs.cpp index a18fd53..41ca038 100644 --- a/src/LayerWgs.cpp +++ b/src/LayerWgs.cpp @@ -8,8 +8,13 @@ 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, bool final) : Layer(net, final) { - + std::string fname_weights, bool batchnorm, bool additional_bias, bool final, bool deConv, int groups) : Layer(net, final) { + + if(deConv) + inputs = inputs/groups; + else + outputs = outputs/groups; + this->inputs = inputs; this->outputs = outputs; this->weights_path = std::string(fname_weights); diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 0f2d621..2b7377d 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -245,6 +245,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { checkNULL(lRTconv); lRTconv->setStride(DimsHW{l->strideH, l->strideW}); lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); + lRTconv->setNbGroups(l->groups); lRT = (ILayer*) lRTconv; } else { IDeconvolutionLayer *lRTconv = networkRT->addDeconvolution(*input, @@ -252,6 +253,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { checkNULL(lRTconv); lRTconv->setStride(DimsHW{l->strideH, l->strideW}); lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); + lRTconv->setNbGroups(l->groups); lRT = (ILayer*) lRTconv; Dims d = lRTconv->getOutput(0)->getDimensions(); diff --git a/tests/dla34_cnet/dla34_cnet.cpp b/tests/dla34_cnet/dla34_cnet.cpp new file mode 100644 index 0000000..bccbc09 --- /dev/null +++ b/tests/dla34_cnet/dla34_cnet.cpp @@ -0,0 +1,527 @@ +#include +#include "tkdnn.h" + +const char *input_bin = "../tests/dla34_cnet/debug/input.bin"; +const char *conv1_bin = "../tests/dla34_cnet/layers/base-base_layer-0.bin"; +const char *conv2_bin = "../tests/dla34_cnet/layers/base-level0-0.bin"; +const char *conv3_bin = "../tests/dla34_cnet/layers/base-level1-0.bin"; +// s - stage, t - tree +const char *s1_t1_conv1_bin = "../tests/dla34_cnet/layers/base-level2-tree1-conv1.bin"; +const char *s1_t1_conv2_bin = "../tests/dla34_cnet/layers/base-level2-tree1-conv2.bin"; +const char *s1_t1_project = "../tests/dla34_cnet/layers/base-level2-project-0.bin"; +const char *s1_t2_conv1_bin = "../tests/dla34_cnet/layers/base-level2-tree2-conv1.bin"; +const char *s1_t2_conv2_bin = "../tests/dla34_cnet/layers/base-level2-tree2-conv2.bin"; +const char *s1_root_conv1_bin = "../tests/dla34_cnet/layers/base-level2-root-conv.bin"; +const char *s2_t1_t1_conv1_bin = "../tests/dla34_cnet/layers/base-level3-tree1-tree1-conv1.bin"; +const char *s2_t1_t1_conv2_bin = "../tests/dla34_cnet/layers/base-level3-tree1-tree1-conv2.bin"; +const char *s2_t1_t1_project = "../tests/dla34_cnet/layers/base-level3-tree1-project-0.bin"; +const char *s2_t1_t2_conv1_bin = "../tests/dla34_cnet/layers/base-level3-tree1-tree2-conv1.bin"; +const char *s2_t1_t2_conv2_bin = "../tests/dla34_cnet/layers/base-level3-tree1-tree2-conv2.bin"; +const char *s2_t1_root_conv1_bin = "../tests/dla34_cnet/layers/base-level3-tree1-root-conv.bin"; +const char *s2_t2_t1_conv1_bin = "../tests/dla34_cnet/layers/base-level3-tree2-tree1-conv1.bin"; +const char *s2_t2_t1_conv2_bin = "../tests/dla34_cnet/layers/base-level3-tree2-tree1-conv2.bin"; +const char *s2_t2_t2_conv1_bin = "../tests/dla34_cnet/layers/base-level3-tree2-tree2-conv1.bin"; +const char *s2_t2_t2_conv2_bin = "../tests/dla34_cnet/layers/base-level3-tree2-tree2-conv2.bin"; +const char *s2_t2_root_conv1_bin = "../tests/dla34_cnet/layers/base-level3-tree2-root-conv.bin"; +const char *s3_t1_t1_conv1_bin = "../tests/dla34_cnet/layers/base-level4-tree1-tree1-conv1.bin"; +const char *s3_t1_t1_conv2_bin = "../tests/dla34_cnet/layers/base-level4-tree1-tree1-conv2.bin"; +const char *s3_t1_t1_project = "../tests/dla34_cnet/layers/base-level4-tree1-project-0.bin"; +const char *s3_t1_t2_conv1_bin = "../tests/dla34_cnet/layers/base-level4-tree1-tree2-conv1.bin"; +const char *s3_t1_t2_conv2_bin = "../tests/dla34_cnet/layers/base-level4-tree1-tree2-conv2.bin"; +const char *s3_t1_root_conv1_bin = "../tests/dla34_cnet/layers/base-level4-tree1-root-conv.bin"; +const char *s3_t2_t1_conv1_bin = "../tests/dla34_cnet/layers/base-level4-tree2-tree1-conv1.bin"; +const char *s3_t2_t1_conv2_bin = "../tests/dla34_cnet/layers/base-level4-tree2-tree1-conv2.bin"; +const char *s3_t2_t2_conv1_bin = "../tests/dla34_cnet/layers/base-level4-tree2-tree2-conv1.bin"; +const char *s3_t2_t2_conv2_bin = "../tests/dla34_cnet/layers/base-level4-tree2-tree2-conv2.bin"; +const char *s3_t2_root_conv1_bin = "../tests/dla34_cnet/layers/base-level4-tree2-root-conv.bin"; +const char *s4_t1_conv1_bin = "../tests/dla34_cnet/layers/base-level5-tree1-conv1.bin"; +const char *s4_t1_conv2_bin = "../tests/dla34_cnet/layers/base-level5-tree1-conv2.bin"; +const char *s4_t1_project = "../tests/dla34_cnet/layers/base-level5-project-0.bin"; +const char *s4_t2_conv1_bin = "../tests/dla34_cnet/layers/base-level5-tree2-conv1.bin"; +const char *s4_t2_conv2_bin = "../tests/dla34_cnet/layers/base-level5-tree2-conv2.bin"; +const char *s4_root_conv1_bin = "../tests/dla34_cnet/layers/base-level5-root-conv.bin"; + +//final +// const char *fc_bin = "../tests/dla34_cnet/layers/output.bin"; + +const char *ida_0_p_1_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_0-proj_1-conv.bin"; +const char *ida_0_p_1_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_0-proj_1-conv-conv_offset_mask.bin"; +const char *ida_0_up_1_deconv_bin = "../tests/dla34_cnet/layers/dla_up-ida_0-up_1.bin"; +const char *ida_0_n_1_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_0-node_1-conv.bin"; +const char *ida_0_n_1_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_0-node_1-conv-conv_offset_mask.bin"; + +const char *ida_1_p_1_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-proj_1-conv.bin"; +const char *ida_1_p_1_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-proj_1-conv-conv_offset_mask.bin"; +const char *ida_1_up_1_deconv_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-up_1.bin"; +const char *ida_1_n_1_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-node_1-conv.bin"; +const char *ida_1_n_1_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-node_1-conv-conv_offset_mask.bin"; +const char *ida_1_p_2_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-proj_2-conv.bin"; +const char *ida_1_p_2_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-proj_2-conv-conv_offset_mask.bin"; +const char *ida_1_up_2_deconv_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-up_2.bin"; +const char *ida_1_n_2_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-node_2-conv.bin"; +const char *ida_1_n_2_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_1-node_2-conv-conv_offset_mask.bin"; + +const char *ida_2_p_1_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-proj_1-conv.bin"; +const char *ida_2_p_1_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-proj_1-conv-conv_offset_mask.bin"; +const char *ida_2_up_1_deconv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-up_1.bin"; +const char *ida_2_n_1_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-node_1-conv.bin"; +const char *ida_2_n_1_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-node_1-conv-conv_offset_mask.bin"; +const char *ida_2_p_2_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-proj_2-conv.bin"; +const char *ida_2_p_2_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-proj_2-conv-conv_offset_mask.bin"; +const char *ida_2_up_2_deconv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-up_2.bin"; +const char *ida_2_n_2_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-node_2-conv.bin"; +const char *ida_2_n_2_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-node_2-conv-conv_offset_mask.bin"; +const char *ida_2_p_3_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-proj_3-conv.bin"; +const char *ida_2_p_3_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-proj_3-conv-conv_offset_mask.bin"; +const char *ida_2_up_3_deconv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-up_3.bin"; +const char *ida_2_n_3_dcn_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-node_3-conv.bin"; +const char *ida_2_n_3_conv_bin = "../tests/dla34_cnet/layers/dla_up-ida_2-node_3-conv-conv_offset_mask.bin"; + +const char *ida_up_p_1_dcn_bin = "../tests/dla34_cnet/layers/ida_up-proj_1-conv.bin"; +const char *ida_up_p_1_conv_bin = "../tests/dla34_cnet/layers/ida_up-proj_1-conv-conv_offset_mask.bin"; +const char *ida_up_up_1_deconv_bin = "../tests/dla34_cnet/layers/ida_up-up_1.bin"; +const char *ida_up_n_1_dcn_bin = "../tests/dla34_cnet/layers/ida_up-node_1-conv.bin"; +const char *ida_up_n_1_conv_bin = "../tests/dla34_cnet/layers/ida_up-node_1-conv-conv_offset_mask.bin"; +const char *ida_up_p_2_dcn_bin = "../tests/dla34_cnet/layers/ida_up-proj_2-conv.bin"; +const char *ida_up_p_2_conv_bin = "../tests/dla34_cnet/layers/ida_up-proj_2-conv-conv_offset_mask.bin"; +const char *ida_up_up_2_deconv_bin = "../tests/dla34_cnet/layers/ida_up-up_2.bin"; +const char *ida_up_n_2_dcn_bin = "../tests/dla34_cnet/layers/ida_up-node_2-conv.bin"; +const char *ida_up_n_2_conv_bin = "../tests/dla34_cnet/layers/ida_up-node_2-conv-conv_offset_mask.bin"; + +const char *hm_conv1_bin = "../tests/dla34_cnet/layers/hm-0.bin"; +const char *hm_conv2_bin = "../tests/dla34_cnet/layers/hm-2.bin"; +const char *wh_conv1_bin = "../tests/dla34_cnet/layers/wh-0.bin"; +const char *wh_conv2_bin = "../tests/dla34_cnet/layers/wh-2.bin"; +const char *reg_conv1_bin = "../tests/dla34_cnet/layers/reg-0.bin"; +const char *reg_conv2_bin = "../tests/dla34_cnet/layers/reg-2.bin"; + +const char *output_bin[]={ +"../tests/dla34_cnet/debug/hm.bin", +"../tests/dla34_cnet/debug/wh.bin", +"../tests/dla34_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::Layer *last1, *last2, *last3, *last4; + tk::dnn::Layer *base1, *base2, *base3, *base4, *base5, *base6, *ida1, *ida2_1, *ida2_2, *ida3_1, *ida3_2, *ida3_3, *idaup_1, *idaup_2; + + tk::dnn::Conv2d conv1(&net, 16, 7, 7, 1, 1, 3, 3, conv1_bin, true); + tk::dnn::Activation relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d conv2(&net, 16, 3, 3, 1, 1, 1, 1, conv2_bin, true); + tk::dnn::Activation relu2(&net, CUDNN_ACTIVATION_RELU); + base1 = &relu2; + + tk::dnn::Conv2d conv3(&net, 32, 3, 3, 2, 2, 1, 1, conv3_bin, true); + tk::dnn::Activation relu3(&net, CUDNN_ACTIVATION_RELU); + base2 = &relu3; + + // level 2 + // tree 1 + tk::dnn::Conv2d s1_t1_conv1(&net, 64, 3, 3, 2, 2, 1, 1, s1_t1_conv1_bin, true); + tk::dnn::Activation s1_t1_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s1_t1_conv2(&net, 64, 3, 3, 1, 1, 1, 1, s1_t1_conv2_bin, true); + last2 = &s1_t1_conv2; + + // get the basicblock input and apply maxpool conv2d and relu + tk::dnn::Layer *route_s1_t1_layers[1] = { base2 }; + tk::dnn::Route route_s1_t1(&net, route_s1_t1_layers, 1); + // downsample + tk::dnn::Pooling s1_t1_maxpool1(&net, 2, 2, 2, 2, 0, 0, tk::dnn::POOLING_MAX); + // project + tk::dnn::Conv2d s1_t1_residual1_conv1(&net, 64, 1, 1, 1, 1, 0, 0, s1_t1_project, true); + + tk::dnn::Shortcut s1_t1_s1(&net, last2); + tk::dnn::Activation s1_t1_relu(&net, CUDNN_ACTIVATION_RELU); + + last1 = &s1_t1_relu; + // tree 2 + tk::dnn::Conv2d s1_t2_conv1(&net, 64, 3, 3, 1, 1, 1, 1, s1_t2_conv1_bin, true); + tk::dnn::Activation s1_t2_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s1_t2_conv2(&net, 64, 3, 3, 1, 1, 1, 1, s1_t2_conv2_bin, true); + + tk::dnn::Shortcut s1_t2_s1(&net, last1); + tk::dnn::Activation s1_t2_relu(&net, CUDNN_ACTIVATION_RELU); + last2 = &s1_t2_relu; + + // root + // join last1 and net in single input 128, 56, 56 + tk::dnn::Layer *route_s1_root_layers[2] = { last2, last1 }; + tk::dnn::Route route_s1_root(&net, route_s1_root_layers, 2); + tk::dnn::Conv2d s1_root_conv1(&net, 64, 1, 1, 1, 1, 0, 0, s1_root_conv1_bin, true); + tk::dnn::Activation s1_root_relu(&net, CUDNN_ACTIVATION_RELU); + + base3 = &s1_root_relu; + + // level 3 + // tree 1 + // tree 1 + tk::dnn::Conv2d s2_t1_t1_conv1(&net, 128, 3, 3, 2, 2, 1, 1, s2_t1_t1_conv1_bin, true); + tk::dnn::Activation s2_t1_t1_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s2_t1_t1_conv2(&net, 128, 3, 3, 1, 1, 1, 1, s2_t1_t1_conv2_bin, true); + last2 = &s2_t1_t1_conv2; + + // get the basicblock input and apply maxpool conv2d and relu + tk::dnn::Layer *route_s2_t1_t1_layers[1] = { base3 }; + tk::dnn::Route route_s2_t1_t1(&net, route_s2_t1_t1_layers, 1); + // downsample + tk::dnn::Pooling s2_t1_t1_maxpool1(&net, 2, 2, 2, 2, 0, 0, tk::dnn::POOLING_MAX); + last4 = &s2_t1_t1_maxpool1; + // project + tk::dnn::Conv2d s2_t1_t1_residual1_conv1(&net, 128, 1, 1, 1, 1, 0, 0, s2_t1_t1_project, true); + + tk::dnn::Shortcut s2_t1_t1_s1(&net, last2); + tk::dnn::Activation s2_t1_t1_relu(&net, CUDNN_ACTIVATION_RELU); + + last1 = &s2_t1_t1_relu; + + // tree 2 + tk::dnn::Conv2d s2_t1_t2_conv1(&net, 128, 3, 3, 1, 1, 1, 1, s2_t1_t2_conv1_bin, true); + tk::dnn::Activation s2_t1_t2_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s2_t1_t2_conv2(&net, 128, 3, 3, 1, 1, 1, 1, s2_t1_t2_conv2_bin, true); + + tk::dnn::Shortcut s2_t1_t2_s1(&net, last1); + tk::dnn::Activation s2_t1_t2_relu(&net, CUDNN_ACTIVATION_RELU); + last2 = &s2_t1_t2_relu; + + // root + // join last1 and net in single input 128, 56, 56 + tk::dnn::Layer *route_s2_t1_root_layers[2] = { last2, last1 }; + tk::dnn::Route route_s2_t1_root(&net, route_s2_t1_root_layers, 2); + tk::dnn::Conv2d s2_t1_root_conv1(&net, 128, 1, 1, 1, 1, 0, 0, s2_t1_root_conv1_bin, true); + tk::dnn::Activation s2_t1_root_relu(&net, CUDNN_ACTIVATION_RELU); + + last1 = &s2_t1_root_relu; + last3 = &s2_t1_root_relu; + // tree 2 + // tree 1 + tk::dnn::Conv2d s2_t2_t1_conv1(&net, 128, 3, 3, 1, 1, 1, 1, s2_t2_t1_conv1_bin, true); + tk::dnn::Activation s2_t2_t1_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s2_t2_t1_conv2(&net, 128, 3, 3, 1, 1, 1, 1, s2_t2_t1_conv2_bin, true); + tk::dnn::Shortcut s2_t2_t1_s1(&net, last1); + tk::dnn::Activation s2_t2_t1_relu(&net, CUDNN_ACTIVATION_RELU); + + last1 = &s2_t2_t1_relu; + + // tree 2 + tk::dnn::Conv2d s2_t2_t2_conv1(&net, 128, 3, 3, 1, 1, 1, 1, s2_t2_t2_conv1_bin, true); + tk::dnn::Activation s2_t2_t2_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s2_t2_t2_conv2(&net, 128, 3, 3, 1, 1, 1, 1, s2_t2_t2_conv2_bin, true); + + tk::dnn::Shortcut s2_t2_t2_s1(&net, last1); + tk::dnn::Activation s2_t2_t2_relu(&net, CUDNN_ACTIVATION_RELU); + last2 = &s2_t2_t2_relu; + + // root + // join last1 and net in single input 128, 56, 56 + tk::dnn::Layer *route_s2_t2_root_layers[4] = { last2, last1, last4, last3}; + tk::dnn::Route route_s2_t2_root(&net, route_s2_t2_root_layers, 4); + tk::dnn::Conv2d s2_t2_root_conv1(&net, 128, 1, 1, 1, 1, 0, 0, s2_t2_root_conv1_bin, true); + tk::dnn::Activation s2_t2_root_relu(&net, CUDNN_ACTIVATION_RELU); + + base4 = &s2_t2_root_relu; + + // level 4 + // tree 1 + // tree 1 + tk::dnn::Conv2d s3_t1_t1_conv1(&net, 256, 3, 3, 2, 2, 1, 1, s3_t1_t1_conv1_bin, true); + tk::dnn::Activation s3_t1_t1_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s3_t1_t1_conv2(&net, 256, 3, 3, 1, 1, 1, 1, s3_t1_t1_conv2_bin, true); + last2 = &s3_t1_t1_conv2; + + // get the basicblock input and apply maxpool conv2d and relu + tk::dnn::Layer *route_s3_t1_t1_layers[1] = { base4 }; + tk::dnn::Route route_s3_t1_t1(&net, route_s3_t1_t1_layers, 1); + // downsample + tk::dnn::Pooling s3_t1_t1_maxpool1(&net, 2, 2, 2, 2, 0, 0, tk::dnn::POOLING_MAX); + last4 = &s3_t1_t1_maxpool1; + // project + tk::dnn::Conv2d s3_t1_t1_residual1_conv1(&net, 256, 1, 1, 1, 1, 0, 0, s3_t1_t1_project, true); + + tk::dnn::Shortcut s3_t1_t1_s1(&net, last2); + tk::dnn::Activation s3_t1_t1_relu(&net, CUDNN_ACTIVATION_RELU); + + last1 = &s3_t1_t1_relu; + + // tree 2 + tk::dnn::Conv2d s3_t1_t2_conv1(&net, 256, 3, 3, 1, 1, 1, 1, s3_t1_t2_conv1_bin, true); + tk::dnn::Activation s3_t1_t2_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s3_t1_t2_conv2(&net, 256, 3, 3, 1, 1, 1, 1, s3_t1_t2_conv2_bin, true); + + tk::dnn::Shortcut s3_t1_t2_s1(&net, last1); + tk::dnn::Activation s3_t1_t2_relu(&net, CUDNN_ACTIVATION_RELU); + last2 = &s3_t1_t2_relu; + + // root + // join last1 and net in single input 256, 56, 56 + tk::dnn::Layer *route_s3_t1_root_layers[2] = { last2, last1 }; + tk::dnn::Route route_s3_t1_root(&net, route_s3_t1_root_layers, 2); + tk::dnn::Conv2d s3_t1_root_conv1(&net, 256, 1, 1, 1, 1, 0, 0, s3_t1_root_conv1_bin, true); + tk::dnn::Activation s3_t1_root_relu(&net, CUDNN_ACTIVATION_RELU); + + last1 = &s3_t1_root_relu; + last3 = &s3_t1_root_relu; + // tree 2 + // tree 1 + tk::dnn::Conv2d s3_t2_t1_conv1(&net, 256, 3, 3, 1, 1, 1, 1, s3_t2_t1_conv1_bin, true); + tk::dnn::Activation s3_t2_t1_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s3_t2_t1_conv2(&net, 256, 3, 3, 1, 1, 1, 1, s3_t2_t1_conv2_bin, true); + tk::dnn::Shortcut s3_t2_t1_s1(&net, last1); + tk::dnn::Activation s3_t2_t1_relu(&net, CUDNN_ACTIVATION_RELU); + + last1 = &s3_t2_t1_relu; + + // tree 2 + tk::dnn::Conv2d s3_t2_t2_conv1(&net, 256, 3, 3, 1, 1, 1, 1, s3_t2_t2_conv1_bin, true); + tk::dnn::Activation s3_t2_t2_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s3_t2_t2_conv2(&net, 256, 3, 3, 1, 1, 1, 1, s3_t2_t2_conv2_bin, true); + + tk::dnn::Shortcut s3_t2_t2_s1(&net, last1); + tk::dnn::Activation s3_t2_t2_relu(&net, CUDNN_ACTIVATION_RELU); + last2 = &s3_t2_t2_relu; + + // root + // join last1 and net in single input 256, 56, 56 + tk::dnn::Layer *route_s3_t2_root_layers[4] = { last2, last1, last4, last3}; + tk::dnn::Route route_s3_t2_root(&net, route_s3_t2_root_layers, 4); + tk::dnn::Conv2d s3_t2_root_conv1(&net, 256, 1, 1, 1, 1, 0, 0, s3_t2_root_conv1_bin, true); + tk::dnn::Activation s3_t2_root_relu(&net, CUDNN_ACTIVATION_RELU); + + base5 = &s3_t2_root_relu; + + // level 5 + // tree 1 + tk::dnn::Conv2d s4_t1_conv1(&net, 512, 3, 3, 2, 2, 1, 1, s4_t1_conv1_bin, true); + tk::dnn::Activation s4_t1_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s4_t1_conv2(&net, 512, 3, 3, 1, 1, 1, 1, s4_t1_conv2_bin, true); + last2 = &s4_t1_conv2; + + // get the basicblock input and apply maxpool conv2d and relu + tk::dnn::Layer *route_s4_t1_layers[1] = { base5 }; + tk::dnn::Route route_s4_t1(&net, route_s4_t1_layers, 1); + // downsample + tk::dnn::Pooling s4_t1_maxpool1(&net, 2, 2, 2, 2, 0, 0, tk::dnn::POOLING_MAX); + last4 = &s4_t1_maxpool1; + // project + tk::dnn::Conv2d s4_t1_residual1_conv1(&net, 512, 1, 1, 1, 1, 0, 0, s4_t1_project, true); + + tk::dnn::Shortcut s4_t1_s1(&net, last2); + tk::dnn::Activation s4_t1_relu(&net, CUDNN_ACTIVATION_RELU); + + last1 = &s4_t1_relu; + + // tree 2 + tk::dnn::Conv2d s4_t2_conv1(&net, 512, 3, 3, 1, 1, 1, 1, s4_t2_conv1_bin, true); + tk::dnn::Activation s4_t2_relu1(&net, CUDNN_ACTIVATION_RELU); + + tk::dnn::Conv2d s4_t2_conv2(&net, 512, 3, 3, 1, 1, 1, 1, s4_t2_conv2_bin, true); + + tk::dnn::Shortcut s4_t2_s1(&net, last1); + tk::dnn::Activation s4_t2_relu(&net, CUDNN_ACTIVATION_RELU); + last2 = &s4_t2_relu; + + // root + // join last1 and net in single input 128, 56, 56 + tk::dnn::Layer *route_s4_root_layers[3] = { last2, last1, last4 }; + tk::dnn::Route route_s4_root(&net, route_s4_root_layers, 3); + tk::dnn::Conv2d s4_root_conv1(&net, 512, 1, 1, 1, 1, 0, 0, s4_root_conv1_bin, true); + tk::dnn::Activation s4_root_relu(&net, CUDNN_ACTIVATION_RELU); + + base6 = &s4_root_relu; + + //final + // tk::dnn::Pooling avgpool(&net, 7, 7, 7, 7, 0, 0, tk::dnn::POOLING_AVERAGE); + // tk::dnn::Dense fc(&net, 1000, fc_bin); + + //ida 0 + tk::dnn::DeformConv2d ida_0_p_1_dcn(&net, 256, 1, 3, 3, 1, 1, 1, 1, ida_0_p_1_dcn_bin, ida_0_p_1_conv_bin, true); + tk::dnn::Activation ida_0_p_1_relu(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d ida_0_up_1_deconv(&net, 256, 4, 4, 2, 2, 1, 1, ida_0_up_1_deconv_bin, false, 256); + tk::dnn::Shortcut ida_0_shortcut(&net, base5); + tk::dnn::DeformConv2d ida_0_n_1_dcn(&net, 256, 1, 3, 3, 1, 1, 1, 1, ida_0_n_1_dcn_bin, ida_0_n_1_conv_bin, true); + tk::dnn::Activation ida_0_n_1_relu(&net, CUDNN_ACTIVATION_RELU); + ida1 = &ida_0_n_1_relu; + + //ida1-1 + tk::dnn::Layer *route_ida1_layers_1[1] = { base5 }; + tk::dnn::Route route_ida1_1(&net, route_ida1_layers_1, 1); + + tk::dnn::DeformConv2d ida_1_p_1_dcn(&net, 128, 1, 3, 3, 1, 1, 1, 1, ida_1_p_1_dcn_bin, ida_1_p_1_conv_bin, true); + tk::dnn::Activation ida_1_p_1_relu(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d ida_1_up_1_deconv(&net, 128, 4, 4, 2, 2, 1, 1, ida_1_up_1_deconv_bin, false, 128); + tk::dnn::Shortcut ida_1_shortcut1(&net, base4); + tk::dnn::DeformConv2d ida_1_n_1_dcn(&net, 128, 1, 3, 3, 1, 1, 1, 1, ida_1_n_1_dcn_bin, ida_1_n_1_conv_bin, true); + tk::dnn::Activation ida_1_n_1_relu(&net, CUDNN_ACTIVATION_RELU); + ida2_1 = &ida_1_n_1_relu; + + //ida1-2 + tk::dnn::Layer *route_ida1_layers_2[1] = { ida1 }; + tk::dnn::Route route_ida1_2(&net, route_ida1_layers_2, 1); + + tk::dnn::DeformConv2d ida_1_p_2_dcn(&net, 128, 1, 3, 3, 1, 1, 1, 1, ida_1_p_2_dcn_bin, ida_1_p_2_conv_bin, true); + tk::dnn::Activation ida_1_p_2_relu(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d ida_1_up_2_deconv(&net, 128, 4, 4, 2, 2, 1, 1, ida_1_up_2_deconv_bin, false, 128); + tk::dnn::Shortcut ida_1_shortcut2(&net, ida2_1); + tk::dnn::DeformConv2d ida_1_n_2_dcn(&net, 128, 1, 3, 3, 1, 1, 1, 1, ida_1_n_2_dcn_bin, ida_1_n_2_conv_bin, true); + tk::dnn::Activation ida_1_n_2_relu(&net, CUDNN_ACTIVATION_RELU); + ida2_2 = &ida_1_n_2_relu; + + //ida2-1 + tk::dnn::Layer *route_ida2_layers_1[1] = { base4 }; + tk::dnn::Route route_ida2_1(&net, route_ida2_layers_1, 1); + + tk::dnn::DeformConv2d ida_2_p_1_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_2_p_1_dcn_bin, ida_2_p_1_conv_bin, true); + tk::dnn::Activation ida_2_p_1_relu(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d ida_2_up_1_deconv(&net, 64, 4, 4, 2, 2, 1, 1, ida_2_up_1_deconv_bin, false, 64); + tk::dnn::Shortcut ida_2_shortcut1(&net, base3); + tk::dnn::DeformConv2d ida_2_n_1_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_2_n_1_dcn_bin, ida_2_n_1_conv_bin, true); + tk::dnn::Activation ida_2_n_1_relu(&net, CUDNN_ACTIVATION_RELU); + ida3_1 = &ida_2_n_1_relu; + + //ida2-2 + tk::dnn::Layer *route_ida2_layers_2[1] = { ida2_1 }; + tk::dnn::Route route_ida2_2(&net, route_ida2_layers_2, 1); + + tk::dnn::DeformConv2d ida_2_p_2_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_2_p_2_dcn_bin, ida_2_p_2_conv_bin, true); + tk::dnn::Activation ida_2_p_2_relu(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d ida_2_up_2_deconv(&net, 64, 4, 4, 2, 2, 1, 1, ida_2_up_2_deconv_bin, false, 64); + tk::dnn::Shortcut ida_2_shortcut2(&net, ida3_1); + tk::dnn::DeformConv2d ida_2_n_2_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_2_n_2_dcn_bin, ida_2_n_2_conv_bin, true); + tk::dnn::Activation ida_2_n_2_relu(&net, CUDNN_ACTIVATION_RELU); + ida3_2 = &ida_2_n_2_relu; + + //ida2-3 + tk::dnn::Layer *route_ida2_layers_3[1] = { ida2_2 }; + tk::dnn::Route route_ida2_3(&net, route_ida2_layers_3, 1); + + tk::dnn::DeformConv2d ida_2_p_3_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_2_p_3_dcn_bin, ida_2_p_3_conv_bin, true); + tk::dnn::Activation ida_2_p_3_relu(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d ida_2_up_3_deconv(&net, 64, 4, 4, 2, 2, 1, 1, ida_2_up_3_deconv_bin, false, 64); + tk::dnn::Shortcut ida_2_shortcut3(&net, ida3_2); + tk::dnn::DeformConv2d ida_2_n_3_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_2_n_3_dcn_bin, ida_2_n_3_conv_bin, true); + tk::dnn::Activation ida_2_n_3_relu(&net, CUDNN_ACTIVATION_RELU); + ida3_3 = &ida_2_n_3_relu; + + //idaup-1 + tk::dnn::Layer *route_idaup_layers_1[1] = { ida2_2 }; + tk::dnn::Route route_idaup_1(&net, route_idaup_layers_1, 1); + + tk::dnn::DeformConv2d idaup_p_1_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_up_p_1_dcn_bin, ida_up_p_1_conv_bin, true); + tk::dnn::Activation idaup_p_1_relu(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d idaup_up_1_deconv(&net, 64, 4, 4, 2, 2, 1, 1, ida_up_up_1_deconv_bin, false, 64); + tk::dnn::Shortcut idaup_shortcut1(&net, ida3_3); + tk::dnn::DeformConv2d idaup_n_1_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_up_n_1_dcn_bin, ida_up_n_1_conv_bin, true); + tk::dnn::Activation idaup_n_1_relu(&net, CUDNN_ACTIVATION_RELU); + idaup_1 = &idaup_n_1_relu; + + //idaup-2 + tk::dnn::Layer *route_idaup_layers_2[1] = { ida1 }; + tk::dnn::Route route_idaup_2(&net, route_idaup_layers_2, 1); + + tk::dnn::DeformConv2d idaup_p_2_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_up_p_2_dcn_bin, ida_up_p_2_conv_bin, true); + tk::dnn::Activation idaup_p_2_relu(&net, CUDNN_ACTIVATION_RELU); + tk::dnn::DeConv2d idaup_up_2_deconv(&net, 64, 8, 8, 4, 4, 2, 2, ida_up_up_2_deconv_bin, false, 64); + tk::dnn::Shortcut idaup_shortcut2(&net, idaup_1); + tk::dnn::DeformConv2d idaup_n_2_dcn(&net, 64, 1, 3, 3, 1, 1, 1, 1, ida_up_n_2_dcn_bin, ida_up_n_2_conv_bin, true); + tk::dnn::Activation idaup_n_2_relu(&net, CUDNN_ACTIVATION_RELU); + idaup_2 = &idaup_n_2_relu; + + tk::dnn::Layer *route_1_0_layers[1] = { idaup_2 }; + + // hm + tk::dnn::Conv2d *hm_conv1 = new tk::dnn::Conv2d(&net, 256, 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, false, true); + int kernel = 3; + int pad = (kernel - 1)/2; + tk::dnn::Activation *hm_sig = new tk::dnn::Activation(&net, CUDNN_ACTIVATION_SIGMOID); + tk::dnn::Pooling *hmax = new tk::dnn::Pooling(&net, kernel, kernel, 1, 1, pad, pad, tk::dnn::POOLING_MAX, true); + + // // wh + 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, 256, 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, false, true); + + // // reg + 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, 256, 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, false, true); + + + // 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, "dla34_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(); + } + + tk::dnn::dataDim_t dim2 = dim; + printCenteredTitle(" TENSORRT inference ", '=', 30); + { + dim2.print(); + TIMER_START + netRT.infer(dim2, data); + TIMER_STOP + dim2.print(); + } + + tk::dnn::Layer *outs[3] = { hm, wh, reg }; + int out_count = 1; + 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); + + dnnType *cudnn_out, *rt_out; + cudnn_out = outs[i]->dstData; + rt_out = (dnnType *)netRT.buffersRT[i+out_count]; + // there is the maxpool. It isn't an output but it is necessary for the process section + if(i==0) + out_count ++; + + 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; +}