From f6527f51e3b0864553ea864944bb8ca9e7742da9 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Wed, 11 Mar 2020 18:37:26 +0100 Subject: [PATCH] Add csresnext50-panet-spp test. Works on CUDNN. Does not work with TensorRT Signed-off-by: Francesco Gatti --- CMakeLists.txt | 3 + include/tkDNN/Layer.h | 5 +- include/tkDNN/Yolo3Detection.h | 2 +- include/tkDNN/kernels.h | 50 +- src/CenternetDetection.cpp | 6 +- src/Pooling.cpp | 24 +- src/Shortcut.cpp | 2 +- src/kernels/pooling.cu | 52 ++ .../csresnext50-panet-spp.cpp | 560 ++++++++++++++++++ 9 files changed, 666 insertions(+), 38 deletions(-) create mode 100644 src/kernels/pooling.cu create mode 100644 tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 139bd9c..6ee23c2 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -100,6 +100,9 @@ target_link_libraries(test_mobilenetv2ssd tkDNN) add_executable(test_resnet101 tests/resnet101/resnet101.cpp) target_link_libraries(test_resnet101 tkDNN) +add_executable(test_csresnext50-panet-spp tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp) +target_link_libraries(test_csresnext50-panet-spp 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 695b8d8..1947309 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -321,12 +321,13 @@ public: int winH, winW; int strideH, strideW; int paddingH, paddingW; + bool test; tkdnnPoolingMode_t pool_mode; Pooling(Network *net, int winH, int winW, int strideH, int strideW, int paddingH = 0, int paddingW = 0, - tkdnnPoolingMode_t pool_mode = POOLING_MAX, bool final = false); + tkdnnPoolingMode_t pool_mode = POOLING_MAX, bool final = false, bool test=false); virtual ~Pooling(); virtual layerType_t getLayerType() { return LAYER_POOLING; }; @@ -473,7 +474,7 @@ public: dnnType *predictions; - static const int MAX_DETECTIONS = 1024; + static const int MAX_DETECTIONS = 4096; static Yolo::detection *allocateDetections(int nboxes, int classes); static void mergeDetections(Yolo::detection *dets, int ndets, int classes); }; diff --git a/include/tkDNN/Yolo3Detection.h b/include/tkDNN/Yolo3Detection.h index 51994e1..80d179e 100644 --- a/include/tkDNN/Yolo3Detection.h +++ b/include/tkDNN/Yolo3Detection.h @@ -37,7 +37,7 @@ class Yolo3Detection { int classes = 0; int num = 0; int n_masks = 0; - float thresh = 0.05; + float thresh = 0.3; cv::Scalar colors[256]; // this is filled with results diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index d3a5d0d..13fda1f 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -3,33 +3,34 @@ #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 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)); +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 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)); -void fill(dnnType* data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0)); +void fill(dnnType *data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0)); -void resizeForward( dnnType* srcData, dnnType* dstData, int n, int i_c, int i_h, int i_w, - int o_c, int o_h, int o_w, cudaStream_t stream = cudaStream_t(0)); +void resizeForward(dnnType *srcData, dnnType *dstData, int n, int i_c, int i_h, int i_w, + int o_c, int o_h, int o_w, cudaStream_t stream = cudaStream_t(0)); -void reorgForward( dnnType* srcData, dnnType* dstData, - int n, int c, int h, int w, int stride, cudaStream_t stream = cudaStream_t(0)); -void softmaxForward(float *input, int n, int batch, int batch_offset, +void reorgForward(dnnType *srcData, dnnType *dstData, + int n, int c, int h, int w, int stride, cudaStream_t stream = cudaStream_t(0)); + +void MaxPoolingForward(dnnType *srcData, dnnType *dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream = cudaStream_t(0)); + +void softmaxForward(float *input, int n, int batch, int batch_offset, int groups, int group_offset, int stride, float temp, float *output, cudaStream_t stream = cudaStream_t(0)); - -void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int w1, int s1, - int n2, int c2, int h2, int w2, int s2, +void shortcutForward(dnnType *srcData, dnnType *dstData, int n1, int c1, int h1, int w1, int s1, + int n2, int c2, int h2, int w2, int s2, cudaStream_t stream = cudaStream_t(0)); -void upsampleForward(dnnType* srcData, dnnType* dstData, - int n, int c, int h, int w, int s, int forward, float scale, +void upsampleForward(dnnType *srcData, dnnType *dstData, + int n, int c, int h, int w, int s, int forward, float scale, cudaStream_t stream = cudaStream_t(0)); -void float2half(float* srcData, __half* dstData, int size, const cudaStream_t stream = cudaStream_t(0)); - +void float2half(float *srcData, __half *dstData, int size, const cudaStream_t stream = cudaStream_t(0)); // void modulated_deformable_im2col_cuda(cudaStream_t stream, // const float *data_im, const float *data_offset, const float *data_mask, @@ -39,13 +40,12 @@ void float2half(float* srcData, __half* dstData, int size, const cudaStream_t st // const int dilation_h, const int dilation_w, // const int deformable_group, float *data_col); void modulated_deformable_im2col_cuda(cudaStream_t stream, - const float *data_im, const float *data_offset, const float *data_mask, - const int batch_size, const int channels, const int height_im, const int width_im, - const int height_col, const int width_col, - const int deformable_group, float *data_col); + const float *data_im, const float *data_offset, const float *data_mask, + const int batch_size, const int channels, const int height_im, const int width_im, + const int height_col, const int width_col, + const int deformable_group, float *data_col); - -void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, +void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, float *input, float *weight, float *bias, float *ones, float *offset, float *mask, @@ -54,7 +54,7 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, 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 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)); diff --git a/src/CenternetDetection.cpp b/src/CenternetDetection.cpp index 512d0e0..df44310 100644 --- a/src/CenternetDetection.cpp +++ b/src/CenternetDetection.cpp @@ -229,7 +229,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) { sz_old = sz; cv::cuda::GpuMat im_Orig; im_Orig = cv::cuda::GpuMat(imageORIG); - cv::cuda::resize (im_Orig, imageF1_d, cv::Size(new_width, new_height)); + // cv::cuda::resize (im_Orig, imageF1_d, cv::Size(new_width, new_height)); checkCuda( cudaDeviceSynchronize() ); sz = imageF1_d.size(); @@ -238,7 +238,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) { std::cout << " TIME resize: " << std::chrono::duration_cast(end_t - step_t).count() << " us" << std::endl; step_t = end_t; - cv::cuda::warpAffine(imageF1_d, imageF2_d, trans, cv::Size(inp_width, inp_height), cv::INTER_LINEAR ); + // cv::cuda::warpAffine(imageF1_d, imageF2_d, trans, cv::Size(inp_width, inp_height), cv::INTER_LINEAR ); checkCuda( cudaDeviceSynchronize() ); end_t = std::chrono::steady_clock::now(); std::cout << " TIME warpAffine: " << std::chrono::duration_cast(end_t - step_t).count() << " us" << std::endl; @@ -251,7 +251,7 @@ void CenternetDetection::update(cv::Mat &imageORIG) { step_t = end_t; dim2 = dim; - cv::cuda::split(imageF1_d,bgr);//split source + // cv::cuda::split(imageF1_d,bgr);//split source end_t = std::chrono::steady_clock::now(); std::cout << " TIME split: " << std::chrono::duration_cast(end_t - step_t).count() << " us" << std::endl; step_t = end_t; diff --git a/src/Pooling.cpp b/src/Pooling.cpp index cc498df..272edca 100644 --- a/src/Pooling.cpp +++ b/src/Pooling.cpp @@ -7,7 +7,7 @@ namespace tk { namespace dnn { Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, int paddingH, int paddingW, - tkdnnPoolingMode_t pool_mode, bool final) : + tkdnnPoolingMode_t pool_mode, bool final, bool test) : Layer(net, final) { this->winH = winH; @@ -17,6 +17,7 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, this->pool_mode = pool_mode; this->paddingH = paddingH; this->paddingW = paddingW; + this->test = test; checkCUDNN( cudnnCreatePoolingDescriptor(&poolingDesc) ); @@ -111,11 +112,22 @@ dnnType* Pooling::infer(dataDim_t &dim, dnnType* srcData) { poolDst = tmpOutputData; } - dnnType alpha = dnnType(1); - dnnType beta = dnnType(0); - checkCUDNN( cudnnPoolingForward(net->cudnnHandle, poolingDesc, - &alpha, srcTensorDesc, poolSrc, - &beta, dstTensorDesc, poolDst) ); + + + if(this->test) + { + MaxPoolingForward(poolSrc, poolDst, dim.n, dim.c, dim.h, dim.w, this->strideH, this->strideW, this->winH, this->winH-1); + } + else + { + dnnType alpha = dnnType(1); + dnnType beta = dnnType(0); + checkCUDNN( cudnnPoolingForward(net->cudnnHandle, poolingDesc, + &alpha, srcTensorDesc, poolSrc, + &beta, dstTensorDesc, poolDst) ); + + } + //update dim dim = output_dim; diff --git a/src/Shortcut.cpp b/src/Shortcut.cpp index bcd2a00..78a2f23 100644 --- a/src/Shortcut.cpp +++ b/src/Shortcut.cpp @@ -10,7 +10,7 @@ Shortcut::Shortcut(Network *net, Layer *backLayer) : Layer(net) { this->backLayer = backLayer; checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) ); - if( backLayer->output_dim.c != input_dim.c || + 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"); diff --git a/src/kernels/pooling.cu b/src/kernels/pooling.cu new file mode 100644 index 0000000..cd80a6d --- /dev/null +++ b/src/kernels/pooling.cu @@ -0,0 +1,52 @@ +#include "kernels.h" + +__global__ void forward_maxpool_layer_kernel(int n, int in_h, int in_w, int in_c, int stride_x, int stride_y, int size, int pad, float *input, float *output) +{ + int h = (in_h + pad - size) / stride_y + 1; + int w = (in_w + pad - size) / stride_x + 1; + int c = in_c; + + int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x; + if(id >= n) return; + + int j = id % w; + id /= w; + int i = id % h; + id /= h; + int k = id % c; + id /= c; + int b = id; + + int w_offset = -pad / 2; + int h_offset = -pad / 2; + + int out_index = j + w*(i + h*(k + c*b)); + float max = -9999999; + int max_i = -1; + int l, m; + for(l = 0; l < size; ++l){ + for(m = 0; m < size; ++m){ + int cur_h = h_offset + i*stride_y + l; + int cur_w = w_offset + j*stride_x + m; + int index = cur_w + in_w*(cur_h + in_h*(k + b*in_c)); + int valid = (cur_h >= 0 && cur_h < in_h && + cur_w >= 0 && cur_w < in_w); + float val = (valid != 0) ? input[index] : -9999999; + max_i = (val > max) ? index : max_i; + max = (val > max) ? val : max; + } + } + output[out_index] = max; +} + +void MaxPoolingForward(dnnType* srcData, dnnType* dstData, int n, int c, int h, int w, int stride_x, int stride_y, int size, int padding, cudaStream_t stream) +{ + + int tot_size = n*c*h*w; + + int blocks = (tot_size+255)/256; + int threads = 256; + + forward_maxpool_layer_kernel<<>>(tot_size, h, w, c, stride_x, stride_y, size, padding, srcData, dstData); +} + diff --git a/tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp b/tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp new file mode 100644 index 0000000..dd4e2fb --- /dev/null +++ b/tests/csresnext50-panet-spp/csresnext50-panet-spp.cpp @@ -0,0 +1,560 @@ +#include +#include +#include "tkdnn.h" + +int main() +{ + + // Network layout + tk::dnn::dataDim_t dim(1, 3, 416, 416, 1); + tk::dnn::Network net(dim); + + // create csresnext50-panet-spp model + std::string bin_path = "../tests/csresnext50-panet-spp"; + int classes = 80; + tk::dnn::Yolo *yolo[3]; + + std::string input_bin = bin_path + "/layers/input.bin"; + std::string output_bin = bin_path + "/debug/layer137_out.bin"; + std::vector output_bins = { + bin_path + "/debug/layer115_out.bin", + bin_path + "/debug/layer126_out.bin", + bin_path + "/debug/layer137_out.bin"}; + std::string c0_bin = bin_path + "/layers/c0.bin"; + std::string c2_bin = bin_path + "/layers/c2.bin"; + std::string c4_bin = bin_path + "/layers/c4.bin"; + std::string c5_bin = bin_path + "/layers/c5.bin"; + std::string c6_bin = bin_path + "/layers/c6.bin"; + std::string c7_bin = bin_path + "/layers/c7.bin"; + std::string c9_bin = bin_path + "/layers/c9.bin"; + std::string c10_bin = bin_path + "/layers/c10.bin"; + std::string c11_bin = bin_path + "/layers/c11.bin"; + std::string c13_bin = bin_path + "/layers/c13.bin"; + std::string c14_bin = bin_path + "/layers/c14.bin"; + std::string c15_bin = bin_path + "/layers/c15.bin"; + std::string c17_bin = bin_path + "/layers/c17.bin"; + std::string c19_bin = bin_path + "/layers/c19.bin"; + std::string c20_bin = bin_path + "/layers/c20.bin"; + std::string c21_bin = bin_path + "/layers/c21.bin"; + std::string c23_bin = bin_path + "/layers/c23.bin"; + std::string c24_bin = bin_path + "/layers/c24.bin"; + std::string c25_bin = bin_path + "/layers/c25.bin"; + std::string c26_bin = bin_path + "/layers/c26.bin"; + std::string c28_bin = bin_path + "/layers/c28.bin"; + std::string c29_bin = bin_path + "/layers/c29.bin"; + std::string c30_bin = bin_path + "/layers/c30.bin"; + std::string c32_bin = bin_path + "/layers/c32.bin"; + std::string c33_bin = bin_path + "/layers/c33.bin"; + std::string c34_bin = bin_path + "/layers/c34.bin"; + std::string c36_bin = bin_path + "/layers/c36.bin"; + std::string c38_bin = bin_path + "/layers/c38.bin"; + std::string c39_bin = bin_path + "/layers/c39.bin"; + std::string c40_bin = bin_path + "/layers/c40.bin"; + std::string c42_bin = bin_path + "/layers/c42.bin"; + std::string c43_bin = bin_path + "/layers/c43.bin"; + std::string c44_bin = bin_path + "/layers/c44.bin"; + std::string c45_bin = bin_path + "/layers/c45.bin"; + std::string c47_bin = bin_path + "/layers/c47.bin"; + std::string c48_bin = bin_path + "/layers/c48.bin"; + std::string c49_bin = bin_path + "/layers/c49.bin"; + std::string c51_bin = bin_path + "/layers/c51.bin"; + std::string c52_bin = bin_path + "/layers/c52.bin"; + std::string c53_bin = bin_path + "/layers/c53.bin"; + std::string c55_bin = bin_path + "/layers/c55.bin"; + std::string c56_bin = bin_path + "/layers/c56.bin"; + std::string c57_bin = bin_path + "/layers/c57.bin"; + std::string c59_bin = bin_path + "/layers/c59.bin"; + std::string c60_bin = bin_path + "/layers/c60.bin"; + std::string c61_bin = bin_path + "/layers/c61.bin"; + std::string c63_bin = bin_path + "/layers/c63.bin"; + std::string c65_bin = bin_path + "/layers/c65.bin"; + std::string c66_bin = bin_path + "/layers/c66.bin"; + std::string c67_bin = bin_path + "/layers/c67.bin"; + std::string c69_bin = bin_path + "/layers/c69.bin"; + std::string c70_bin = bin_path + "/layers/c70.bin"; + std::string c71_bin = bin_path + "/layers/c71.bin"; + std::string c72_bin = bin_path + "/layers/c72.bin"; + std::string c74_bin = bin_path + "/layers/c74.bin"; + std::string c75_bin = bin_path + "/layers/c75.bin"; + std::string c76_bin = bin_path + "/layers/c76.bin"; + std::string c78_bin = bin_path + "/layers/c78.bin"; + std::string c80_bin = bin_path + "/layers/c80.bin"; + std::string c81_bin = bin_path + "/layers/c81.bin"; + std::string c82_bin = bin_path + "/layers/c82.bin"; + std::string c83_bin = bin_path + "/layers/c83.bin"; + std::string c90_bin = bin_path + "/layers/c90.bin"; + std::string c91_bin = bin_path + "/layers/c91.bin"; + std::string c92_bin = bin_path + "/layers/c92.bin"; + std::string c93_bin = bin_path + "/layers/c93.bin"; + std::string c96_bin = bin_path + "/layers/c96.bin"; + std::string c98_bin = bin_path + "/layers/c98.bin"; + std::string c99_bin = bin_path + "/layers/c99.bin"; + std::string c100_bin = bin_path + "/layers/c100.bin"; + std::string c101_bin = bin_path + "/layers/c101.bin"; + std::string c102_bin = bin_path + "/layers/c102.bin"; + std::string c103_bin = bin_path + "/layers/c103.bin"; + std::string c106_bin = bin_path + "/layers/c106.bin"; + std::string c108_bin = bin_path + "/layers/c108.bin"; + std::string c109_bin = bin_path + "/layers/c109.bin"; + std::string c110_bin = bin_path + "/layers/c110.bin"; + std::string c111_bin = bin_path + "/layers/c111.bin"; + std::string c112_bin = bin_path + "/layers/c112.bin"; + std::string c113_bin = bin_path + "/layers/c113.bin"; + std::string c114_bin = bin_path + "/layers/c114.bin"; + std::string c117_bin = bin_path + "/layers/c117.bin"; + std::string c119_bin = bin_path + "/layers/c119.bin"; + std::string c120_bin = bin_path + "/layers/c120.bin"; + std::string c121_bin = bin_path + "/layers/c121.bin"; + std::string c122_bin = bin_path + "/layers/c122.bin"; + std::string c123_bin = bin_path + "/layers/c123.bin"; + std::string c124_bin = bin_path + "/layers/c124.bin"; + std::string c125_bin = bin_path + "/layers/c125.bin"; + std::string c128_bin = bin_path + "/layers/c128.bin"; + std::string c130_bin = bin_path + "/layers/c130.bin"; + std::string c131_bin = bin_path + "/layers/c131.bin"; + std::string c132_bin = bin_path + "/layers/c132.bin"; + std::string c133_bin = bin_path + "/layers/c133.bin"; + std::string c134_bin = bin_path + "/layers/c134.bin"; + std::string c135_bin = bin_path + "/layers/c135.bin"; + std::string c136_bin = bin_path + "/layers/c136.bin"; + std::string g115_bin = bin_path + "/layers/g115.bin"; + std::string g126_bin = bin_path + "/layers/g126.bin"; + std::string g137_bin = bin_path + "/layers/g137.bin"; + + tk::dnn::Conv2d c0(&net, 64, 7, 7, 2, 2, 3, 3, c0_bin, true); + tk::dnn::Activation a0(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Pooling p1(&net, 2, 2, 2, 2, 0,0, tk::dnn::POOLING_MAX); + + tk::dnn::Conv2d c2(&net, 128, 1, 1, 1, 1, 0, 0, c2_bin, true); + tk::dnn::Activation a2(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Layer *r3_layers[1] = {&p1}; + tk::dnn::Route r3(&net, r3_layers, 1); + + tk::dnn::Conv2d c4(&net, 64, 1, 1, 1, 1, 0, 0, c4_bin, true); + tk::dnn::Activation a4(&net, tk::dnn::ACTIVATION_LEAKY); + + // //1-1 + tk::dnn::Conv2d c5(&net, 128, 1, 1, 1, 1, 0, 0, c5_bin, true); + tk::dnn::Activation a5(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c6(&net, 128, 3, 3, 1, 1, 1, 1, c6_bin, true, false, false, 32, false); + tk::dnn::Activation a6(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c7(&net, 128, 1, 1, 1, 1, 0, 0, c7_bin, true); + + tk::dnn::Shortcut s8(&net, &a4); + tk::dnn::Activation a8(&net, tk::dnn::ACTIVATION_LEAKY); + + //1-2 + tk::dnn::Conv2d c9(&net, 128, 1, 1, 1, 1, 0, 0, c9_bin, true); + tk::dnn::Activation a9(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c10(&net, 128, 3, 3, 1, 1, 1, 1, c10_bin, true, false, false, 32); + tk::dnn::Activation a10(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c11(&net, 128, 1, 1, 1, 1, 0, 0, c11_bin, true); + + tk::dnn::Shortcut s12(&net, &a8); + tk::dnn::Activation a12(&net, tk::dnn::ACTIVATION_LEAKY); + + //1-3 + tk::dnn::Conv2d c13(&net, 128, 1, 1, 1, 1, 0, 0, c13_bin, true); + tk::dnn::Activation a13(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c14(&net, 128, 3, 3, 1, 1, 1, 1, c14_bin, true, false, false, 32); + tk::dnn::Activation a14(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c15(&net, 128, 1, 1, 1, 1, 0, 0, c15_bin, true); + + tk::dnn::Shortcut s16(&net, &a12); + tk::dnn::Activation a16(&net, tk::dnn::ACTIVATION_LEAKY); + + // //1-T + tk::dnn::Conv2d c17(&net, 128, 1, 1, 1, 1, 0, 0, c17_bin, true); + tk::dnn::Activation a17(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Layer *r18_layers[2] = {&a17, &a2}; + tk::dnn::Route r18(&net, r18_layers, 2); + + tk::dnn::Conv2d c19(&net, 256, 1, 1, 1, 1, 0, 0, c19_bin, true); + tk::dnn::Activation a19(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c20(&net, 256, 3, 3, 2, 2, 1, 1, c20_bin, true, false, false, 32); + tk::dnn::Activation a20(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c21(&net, 256, 1, 1, 1, 1, 0, 0, c21_bin, true); + + tk::dnn::Layer *r22_layers[2] = {&a20}; + tk::dnn::Route r22(&net, r22_layers, 1); + + tk::dnn::Conv2d c23(&net, 256, 1, 1, 1, 1, 0, 0, c23_bin, true); + + //2-1 + tk::dnn::Conv2d c24(&net, 256, 1, 1, 1, 1, 0, 0, c24_bin, true); + tk::dnn::Activation a24(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c25(&net, 256, 3, 3, 1, 1, 1, 1, c25_bin, true, false, false, 32); + tk::dnn::Activation a25(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c26(&net, 256, 1, 1, 1, 1, 0, 0, c26_bin, true); + + tk::dnn::Shortcut s27(&net, &c23); + tk::dnn::Activation a27(&net, tk::dnn::ACTIVATION_LEAKY); + + //2-2 + tk::dnn::Conv2d c28(&net, 256, 1, 1, 1, 1, 0, 0, c28_bin, true); + tk::dnn::Activation a28(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c29(&net, 256, 3, 3, 1, 1, 1, 1, c29_bin, true, false, false, 32); + tk::dnn::Activation a29(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c30(&net, 256, 1, 1, 1, 1, 0, 0, c30_bin, true); + + tk::dnn::Shortcut s31(&net, &a27); + tk::dnn::Activation a31(&net, tk::dnn::ACTIVATION_LEAKY); + + //2-3 + tk::dnn::Conv2d c32(&net, 256, 1, 1, 1, 1, 0, 0, c32_bin, true); + tk::dnn::Activation a32(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c33(&net, 256, 3, 3, 1, 1, 1, 1, c33_bin, true, false, false, 32); + tk::dnn::Activation a33(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c34(&net, 256, 1, 1, 1, 1, 0, 0, c34_bin, true); + + tk::dnn::Shortcut s35(&net, &a31); + tk::dnn::Activation a35(&net, tk::dnn::ACTIVATION_LEAKY); + + // //2-T + tk::dnn::Conv2d c36(&net, 256, 1, 1, 1, 1, 0, 0, c36_bin, true); + tk::dnn::Activation a36(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Layer *r37_layers[2] = {&a36, &c21}; + tk::dnn::Route r37(&net, r37_layers, 2); + + tk::dnn::Conv2d c38(&net, 512, 1, 1, 1, 1, 0, 0, c38_bin, true); + tk::dnn::Activation a38(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c39(&net, 512, 3, 3, 2, 2, 1, 1, c39_bin, true, false, false, 32); + tk::dnn::Activation a39(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c40(&net, 512, 1, 1, 1, 1, 0, 0, c40_bin, true); + + tk::dnn::Layer *r41_layers[2] = {&a39}; + tk::dnn::Route r41(&net, r41_layers, 1); + + tk::dnn::Conv2d c42(&net, 512, 1, 1, 1, 1, 0, 0, c42_bin, true); + + //3-1 + tk::dnn::Conv2d c43(&net, 512, 1, 1, 1, 1, 0, 0, c43_bin, true); + tk::dnn::Activation a43(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c44(&net, 512, 3, 3, 1, 1, 1, 1, c44_bin, true, false, false, 32); + tk::dnn::Activation a44(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c45(&net, 512, 1, 1, 1, 1, 0, 0, c45_bin, true); + + tk::dnn::Shortcut s46(&net, &c42); + tk::dnn::Activation a46(&net, tk::dnn::ACTIVATION_LEAKY); + + //3-2 + tk::dnn::Conv2d c47(&net, 512, 1, 1, 1, 1, 0, 0, c47_bin, true); + tk::dnn::Activation a47(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c48(&net, 512, 3, 3, 1, 1, 1, 1, c48_bin, true, false, false, 32); + tk::dnn::Activation a48(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c49(&net, 512, 1, 1, 1, 1, 0, 0, c49_bin, true); + + tk::dnn::Shortcut s50(&net, &a46); + tk::dnn::Activation a50(&net, tk::dnn::ACTIVATION_LEAKY); + + //3-3 + tk::dnn::Conv2d c51(&net, 512, 1, 1, 1, 1, 0, 0, c51_bin, true); + tk::dnn::Activation a51(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c52(&net, 512, 3, 3, 1, 1, 1, 1, c52_bin, true, false, false, 32); + tk::dnn::Activation a52(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c53(&net, 512, 1, 1, 1, 1, 0, 0, c53_bin, true); + + tk::dnn::Shortcut s54(&net, &a50); + tk::dnn::Activation a54(&net, tk::dnn::ACTIVATION_LEAKY); + + //3-4 + tk::dnn::Conv2d c55(&net, 512, 1, 1, 1, 1, 0, 0, c55_bin, true); + tk::dnn::Activation a55(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c56(&net, 512, 3, 3, 1, 1, 1, 1, c56_bin, true, false, false, 32); + tk::dnn::Activation a56(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c57(&net, 512, 1, 1, 1, 1, 0, 0, c57_bin, true); + + tk::dnn::Shortcut s58(&net, &a54); + tk::dnn::Activation a58(&net, tk::dnn::ACTIVATION_LEAKY); + + //3-5 + tk::dnn::Conv2d c59(&net, 512, 1, 1, 1, 1, 0, 0, c59_bin, true); + tk::dnn::Activation a59(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c60(&net, 512, 3, 3, 1, 1, 1, 1, c60_bin, true, false, false, 32); + tk::dnn::Activation a60(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c61(&net, 512, 1, 1, 1, 1, 0, 0, c61_bin, true); + + tk::dnn::Shortcut s62(&net, &a58); + tk::dnn::Activation a62(&net, tk::dnn::ACTIVATION_LEAKY); + + //3-T + tk::dnn::Conv2d c63(&net, 512, 1, 1, 1, 1, 0, 0, c63_bin, true); + tk::dnn::Activation a63(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Layer *r64_layers[2] = {&a63, &c40}; + tk::dnn::Route r64(&net, r64_layers, 2); + + tk::dnn::Conv2d c65(&net, 1024, 1, 1, 1, 1, 0, 0, c65_bin, true); + tk::dnn::Activation a65(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c66(&net, 1024, 3, 3, 2, 2, 1, 1, c66_bin, true, false, false, 32); + tk::dnn::Activation a66(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c67(&net, 1024, 1, 1, 1, 1, 0, 0, c67_bin, true); + tk::dnn::Activation a67(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Layer *r68_layers[2] = {&a66}; + tk::dnn::Route r68(&net, r68_layers, 1); + + tk::dnn::Conv2d c69(&net, 1024, 1, 1, 1, 1, 0, 0, c69_bin, true); + tk::dnn::Activation a69(&net, tk::dnn::ACTIVATION_LEAKY); + + //4-1 + tk::dnn::Conv2d c70(&net, 1024, 1, 1, 1, 1, 0, 0, c70_bin, true); + tk::dnn::Activation a70(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c71(&net, 1024, 3, 3, 1, 1, 1, 1, c71_bin, true, false, false, 32); + tk::dnn::Activation a71(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c72(&net, 1024, 1, 1, 1, 1, 0, 0, c72_bin, true); + + tk::dnn::Shortcut s73(&net, &a69); + tk::dnn::Activation a73(&net, tk::dnn::ACTIVATION_LEAKY); + + //4-2 + tk::dnn::Conv2d c74(&net, 1024, 1, 1, 1, 1, 0, 0, c74_bin, true); + tk::dnn::Activation a74(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c75(&net, 1024, 3, 3, 1, 1, 1, 1, c75_bin, true, false, false, 32); + tk::dnn::Activation a75(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c76(&net, 1024, 1, 1, 1, 1, 0, 0, c76_bin, true); + + tk::dnn::Shortcut s77(&net, &a73); + tk::dnn::Activation a77(&net, tk::dnn::ACTIVATION_LEAKY); + + //4-T + tk::dnn::Conv2d c78(&net, 1024, 1, 1, 1, 1, 0, 0, c78_bin, true); + tk::dnn::Activation a78(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Layer *r79_layers[2] = {&a78, &a67}; + tk::dnn::Route r79(&net, r79_layers, 2); + + tk::dnn::Conv2d c80(&net, 2048, 1, 1, 1, 1, 0, 0, c80_bin, true); + tk::dnn::Activation a80(&net, tk::dnn::ACTIVATION_LEAKY); + + // //////////////////// + + tk::dnn::Conv2d c81(&net, 512, 1, 1, 1, 1, 0, 0, c81_bin, true); + tk::dnn::Activation a81(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c82(&net, 1024, 3, 3, 1, 1, 1, 1, c82_bin, true); + tk::dnn::Activation a82(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c83(&net, 512, 1, 1, 1, 1, 0, 0, c83_bin, true); + tk::dnn::Activation a83(&net, tk::dnn::ACTIVATION_LEAKY); + + // //SPP + tk::dnn::Pooling p84(&net, 5, 5, 1, 1,0,0, tk::dnn::POOLING_MAX, false, true); + tk::dnn::Layer *r85_layers[1] = {&a83}; + tk::dnn::Route r85(&net, r85_layers, 1); + + tk::dnn::Pooling p86(&net, 9, 9, 1, 1,0,0, tk::dnn::POOLING_MAX, false, true); + tk::dnn::Layer *r87_layers[1] = {&a83}; + tk::dnn::Route r87(&net, r87_layers, 1); + + tk::dnn::Pooling p88(&net, 13, 13, 1, 1, 12, 12, tk::dnn::POOLING_MAX, false, true); + tk::dnn::Layer *r89_layers[4] = {&p88, &p86, &p84, &a83}; + tk::dnn::Route r89(&net, r89_layers, 4); + // //END SPP + + tk::dnn::Conv2d c90(&net, 512, 1, 1, 1, 1, 0, 0, c90_bin, true); + tk::dnn::Activation a90(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c91(&net, 1024, 3, 3, 1, 1, 1, 1, c91_bin, true); + tk::dnn::Activation a91(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c92(&net, 512, 1, 1, 1, 1, 0, 0, c92_bin, true); + tk::dnn::Activation a92(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c93(&net, 256, 1, 1, 1, 1, 0, 0, c93_bin, true); + tk::dnn::Activation a93(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Upsample u94(&net, 2); + tk::dnn::Layer *r95_layers[1] = {&a65}; + tk::dnn::Route r95(&net, r95_layers, 1); + tk::dnn::Conv2d c96(&net, 256, 1, 1, 1, 1, 0, 0, c96_bin, true); + tk::dnn::Activation a96(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Layer *r97_layers[2] = {&a96,&u94}; + tk::dnn::Route r97(&net, r97_layers, 2); + + tk::dnn::Conv2d c98(&net, 256, 1, 1, 1, 1, 0, 0, c98_bin, true); + tk::dnn::Activation a98(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c99(&net, 512, 3, 3, 1, 1, 1, 1, c99_bin, true); + tk::dnn::Activation a99(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c100(&net, 256, 1, 1, 1, 1, 0, 0, c100_bin, true); + tk::dnn::Activation a100(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c101(&net, 512, 3, 3, 1, 1, 1, 1, c101_bin, true); + tk::dnn::Activation a101(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c102(&net, 256, 1, 1, 1, 1, 0, 0, c102_bin, true); + tk::dnn::Activation a102(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c103(&net, 128, 1, 1, 1, 1, 0, 0, c103_bin, true); + tk::dnn::Activation a103(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Upsample u104(&net, 2); + tk::dnn::Layer *r105_layers[1] = {&a38}; + tk::dnn::Route r105(&net, r105_layers, 1); + tk::dnn::Conv2d c106(&net, 128, 1, 1, 1, 1, 0, 0, c106_bin, true); + tk::dnn::Activation a106(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Layer *r107_layers[2] = {&a106,&u104}; + tk::dnn::Route r107(&net, r107_layers, 2); + + + tk::dnn::Conv2d c108(&net, 128, 1, 1, 1, 1, 0, 0, c108_bin, true); + tk::dnn::Activation a108(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c109(&net, 256, 3, 3, 1, 1, 1, 1, c109_bin, true); + tk::dnn::Activation a109(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c110(&net, 128, 1, 1, 1, 1, 0, 0, c110_bin, true); + tk::dnn::Activation a110(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c111(&net, 256, 3, 3, 1, 1, 1, 1, c111_bin, true); + tk::dnn::Activation a111(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c112(&net, 128, 1, 1, 1, 1, 0, 0, c112_bin, true); + tk::dnn::Activation a112(&net, tk::dnn::ACTIVATION_LEAKY); + + // ########################### + + tk::dnn::Conv2d c113(&net, 256, 3, 3, 1, 1, 1, 1, c113_bin, true); + tk::dnn::Activation a113(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c114(&net, 255, 1, 1, 1, 1, 0, 0, c114_bin, false); + tk::dnn::Yolo yolo115(&net, classes, 3, g115_bin); + + tk::dnn::Layer *r116_layers[1] = {&a112}; + tk::dnn::Route r116(&net, r116_layers, 1); + tk::dnn::Conv2d c117(&net, 256, 3, 3, 2, 2, 1, 1, c117_bin, true); + tk::dnn::Activation a117(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Layer *r118_layers[2] = {&a117,&a102}; + tk::dnn::Route r118(&net, r118_layers, 2); + + tk::dnn::Conv2d c119(&net, 256, 1, 1, 1, 1, 0, 0, c119_bin, true); + tk::dnn::Activation a119(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c120(&net, 512, 3, 3, 1, 1, 1, 1, c120_bin, true); + tk::dnn::Activation a120(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c121(&net, 256, 1, 1, 1, 1, 0, 0, c121_bin, true); + tk::dnn::Activation a121(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c122(&net, 512, 3, 3, 1, 1, 1, 1, c122_bin, true); + tk::dnn::Activation a122(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c123(&net, 256, 1, 1, 1, 1, 0, 0, c123_bin, true); + tk::dnn::Activation a123(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Conv2d c124(&net, 512, 3, 3, 1, 1, 1, 1, c124_bin, true); + tk::dnn::Activation a124(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c125(&net, 255, 1, 1, 1, 1, 0, 0, c125_bin, false); + tk::dnn::Yolo yolo126(&net, classes, 3, g126_bin); + + tk::dnn::Layer *r127_layers[1] = {&a123}; + tk::dnn::Route r127(&net, r127_layers, 1); + tk::dnn::Conv2d c128(&net, 512, 3, 3, 2, 2, 1, 1, c128_bin, true); + tk::dnn::Activation a128(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Layer *r129_layers[2] = {&a128,&a92}; + tk::dnn::Route r129(&net, r129_layers, 2); + + tk::dnn::Conv2d c130(&net, 512, 1, 1, 1, 1, 0, 0, c130_bin, true); + tk::dnn::Activation a130(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c131(&net, 1024, 3, 3, 1, 1, 1, 1, c131_bin, true); + tk::dnn::Activation a131(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c132(&net, 512, 1, 1, 1, 1, 0, 0, c132_bin, true); + tk::dnn::Activation a132(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c133(&net, 1024, 3, 3, 1, 1, 1, 1, c133_bin, true); + tk::dnn::Activation a133(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c134(&net, 512, 1, 1, 1, 1, 0, 0, c134_bin, true); + tk::dnn::Activation a134(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Conv2d c135(&net, 1024, 3, 3, 1, 1, 1, 1, c135_bin, true); + tk::dnn::Activation a135(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c136(&net, 255, 1, 1, 1, 1, 0, 0, c136_bin, false); + tk::dnn::Yolo yolo137(&net, classes, 3, g137_bin); + + + + // yolo[0] = &yolo115; + // yolo[1] = &yolo126; + // yolo[2] = &yolo137; + + // // fill classes names + // for (int i = 0; i < 3; i++) + // { + // yolo[i]->classesNames = {"person", "bicycle", "car", "motorbike", "aeroplane", "bus", "train", "truck", "boat", "traffic light", "fire hydrant", "stop sign", "parking meter", "bench", "bird", "cat", "dog", "horse", "sheep", "cow", "elephant", "bear", "zebra", "giraffe", "backpack", "umbrella", "handbag", "tie", "suitcase", "frisbee", "skis", "snowboard", "sports ball", "kite", "baseball bat", "baseball glove", "skateboard", "surfboard", "tennis racket", "bottle", "wine glass", "cup", "fork", "knife", "spoon", "bowl", "banana", "apple", "sandwich", "orange", "broccoli", "carrot", "hot dog", "pizza", "donut", "cake", "chair", "sofa", "pottedplant", "bed", "diningtable", "toilet", "tvmonitor", "laptop", "mouse", "remote", "keyboard", "cell phone", "microwave", "oven", "toaster", "sink", "refrigerator", "book", "clock", "vase", "scissors", "teddy bear", "hair drier", "toothbrush"}; + // } + + // Load input + dnnType *data; + dnnType *input_h; + readBinaryFile(input_bin, dim.tot(), &input_h, &data); + + //print network model + net.print(); + + // //convert network to tensorRT + // tk::dnn::NetworkRT netRT(&net, "csresnext50-panet-spp.rt"); + + // the network have 3 outputs + // tk::dnn::dataDim_t out_dim[3]; + // for (int i = 0; i < 3; i++) + // out_dim[i] = yolo[i]->output_dim; + // dnnType *cudnn_out[3], *rt_out[3]; + + tk::dnn::dataDim_t dim1 = dim; //input dim + printCenteredTitle(" CUDNN inference ", '=', 30); + { + dim1.print(); + TIMER_START + net.infer(dim1, data); + TIMER_STOP + dim1.print(); + } + dnnType *cudnn_out = net.layers[net.num_layers-1]->dstData; + tk::dnn::dataDim_t out_dim = net.layers[net.num_layers-1]->output_dim; + dnnType *out1, *out1_h; + int odim1 = out_dim.tot(); + readBinaryFile(output_bin, odim1, &out1_h, &out1); + std::cout << "CUDNN vs correct" << std::endl; + // printDeviceVector(odim1, cudnn_out); + checkResult(odim1, cudnn_out, out1); + + // for (int i = 0; i < 3; i++) + // cudnn_out[i] = yolo[i]->dstData; + + // printCenteredTitle(" compute detections ", '=', 30); + // TIMER_START + // int ndets = 0; + // tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes); + // for (int i = 0; i < 3; i++) + // yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); + // tk::dnn::Yolo::mergeDetections(dets, ndets, classes); + + // for (int j = 0; j < ndets; j++) + // { + // tk::dnn::Yolo::box b = dets[j].bbox; + // int x0 = (b.x - b.w / 2.); + // int x1 = (b.x + b.w / 2.); + // int y0 = (b.y - b.h / 2.); + // int y1 = (b.y + b.h / 2.); + + // int cl = 0; + // for (int c = 0; c < classes; ++c) + // { + // float prob = dets[j].prob[c]; + // if (prob > 0) + // cl = c; + // } + // std::cout << cl << ": " << x0 << " " << y0 << " " << x1 << " " << y1 << "\n"; + // } + // TIMER_STOP + + // tk::dnn::dataDim_t dim2 = dim; + // printCenteredTitle(" TENSORRT inference ", '=', 30); + // { + // dim2.print(); + // TIMER_START + // netRT.infer(dim2, data); + // TIMER_STOP + // dim2.print(); + // } + // for (int i = 0; i < 3; i++) + // rt_out[i] = (dnnType *)netRT.buffersRT[i + 1]; + + // for (int i = 0; i < 3; i++) + // { + // printCenteredTitle((std::string(" YOLO ") + std::to_string(i) + " CHECK RESULTS ").c_str(), '=', 30); + // dnnType *out, *out_h; + // int odim = out_dim[i].tot(); + // readBinaryFile(output_bins[i], odim, &out_h, &out); + // std::cout << "CUDNN vs correct"; + // checkResult(odim, cudnn_out[i], out); + // std::cout << "TRT vs correct"; + // checkResult(odim, rt_out[i], out); + // std::cout << "CUDNN vs TRT "; + // checkResult(odim, cudnn_out[i], rt_out[i]); + // } + return 0; +}