Improve Yolo, mobilenet and shelfnet preprocessing using GPU #56

Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>

CMakeLists.txt

@@ -3,10 +3,10 @@ cmake_minimum_required(VERSION 3.15)
 project (tkDNN)
 set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake)
 if(UNIX)
-set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC -Wno-deprecated-declarations -Wno-unused-variable ")
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++14 -fPIC -Wno-deprecated-declarations")
 endif()
 if(WIN32)
-set(CMAKE_CXX_STANDARD 11)
+set(CMAKE_CXX_STANDARD 14)
 set(CMAKE_CXX_FLAGS "/O2 /FS /EHsc")
 set(CMAKE_WINDOWS_EXPORT_ALL_SYMBOLS ON)
 endif(WIN32)
@@ -140,6 +140,9 @@ target_link_libraries(test_shelfnet_berkeley tkDNN)
 add_executable(test_shelfnet_mapillary tests/shelfnet/shelfnet_mapillary.cpp)
 target_link_libraries(test_shelfnet_mapillary tkDNN)
 
+add_executable(test_shelfnet_coco tests/shelfnet/shelfnet_coco.cpp)
+target_link_libraries(test_shelfnet_coco tkDNN)
+
 # DEMOS
 add_executable(test_rtinference tests/test_rtinference/rtinference.cpp)
 target_link_libraries(test_rtinference tkDNN)

Issues.md

@@ -1 +0,0 @@
-1)error C2131 @ Yolo3Detection.cpp(97) -> expression doesnt evaluate to a constant caused to read of variable outside its lifetime

README.md

@@ -17,10 +17,15 @@ If you use tkDNN in your research, please cite the [following paper](https://iee
 }
 ```
 
-### What's new (20 July 2021)
+### What's new
+#### 20 July 2021
 - [x] Support to sematic segmentation [README](docs/README_seg.md)
 - [x] Support 2D/3D Object Detection and Tracking [README](docs/README_2d3dtracking.md)
-- [ ] Support to TensorRT8 (WIP)
+#### 24 November 2021
+- [x] Support to sematic segmentation on cuda 11
+- [x] Support to TensorRT8 (thanks to [Harshvardhan Chandirasekar](https://github.com/perseusdg)). 
+
+TensorRT8 (and therefore Jetpack 4.6) is currently supported only on the branch tensorrt8 due to [performance issue with TensorRT8](https://docs.nvidia.com/deeplearning/tensorrt/release-notes/tensorrt-8.html)). We will merge it to the master as soon as those issues are fixed (probably in future minor releases).
 
 ## FPS Results
 Inference FPS of yolov4 with tkDNN, average of 1200 images with the same dimension as the input size, on 
@@ -82,7 +87,7 @@ Results for COCO val 2017 (5k images), on RTX 2080Ti, with conf threshold=0.001
 
 ## Dependencies
 This branch works on every NVIDIA GPU that supports the following (latest tested) dependencies:
-* CUDA 11.0 (or >= 10)
+* CUDA 11.0 (or >= 10) [the segmentation only works with CUDA 10 for now]
 * cuDNN 8.0.4 (or >= 7.3)
 * TensorRT 7.2.0 (or >=5)
 * OpenCV 4.5.2 (or >=4)
@@ -168,11 +173,10 @@ For specific details on how to run tkDNN on Windows 10 see [HERE](./docs/windows
 | yolo4_320         | Yolov4 <sup>8</sup>                           | [COCO 2017](http://cocodataset.org/)                          | 80        | 320x320       | [weights](https://cloud.hipert.unimore.it/s/d97CFzYqCPCp5Hg/download)     |
 | yolo4_512         | Yolov4 <sup>8</sup>                           | [COCO 2017](http://cocodataset.org/)                          | 80        | 512x512       | [weights](https://cloud.hipert.unimore.it/s/d97CFzYqCPCp5Hg/download)     |
 | yolo4_608         | Yolov4 <sup>8</sup>                           | [COCO 2017](http://cocodataset.org/)                          | 80        | 608x608       | [weights](https://cloud.hipert.unimore.it/s/d97CFzYqCPCp5Hg/download)     |
-| yolo4_berkeley             | Yolov4 <sup>8</sup>                           | [BDD100K  ](https://bair.berkeley.edu/blog/2018/05/30/bdd/)                          | 10        | 540x320       | [weights](https://cloud.hipert.unimore.it/s/nkWFa5fgb4NTdnB/download)     |
+| yolo4_berkeley             | Yolov4 <sup>8</sup>                           | [BDD100K  ](https://bair.berkeley.edu/blog/2018/05/30/bdd/)                          | 10        | 544x320       | [weights](https://cloud.hipert.unimore.it/s/nkWFa5fgb4NTdnB/download)     |
 | yolo4tiny             | Yolov4 tiny <sup>9</sup>                           | [COCO 2017](http://cocodataset.org/)                          | 80        | 416x416       | [weights](https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download)     |
-| yolo4x             | Yolov4x-mish  <sup>9</sup>                          | [COCO 2017](http://cocodataset.org/)                          | 
+| yolo4x             | Yolov4x-mish  <sup>9</sup>                          | [COCO 2017](http://cocodataset.org/)                          | 80        | 640x640       | [weights](https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download)     |
 | yolo4tiny_512           | Yolov4 tiny <sup>9</sup>                           | [COCO 2017](http://cocodataset.org/)                          | 80        | 512x512       | [weights](https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download)     |
-80        | 640x640       | [weights](https://cloud.hipert.unimore.it/s/5MFjtNtgbDGdJEo/download)     |
 | yolo4x-cps            | Scaled Yolov4 <sup>10</sup>                          | [COCO 2017](http://cocodataset.org/)                          | 80        | 512x512       | [weights](https://cloud.hipert.unimore.it/s/AfzHE4BfTeEm2gH/download)     |
 | shelfnet              | ShelfNet18_realtime<sup>11</sup>                      | [Cityscapes](https://www.cityscapes-dataset.com/)                          | 19        | 1024x1024       | [weights](https://cloud.hipert.unimore.it/s/mEDZMRJaGCFWSJF/download)                                                                   |
 | shelfnet_berkeley              | ShelfNet18_realtime<sup>11</sup>                           | [DeepDrive](https://bdd-data.berkeley.edu/)                          | 20        | 1024x1024       | [weights](https://cloud.hipert.unimore.it/s/m92e7QdD9gYMF7f/download)                                                                   |

demo/demo/demo.cpp

@@ -9,7 +9,6 @@
 #include "Yolo3Detection.h"
 
 bool gRun;
-bool SAVE_RESULT = false;
 
 void sig_handler(int signo) {
     std::cout<<"request gateway stop\n";
@@ -18,43 +17,53 @@ void sig_handler(int signo) {
 
 int main(int argc, char *argv[]) {
 
-    std::cout<<"detection\n";
     signal(SIGINT, sig_handler);
 
-
-    std::string net = "yolo4tiny_fp32.rt";
-    if(argc > 1)
-        net = argv[1]; 
+    // get config file path and read it
     #ifdef __linux__ 
-        std::string input = "../demo/yolo_test.mp4";
+        std::string config_file = "../demo/demoConfig.yaml";
     #elif _WIN32
-        std::string input = "..\\..\\..\\demo\\yolo_test.mp4";
+        std::string config_file = "..\\..\\..\\demo\\demoConfig.yaml";
     #endif
+    if(argc > 1)
+        config_file = argv[1]; 
+    
+    YAML::Node conf =  YAMLloadConf(config_file);
+    if(!conf)
+        FatalError("Problem with config file");
 
-    if(argc > 2)
-        input = argv[2]; 
-    char ntype = 'y';
-    if(argc > 3)
-        ntype = argv[3][0]; 
-    int n_classes = 80;
-    if(argc > 4)
-        n_classes = atoi(argv[4]); 
-    int n_batch = 1;
-    if(argc > 5)
-        n_batch = atoi(argv[5]); 
-    bool show = true;
-    if(argc > 6)
-        show = atoi(argv[6]); 
-    float conf_thresh=0.3;
-    if(argc > 7)
-        conf_thresh = atof(argv[7]);     
+    // read settings from config file
+    std::string net = YAMLgetConf<std::string>(conf, "net", "yolo4tiny_fp32.rt");
+    if(!fileExist(net.c_str()))
+        FatalError("The given network does not exist. Create the rt first.");
 
+    #ifdef __linux__ 
+        std::string input = YAMLgetConf<std::string>(conf, "input", "../demo/yolo_test.mp4");
+    #elif _WIN32
+        std::string input = YAMLgetConf(conf, "win_input", "..\\..\\..\\demo\\yolo_test.mp4");
+    #endif
+    if(!fileExist(input.c_str()))
+        FatalError("The given input video does not exist.");
+    
+    char ntype          = YAMLgetConf<char>(conf, "ntype", 'y');
+    int n_classes       = YAMLgetConf<int>(conf, "n_classes", 80);
+    int n_batch         = YAMLgetConf<int>(conf, "n_batch", 1);
     if(n_batch < 1 || n_batch > 64)
         FatalError("Batch dim not supported");
+    float conf_thresh   = YAMLgetConf<float>(conf, "conf_thresh", 0.3);
+    bool show           = YAMLgetConf<bool>(conf, "show", true);
+    bool save           = YAMLgetConf<bool>(conf, "save", false);
 
-    if(!show)
-        SAVE_RESULT = true;
-
+    std::cout   <<"Net settings - net: "<< net
+                <<", ntype: "<< ntype
+                <<", n_classes: "<< n_classes
+                <<", n_batch: "<< n_batch
+                <<", conf_thresh: "<< conf_thresh<<"\n"; 
+    std::cout   <<"Demo settings - input: "<< input
+                <<", show: "<< show
+                <<", save: "<< save<<"\n\n"; 
+    
+    // create detection network
     tk::dnn::Yolo3Detection yolo;
     tk::dnn::CenternetDetection cnet;
     tk::dnn::MobilenetDetection mbnet;  
@@ -79,8 +88,7 @@ int main(int argc, char *argv[]) {
 
     detNN->init(net, n_classes, n_batch, conf_thresh);
 
-    gRun = true;
-
+    // open video stream
     cv::VideoCapture cap(input);
     if(!cap.isOpened())
         gRun = false; 
@@ -88,19 +96,21 @@ int main(int argc, char *argv[]) {
         std::cout<<"camera started\n";
 
     cv::VideoWriter resultVideo;
-    if(SAVE_RESULT) {
+    if(save) {
         int w = cap.get(cv::CAP_PROP_FRAME_WIDTH);
         int h = cap.get(cv::CAP_PROP_FRAME_HEIGHT);
         resultVideo.open("result.mp4", cv::VideoWriter::fourcc('M','P','4','V'), 30, cv::Size(w, h));
     }
 
-    cv::Mat frame;
     if(show)
         cv::namedWindow("detection", cv::WINDOW_NORMAL);
 
+    cv::Mat frame;
     std::vector<cv::Mat> batch_frame;
     std::vector<cv::Mat> batch_dnn_input;
 
+    // start detection loop
+    gRun = true;
     while(gRun) {
         batch_dnn_input.clear();
         batch_frame.clear();
@@ -128,19 +138,18 @@ int main(int argc, char *argv[]) {
                 cv::waitKey(1);
             }
         }
-        if(n_batch == 1 && SAVE_RESULT)
+        if(n_batch == 1 && save)
             resultVideo << frame;
     }
 
     std::cout<<"detection end\n";   
-    double mean = 0; 
     
+    double mean = 0; 
     std::cout<<COL_GREENB<<"\n\nTime stats:\n";
-   std::cout<<"Min: "<<*std::min_element(detNN->stats.begin(), detNN->stats.end())/n_batch<<" ms\n";    
+    std::cout<<"Min: "<<*std::min_element(detNN->stats.begin(), detNN->stats.end())/n_batch<<" ms\n";    
     std::cout<<"Max: "<<*std::max_element(detNN->stats.begin(), detNN->stats.end())/n_batch<<" ms\n";    
     for(int i=0; i<detNN->stats.size(); i++) mean += detNN->stats[i]; mean /= detNN->stats.size();
     std::cout<<"Avg: "<<mean/n_batch<<" ms\t"<<1000/(mean/n_batch)<<" FPS\n"<<COL_END;   
-    
 
     return 0;
 }

demo/demoConfig.yaml

@@ -0,0 +1,14 @@
+# video input
+input     : "../demo/yolo_test.mp4"
+win_input : "..\\..\\..\\demo\\yolo_test.mp4"
+
+# network config
+net         : "yolo4tiny_fp32.rt"
+ntype       : 'y'
+n_classes   : 80
+n_batch     : 1
+conf_thresh : 0.3
+
+# demo config
+show  : true
+save  : true

docs/README_2d3dtracking.md

@@ -18,6 +18,8 @@ where
 
 * ```<calibration-file>``` is the camera calibration file (opencv format). It is important that the file contains entry "camera_matrix" with sub-entry "rows", "cols", "data". If you do not want to pass the calibration file, pass "NULL" instead.
 
+![demo](https://user-images.githubusercontent.com/11939259/126784875-c4285497-d369-424f-abda-58274cd747ac.gif)
+
 ## Object Detection and Tracking
 
 To run the 3D object detection & tracking demo follow these steps (example with CenterTrack based on DLA34):
@@ -37,6 +39,7 @@ where
 * ```<calibration-file>``` is the camera calibration file (opencv format). It is important that the file contains entry "camera_matrix" with sub-entry "rows", "cols", "data". If you do not want to pass the calibration file, pass "NULL" instead.
 *  ```<2D/3D-flag>``` if set to 0 the demo will be in the 2D mode, while if set to 1 the demo will be in the 3D mode (Default is 1 - 3D mode).
 
+![demo](https://user-images.githubusercontent.com/11939259/126784878-513fa9e8-864a-4c24-b4bd-199737184708.gif)
 
 ## FPS Results
 

docs/README_seg.md

@@ -29,7 +29,7 @@ where
 NB) By default it is used FP32 inference
 NB) The batching is not used to work on more streams, rather to work on more tiles of the same image. Shelfnet never resized the input image, therefore for images greater than 1024x1024 tiles of 1024x1024 are given in input to the network in batch. 
 
-![gif](output.gif "Results on yolo_test.mp4")  
+![demo](https://user-images.githubusercontent.com/11939259/126784236-38d24fc3-02df-4514-81c4-497e87e40b65.gif "Results on yolo_test.mp4")  
 
 For other demo videos refer to [this playlist](https://www.youtube.com/playlist?list=PLv0nEQYDD45y5EdSiywwCGPBmJVUzIWwe).
 

docs/demo.md

@@ -32,21 +32,20 @@ make
 
 Once you have successfully created your rt file, run the demo: 
 ```
-./demo yolo4_fp32.rt ../demo/yolo_test.mp4 y
+./demo <path-to-config>
 ```
-In general the demo program takes 7 parameters:
-```
-./demo <network-rt-file> <path-to-video> <kind-of-network> <number-of-classes> <n-batches> <show-flag> <conf-thresh>
-```
-where
+In general the demo program takes 1 parameter, the ```<path-to-config>``` that is the path to che configuration file. The parameter is optional and its default value is ```"../demo/demoConfig.yaml"```. 
 
-*  ```<network-rt-file>``` is the rt file generated by a test
-*  ```<<path-to-video>``` is the path to a video file or a camera input  
-*  ```<kind-of-network>``` is the type of network. Thee types are currently supported: ```y``` (YOLO family), ```c``` (CenterNet family) and ```m``` (MobileNet-SSD family)
-*  ```<number-of-classes>```is the number of classes the network is trained on
-*  ```<n-batches>``` number of batches to use in inference (N.B. you should first export TKDNN_BATCHSIZE to the required n_batches and create again the rt file for the network).
-*  ```<show-flag>``` if set to 0 the demo will not show the visualization but save the video into result.mp4 (if n-batches ==1)
-*  ```<conf-thresh>``` confidence threshold for the detector. Only bounding boxes with threshold greater than conf-thresh will be displayed.
+The config file is a yaml file with the following attributes:
+*  ```net``` is the rt file generated by a test
+*  ```input``` is the path to a video file or a camera input  (on Linux)
+*  ```win_input``` is the path to a video file or a camera input  (on Windows)
+*  ```ntype``` is the type of network. Thee types are currently supported: ```y``` (YOLO family), ```c``` (CenterNet family) and ```m``` (MobileNet-SSD family)
+*  ```n_classes``` is the number of classes the network is trained on
+*  ```n_batch``` number of batches to use in inference (N.B. you should first export TKDNN_BATCHSIZE to the required n_batches and create again the rt file for the network).
+*  ```conf_thresh``` confidence threshold for the detector. Only bounding boxes with threshold greater than conf-thresh will be displayed.
+*  ```show``` if set to 0 the demo will not show the visualization (if n-batches ==1)
+*  ```save``` if set to 1 the demo will save the video of the demo into result.mp4 (if n-batches ==1)
 
 N.B. By default it is used FP32 inference
 
@@ -61,7 +60,8 @@ To run the demo with FP16 inference follow these steps (example with yolov3):
 export TKDNN_MODE=FP16  # set the half floating point optimization
 rm yolo3_fp16.rt        # be sure to delete(or move) old tensorRT files
 ./test_yolo3            # run the yolo test (is slow)
-./demo yolo3_fp16.rt ../demo/yolo_test.mp4 y
+# set net: yolo3_fp16.rt in the config-file
+./demo 
 ```
 N.B. Using FP16 inference will lead to some errors in the results (first or second decimal). 
 
@@ -86,7 +86,8 @@ export TKDNN_CALIB_LABEL_PATH=../demo/COCO_val2017/all_labels.txt
 export TKDNN_CALIB_IMG_PATH=../demo/COCO_val2017/all_images.txt
 rm yolo3_int8.rt        # be sure to delete(or move) old tensorRT files
 ./test_yolo3            # run the yolo test (is slow)
-./demo yolo3_int8.rt ../demo/yolo_test.mp4 y
+# set net: yolo3_int8.rt in the config-file
+./demo 
 ```
 N.B. 
 

include/tkDNN/DetectionNN.h

@@ -1,6 +1,8 @@
 #ifndef DETECTIONNN_H
 #define DETECTIONNN_H
 
+#include "kernels.h"
+
 #include <iostream>
 #include <signal.h>
 #include <stdlib.h>    
@@ -16,6 +18,7 @@
 #include <opencv2/imgproc/imgproc.hpp>
 
 #include "tkdnn.h"
+#include "utilsNN.h"
 
 //#define OPENCV_CUDACONTRIB //if OPENCV has been compiled with CUDA and contrib.
 
@@ -31,6 +34,7 @@ class DetectionNN {
 
     protected:
         tk::dnn::NetworkRT *netRT = nullptr;
+        uint8_t *frame_d = nullptr;
         dnnType *input_d;
 
         std::vector<cv::Size> originalSize;
@@ -38,6 +42,7 @@ class DetectionNN {
         cv::Scalar colors[256];
 
         int nBatches = 1;
+        int frame_size = 0;
 
 #ifdef OPENCV_CUDACONTRIB
         cv::cuda::GpuMat bgr[3];

include/tkDNN/SegmentationNN.h

@@ -1,6 +1,8 @@
 #ifndef SEGMENTATIONNN_H
 #define SEGMENTATIONNN_H
 
+#include "kernels.h"
+
 #include <iostream>
 #include <signal.h>
 #include <stdlib.h>    
@@ -16,6 +18,7 @@
 #include "tkdnn.h"
 #include "NetworkViz.h"
 #include "kernelsThrust.h"
+#include "utilsNN.h"
 
 namespace tk { namespace dnn {
 
@@ -23,7 +26,9 @@ class SegmentationNN {
 
     protected:
         tk::dnn::NetworkRT *netRT = nullptr;
+        uint8_t *frame_d = nullptr;
         int nBatches = 1;
+        int frame_size = 0;
 
         std::vector<cv::Size> originalSize;
         cv::Mat bgr[3];
@@ -76,18 +81,7 @@ class SegmentationNN {
             cv::copyMakeBorder(frame, frame_cropped, top, bottom, left, right, cv::BORDER_CONSTANT, cv::Scalar(0,0,0) );
 
             tk::dnn::dataDim_t idim = netRT->input_dim;
-
-            resize(frame_cropped, frame_cropped, cv::Size(idim.w, idim.h));
-            
-            cv::split(frame_cropped, bgr);
-            for (int i = 0; i < idim.c; i++){
-                int idx = i * frame_cropped.rows * frame_cropped.cols;
-                int ch = idim.c-1 -i;
-                memcpy((void *)&input[idx + idim.tot()*bi], (void *)bgr[ch].data, frame_cropped.rows * frame_cropped.cols * sizeof(dnnType));
-            }
-
-            checkCuda(cudaMemcpyAsync(input_d+ idim.tot()*bi, input + idim.tot()*bi, idim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
-
+            resizeAndSplit(frame_cropped, &frame_d, frame_size, input_d, netRT, bi, true);
             normalize(input_d + idim.tot()*bi, idim.c, idim.h, idim.w, mean_d, stddev_d);
         }        
 
@@ -182,6 +176,7 @@ class SegmentationNN {
             checkCuda(cudaMemcpyAsync(mean_d, mean.data(), mean.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream));
             checkCuda(cudaMemcpyAsync(stddev_d, stddev.data(), stddev.size() * sizeof(float), cudaMemcpyHostToDevice, netRT->stream));
 
+            return true;
         }
         
         /**

include/tkDNN/kernels.h

@@ -48,4 +48,11 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
                          const int dst_dim, cudaStream_t stream = cudaStream_t(0));
 
 void scalAdd(dnnType* dstData, int size, float alpha, float beta, int inc, cudaStream_t stream = cudaStream_t(0));
+
+void normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev);
+void normalize(float *bgr, const int ch, const int h, const int w, const float mean, const float stddev);
+
+void interleavedToPlanar( uint8_t *d_src, float *d_dst, int s_w, int s_h, int s_c, int d_w, int d_h);
+void interleavedRGBToPlanarBGR( uint8_t *d_src, float *d_dst, int s_w, int s_h, int s_c, int d_w, int d_h);
+
 #endif //KERNELS_H

include/tkDNN/kernelsThrust.h

@@ -31,7 +31,6 @@ void sort(dnnType *src_begin, dnnType *src_end, int *idsrc);
 void topk(dnnType *src_begin, int *idsrc, int K, float *topk_scores,
             int *topk_inds, float *topk_ys, float *topk_xs);
 // void sortAndTopKonDevice(dnnType *src_begin, int *idsrc, float *topk_scores, int *topk_inds, float *topk_ys, float *topk_xs, const int size, const int K, const int n_classes);            
-void normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev);
 void transformDep(float *src_begin, float *src_end, float *dst_begin, float *dst_end);
 void subtractWithThreshold(dnnType *src_begin, dnnType *src_end, dnnType *src2_begin, dnnType *src_out, struct threshold op);
 void topKxyclasses(int *ids_begin, int *ids_end, const int K, const int size, const int wh, int *clses, int *xs, int *ys);

include/tkDNN/utils.h

@@ -21,6 +21,7 @@
 #include <ios>
 #include <chrono>
 
+#include <yaml-cpp/yaml.h>
 
 #define dnnType float
 
@@ -137,4 +138,19 @@ static inline bool isCudaPointer(void *data) {
   cudaPointerAttributes attr;
   return cudaPointerGetAttributes(&attr, data) == 0;
 }
+
+inline YAML::Node YAMLloadConf(const std::string& conf_file) {
+    std::cerr<<"Loading YAML: "<<conf_file<<"\n";
+    return YAML::LoadFile(conf_file);
+}
+
+template<typename T>
+inline T YAMLgetConf(YAML::Node conf, std::string key, T defaultVal) {
+    T val = defaultVal;
+    if(conf && conf[key]) {
+        val = conf[key].as<T>();
+    }
+    return val;
+}
+
 #endif //UTILS_H

include/tkDNN/utilsNN.h

@@ -0,0 +1,10 @@
+#ifndef UTILSNN_H
+#define UTILSNN_H
+
+#include "tkdnn.h"
+#include <opencv2/core/core.hpp>
+
+
+void resizeAndSplit(cv::Mat& frame, uint8_t** frame_d, int& frame_size, dnnType *input_d, tk::dnn::NetworkRT *netRT, const int bi=0, bool BGR=true);
+
+#endif // UTILSNN_H

scripts/checkExecTimes.py

@@ -0,0 +1,37 @@
+import sys
+import pandas as pd 
+
+if len(sys.argv) < 3:
+    print("Error: two csv files are needed, old first new second")
+    exit(1)
+
+old_perf_file = str(sys.argv[1])
+new_perf_file = str(sys.argv[2])
+
+verbose = False
+if len(sys.argv) == 4:
+    verbose = bool(sys.argv[3])
+
+print("Comparing {} vs {}".format(old_perf_file, new_perf_file))
+
+df_old = pd.read_csv (old_perf_file, sep=';', header=None, index_col=0)
+df_new = pd.read_csv (new_perf_file, sep=';', header=None, index_col=0)
+
+for index, row in df_new.iterrows():
+    if index in df_old.index:
+        if verbose:
+            print("New: ",row[1], row[2], row[3])
+            print("Old: ",df_old.loc[index][1], df_old.loc[index][2], df_old.loc[index][3])
+
+        print(index, end=': ')
+        if abs(row[1] - df_old.loc[index][1]) < df_old.loc[index][1]*0.1: 
+            print("similar performance")
+        elif (row[1] < df_old.loc[index][1]): 
+            print('\x1b[3;30;42m' + 'faster' + '\x1b[0m')
+        elif (row[1] > df_old.loc[index][1]): 
+            if row[1] > df_old.loc[index][1] + df_old.loc[index][1] * 0.5 :
+                print('\x1b[3;30;41m' + 'WAY SLOWER' + '\x1b[0m')
+            else:
+                print('\x1b[3;30;41m' + 'slower' + '\x1b[0m')
+
+

src/CenterTrack.cpp

@@ -17,6 +17,8 @@ bool CenterTrack::init(const std::string& tensor_path, const int n_classes, cons
     init_pre_inf();
     init_postprocessing();
     init_visualization(n_classes);
+
+    return true;
 }
 
 bool CenterTrack::init_preprocessing(){
@@ -59,6 +61,8 @@ bool CenterTrack::init_preprocessing(){
     checkCuda( cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot() * nBatches));
     checkCuda( cudaMalloc(&input_pre_inf_d, sizeof(dnnType)*dim.tot()));
     checkCuda( cudaMalloc(&d_ptrs, dim.tot() * sizeof(float)) );
+
+    return true;
 }
 
 bool CenterTrack::init_pre_inf(){
@@ -202,6 +206,8 @@ bool CenterTrack::init_postprocessing(){
     trRes.resize(nBatches);
     countTr.resize(nBatches, 0);
     trackId.resize(nBatches, 0);
+
+    return true;
 }
 
 bool CenterTrack::init_visualization(const int n_classes){
@@ -274,6 +280,8 @@ bool CenterTrack::init_visualization(const int n_classes){
     faceId.push_back({3,0,4,7});
     faceId.push_back({2,3,7,6});
     // ([[0,1,5,4], [1,2,6, 5], [2,3,7,6], [3,0,4,7]]);
+
+    return true;
 }
 
 void CenterTrack::_get_additional_inputs(){
@@ -308,7 +316,7 @@ void CenterTrack::preprocess(cv::Mat &frame, const int bi){
         }
         
         float c[] = {new_width / 2.0f, new_height /2.0f};
-        float s[] = {dim.w, dim.h};
+        float s[] = {float(dim.w), float(dim.h)};
         // float s = new_width >= new_height ? new_width : new_height;
         // ----------- get_affine_transform
         // rot_rad = pi * 0 / 100 --> 0

src/CenternetDetection.cpp

@@ -119,6 +119,7 @@ bool CenternetDetection::init(const std::string& tensor_path, const int n_classe
     dst2.at<float>(2,0)=dst2.at<float>(1,0) + (-dst2.at<float>(0,1)+dst2.at<float>(1,1) );
     dst2.at<float>(2,1)=dst2.at<float>(1,1) + (dst2.at<float>(0,0)-dst2.at<float>(1,0) );
     
+    return true;
 }
 
 

src/CenternetDetection3D.cpp

@@ -167,6 +167,8 @@ bool CenternetDetection3D::init(const std::string& tensor_path, const int n_clas
     faceId.push_back({2,3,7,6});
     faceId.push_back({3,0,4,7});
     // ([[0,1,5,4], [1,2,6, 5], [2,3,7,6], [3,0,4,7]]);
+
+    return true;
 }
 
 void CenternetDetection3D::preprocess(cv::Mat &frame, const int bi){    

src/MobilenetDetection.cpp

@@ -198,7 +198,7 @@ bool MobilenetDetection::init(const std::string& tensor_path, const int n_classe
         "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" , 
+    "toilet" , "tvmonitor" , "laptop" , "mouse" , "remote" , "keyboard" , 
         "cell phone" , "microwave" , "oven" , "toaster" , "sink" , "refrigerator" , 
         "book" , "clock" , "vase" , "scissors" , "teddy bear" , "hair drier" , "toothbrush"};
         classesNames = std::vector<std::string>(classes_names_, std::end(classes_names_));
@@ -207,42 +207,12 @@ bool MobilenetDetection::init(const std::string& tensor_path, const int n_classe
     else{
         FatalError("Number of classes not supported for mobilenet");
     }
-    return 1;
+    return true;
 }
 
 void MobilenetDetection::preprocess(cv::Mat &frame, const int bi){
-#ifdef OPENCV_CUDACONTRIB
-        //move original image on GPU
-        cv::cuda::GpuMat orig_img, frame_nomean;
-        orig_img = cv::cuda::GpuMat(frame);
-
-        //resize image, remove mean, divide by std
-        cv::cuda::resize (orig_img, orig_img, cv::Size(netRT->input_dim.w, netRT->input_dim.h)); 
-        orig_img.convertTo(frame_nomean, CV_32FC3, 1, -127);
-        frame_nomean.convertTo(imagePreproc, CV_32FC3, 1 / 128.0, 0);
-
-        //copy image into tensors
-        cv::cuda::split(imagePreproc, bgr);
-
-        for(int i=0; i < netRT->input_dim.c; i++){
-            int idx = i * imagePreproc.rows * imagePreproc.cols;
-            checkCuda( cudaMemcpy((void *)&input_d[idx + netRT->input_dim.tot()*bi], (void *)bgr[i].data, imagePreproc.rows * imagePreproc.cols* sizeof(float), cudaMemcpyDeviceToDevice) );
-        }
-#else
-        //resize image, remove mean, divide by std
-        cv::Mat frame_nomean;
-        resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
-        frame.convertTo(frame_nomean, CV_32FC3, 1, -127);
-        frame_nomean.convertTo(imagePreproc, CV_32FC3, 1 / 128.0, 0);
-
-        //copy image into tensor and copy it into GPU
-        cv::split(imagePreproc, bgr);
-        for (int i = 0; i < netRT->input_dim.c; i++){
-            int idx = i * imagePreproc.rows * imagePreproc.cols;
-            memcpy((void *)&input[idx + netRT->input_dim.tot()*bi], (void *)bgr[i].data, imagePreproc.rows * imagePreproc.cols * sizeof(dnnType));
-        }
-        checkCuda(cudaMemcpyAsync(input_d+ netRT->input_dim.tot()*bi, input + netRT->input_dim.tot()*bi, netRT->input_dim.tot() * sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
-#endif
+    resizeAndSplit(frame, &frame_d, frame_size, input_d, netRT, bi, false);
+    normalize(input_d +  netRT->input_dim.tot()*bi,  netRT->input_dim.c,  netRT->input_dim.h,  netRT->input_dim.w, 127.0f, 128.0f);
 }
 
 void MobilenetDetection::postprocess(const int bi, const bool mAP){

src/Yolo3Detection.cpp

@@ -57,38 +57,8 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
 } 
 
 void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
-#ifdef OPENCV_CUDACONTRIB
-    cv::cuda::GpuMat orig_img, img_resized;
-    orig_img = cv::cuda::GpuMat(frame);
-    cv::cuda::resize(orig_img, img_resized, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
-
-    img_resized.convertTo(imagePreproc, CV_32FC3, 1/255.0); 
-
-    //split channels
-    cv::cuda::split(imagePreproc,bgr);//split source
-
-    //write channels
-    for(int i=0; i<netRT->input_dim.c; i++) {
-        int size = imagePreproc.rows * imagePreproc.cols;
-        int ch = netRT->input_dim.c-1 -i;
-        bgr[ch].download(bgr_h); //TODO: don't copy back on CPU
-        checkCuda( cudaMemcpy(input_d + i*size + netRT->input_dim.tot()*bi, (float*)bgr_h.data, size*sizeof(dnnType), cudaMemcpyHostToDevice));
-    }
-#else
-    cv::resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
-    frame.convertTo(imagePreproc, CV_32FC3, 1/255.0); 
-
-    //split channels
-    cv::split(imagePreproc,bgr);//split source
-
-    //write channels
-    for(int i=0; i<netRT->input_dim.c; i++) {
-        int idx = i*imagePreproc.rows*imagePreproc.cols;
-        int ch = netRT->input_dim.c-1 -i;
-        memcpy((void*)&input[idx + netRT->input_dim.tot()*bi], (void*)bgr[ch].data, imagePreproc.rows*imagePreproc.cols*sizeof(dnnType));     
-    }
-    checkCuda(cudaMemcpyAsync(input_d + netRT->input_dim.tot()*bi, input + netRT->input_dim.tot()*bi, netRT->input_dim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
-#endif
+    resizeAndSplit(frame, &frame_d, frame_size, input_d, netRT, bi, true);
+    normalize(input_d +  netRT->input_dim.tot()*bi,  netRT->input_dim.c,  netRT->input_dim.h,  netRT->input_dim.w, 0.0f, 255.0f);
 }
 
 void Yolo3Detection::postprocess(const int bi, const bool mAP){

src/kernels/interleaved_to_planar.cu

@@ -0,0 +1,84 @@
+#include "kernels.h"
+
+
+__global__ void interleavedToPlanarKernel(uint8_t *src, float *dst, int s_w, int s_h, int s_c, int d_w, int d_h, float ratio_w, float ratio_h) {
+
+    int x = min( (int)(blockIdx.x * blockDim.x + threadIdx.x), d_w-1);
+    int y = min( (int)(blockIdx.y * blockDim.y + threadIdx.y), d_h-1);
+
+    float sum_r=0, sum_g=0, sum_b=0;
+
+    float x_src = (float) x * ratio_w; // + ratio_w/2;
+    float y_src = (float) y * ratio_h; // + ratio_h/2;
+
+    int r = (int) y_src;
+    int c = (int) x_src;
+    float dr = y_src - r;
+    float dc = x_src - c;
+
+    sum_r = (float) src[(r * s_w + c) * s_c] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c] * (dr) * (dc);
+    sum_g = (float) src[(r * s_w + c) * s_c + 1] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c + 1] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c + 1] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c + 1] * (dr) * (dc);
+    sum_b = (float) src[(r * s_w + c) * s_c + 2] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c + 2] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c + 2] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c + 2] * (dr) * (dc);
+
+    dst[y * d_w + x] 					= sum_r;
+    dst[y * d_w + x + d_w * d_h] 		= sum_g;
+    dst[y * d_w + x + d_w * d_h * 2] 	= sum_b;
+}
+
+__global__ void interleavedRGBToPlanarBGRKernel(uint8_t *src, float *dst, int s_w, int s_h, int s_c, int d_w, int d_h, float ratio_w, float ratio_h) {
+
+    int x = min( (int)(blockIdx.x * blockDim.x + threadIdx.x), d_w-1);
+    int y = min( (int)(blockIdx.y * blockDim.y + threadIdx.y), d_h-1);
+
+    float sum_r=0, sum_g=0, sum_b=0;
+
+    float x_src = (float) x * ratio_w; // + ratio_w/2;
+    float y_src = (float) y * ratio_h; // + ratio_h/2;
+
+    int r = (int) y_src;
+    int c = (int) x_src;
+    float dr = y_src - r;
+    float dc = x_src - c;
+
+    sum_r = (float) src[(r * s_w + c) * s_c] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c] * (dr) * (dc);
+    sum_g = (float) src[(r * s_w + c) * s_c + 1] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c + 1] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c + 1] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c + 1] * (dr) * (dc);
+    sum_b = (float) src[(r * s_w + c) * s_c + 2] * (1.0f - dr) * (1.0f - dc) +
+            (float) src[((r + 1) * s_w + c) * s_c + 2] * (dr) * (1.0f - dc) +
+            (float) src[(r * s_w + c + 1) * s_c + 2] * (1.0f - dr) * (dc) +
+            (float) src[((r + 1) * s_w + c + 1) * s_c + 2] * (dr) * (dc);
+
+    dst[y * d_w + x] 					= sum_b;
+    dst[y * d_w + x + d_w * d_h] 		= sum_g;
+    dst[y * d_w + x + d_w * d_h * 2] 	= sum_r;
+}
+
+void interleavedToPlanar( uint8_t *d_src, float *d_dst, int s_w, int s_h, int s_c, int d_w, int d_h){
+	dim3 dg( ceil( (double)d_w/32 ), ceil( (double)d_h/8 ) );
+	dim3 db( 32, 8);
+
+	interleavedToPlanarKernel<<< dg, db >>>(d_src, d_dst, s_w, s_h, s_c, d_w, d_h, (float)s_w/d_w, (float)s_h/d_h);
+	cudaDeviceSynchronize();
+}
+
+void interleavedRGBToPlanarBGR( uint8_t *d_src, float *d_dst, int s_w, int s_h, int s_c, int d_w, int d_h){
+	dim3 dg( ceil( (double)d_w/32 ), ceil( (double)d_h/8 ) );
+	dim3 db( 32, 8);
+
+	interleavedRGBToPlanarBGRKernel<<< dg, db >>>(d_src, d_dst, s_w, s_h, s_c, d_w, d_h, (float)s_w/d_w, (float)s_h/d_h);
+	cudaDeviceSynchronize();
+}

src/kernels/normalize.cu

@@ -1,4 +1,4 @@
-#include "kernelsThrust.h"
+#include "kernels.h"
 
 __global__
 void normalize_kernel(float *bgr, const int dim, const float *mean, const float *stddev){
@@ -9,8 +9,25 @@ void normalize_kernel(float *bgr, const int dim, const float *mean, const float
     
 }
 
+
+__global__
+void normalize_kernel(float *bgr, const int dim, const float mean, const float stddev){
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+    int j = blockIdx.y;
+    bgr[j*(dim)+i] = bgr[j*(dim)+i] - mean;
+    bgr[j*(dim)+i] = bgr[j*(dim)+i] / stddev;
+    
+}
+
 void normalize(float *bgr, const int ch, const int h, const int w, const float *mean, const float *stddev){
     int num_thread = 256;
     dim3 dimBlock(h*w/num_thread, ch);
     normalize_kernel<<<dimBlock, num_thread, 0>>>(bgr, h*w, mean, stddev);
+}
+
+
+void normalize(float *bgr, const int ch, const int h, const int w, const float mean, const float stddev){
+    int num_thread = 256;
+    dim3 dimBlock(h*w/num_thread, ch);
+    normalize_kernel<<<dimBlock, num_thread, 0>>>(bgr, h*w, mean, stddev);
 }

src/kernels/postprocessing.cu

@@ -46,11 +46,16 @@ void maxElem_kernel(float *src_begin, float *dst_begin, const int n_classes, con
     if (i > size)
         return;
 
-    thrust::device_ptr<float> dPbeg ( &src_begin[i*n_classes] ) ;
-    thrust::device_ptr<float> dPend = dPbeg + n_classes;
-    thrust::device_ptr<float> result = thrust::max_element(thrust::device,dPbeg, dPend);
+    float max = 0;
+    int max_idx = 0;
+    for( int j = i*n_classes; j < i*n_classes + n_classes; ++j ){
+        if( src_begin[j] > max ){
+            max = src_begin[j];
+            max_idx = j;
+        }
+    }
 
-    dst_begin[i] = result - dPbeg;
+    dst_begin[i] = max_idx - i*n_classes;
 }
 
 void maxElem(dnnType *src_begin, dnnType *dst_begin, const int c, const int h, const int w){

src/utilsNN.cpp

@@ -0,0 +1,41 @@
+#include "kernels.h"
+#include "utilsNN.h"
+
+
+void resizeAndSplit(cv::Mat& frame, uint8_t** frame_d, int& frame_size, dnnType *input_d, tk::dnn::NetworkRT *netRT, const int bi, bool BGR){
+    int new_frame_size = sizeof(uint8_t) * frame.cols * frame.rows * frame.channels();
+    if(*frame_d == nullptr){
+        frame_size = new_frame_size;
+        checkCuda(cudaMalloc(frame_d, frame_size));
+    }
+    else{
+        if(new_frame_size > frame_size ){
+            frame_size = new_frame_size;
+            checkCuda(cudaFree(frame_d));
+            checkCuda(cudaMalloc(frame_d, frame_size));
+        }
+    }
+    
+    checkCuda(cudaMemcpyAsync(*frame_d, frame.data, frame_size, cudaMemcpyHostToDevice, netRT->stream));
+    if(BGR){
+        interleavedRGBToPlanarBGR(*frame_d,
+                                input_d + netRT->input_dim.tot() * bi, 
+                                frame.cols, 
+                                frame.rows, 
+                                frame.channels(),
+                                netRT->input_dim.w,
+                                netRT->input_dim.h
+                            );
+    }
+    else{
+        interleavedToPlanar(*frame_d,
+                                input_d + netRT->input_dim.tot() * bi, 
+                                frame.cols, 
+                                frame.rows, 
+                                frame.channels(),
+                                netRT->input_dim.w,
+                                netRT->input_dim.h
+                            );
+
+    }
+}

tests/darknet/names/crowdhuman.names

@@ -0,0 +1,2 @@
+person
+head

tests/darknet/yolo4-csp_crowd.cpp

@@ -0,0 +1,34 @@
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+#include "test.h"
+#include "DarknetParser.h"
+
+int main() {
+    std::string bin_path  = "yolo4-csp_crowd";
+    std::vector<std::string> input_bins = { 
+        bin_path + "/layers/input.bin"
+    };
+    std::vector<std::string> output_bins = {
+        bin_path + "/debug/layer144_out.bin",
+        bin_path + "/debug/layer159_out.bin",
+        bin_path + "/debug/layer174_out.bin"
+    };
+    std::string wgs_path  = bin_path + "/layers";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4-csp_crowd.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/crowdhuman.names";
+    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/RKWfWNmWXfJigsK/download");
+
+    // parse darknet network
+    tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
+    net->print();
+
+    //convert network to tensorRT
+    tk::dnn::NetworkRT *netRT = new tk::dnn::NetworkRT(net, net->getNetworkRTName(bin_path.c_str()));
+    
+    int ret = testInference(input_bins, output_bins, net, netRT);
+    net->releaseLayers();
+    delete net;
+    delete netRT;
+    return ret;
+}

tests/darknet/yolo4_berkeley_f1.cpp

@@ -17,7 +17,7 @@ int main() {
     std::string wgs_path  = bin_path + "/layers";
     std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4_berkeley.cfg";
     std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/berkeley.names";
-    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/M7WJdGoGDaDACnN/download");
+    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/q9dwoqQ5YQqEi7s/download");
 
     // parse darknet network
     tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);

tests/shelfnet/shelfnet_coco.cpp

@@ -0,0 +1,295 @@
+#include <iostream>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/imgproc/imgproc.hpp>
+
+#include "tkdnn.h"
+#include "NetworkViz.h"
+
+
+const char *input_bin = "shelfnet_coco/debug/input.bin";
+
+const char *backbone[] = {
+    "shelfnet_coco/layers/backbone-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer1-0-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer1-0-conv2.bin",
+    "shelfnet_coco/layers/backbone-layer1-1-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer1-1-conv2.bin",
+    "shelfnet_coco/layers/backbone-layer2-0-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer2-0-conv2.bin",
+    "shelfnet_coco/layers/backbone-layer2-0-downsample-0.bin",
+    "shelfnet_coco/layers/backbone-layer2-1-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer2-1-conv2.bin",
+    "shelfnet_coco/layers/backbone-layer3-0-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer3-0-conv2.bin",
+    "shelfnet_coco/layers/backbone-layer3-0-downsample-0.bin",
+    "shelfnet_coco/layers/backbone-layer3-1-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer3-1-conv2.bin",
+    "shelfnet_coco/layers/backbone-layer4-0-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer4-0-conv2.bin",
+    "shelfnet_coco/layers/backbone-layer4-0-downsample-0.bin",
+    "shelfnet_coco/layers/backbone-layer4-1-conv1.bin",
+    "shelfnet_coco/layers/backbone-layer4-1-conv2.bin"};
+
+const char *conv_out[] = {
+    "shelfnet_coco/layers/conv_out-conv-conv.bin",
+    "shelfnet_coco/layers/conv_out-conv_out.bin",
+    "shelfnet_coco/layers/conv_out16-conv-conv.bin",
+    "shelfnet_coco/layers/conv_out16-conv_out.bin",
+    "shelfnet_coco/layers/conv_out32-conv-conv.bin",
+    "shelfnet_coco/layers/conv_out32-conv_out.bin"
+    };
+
+const char *decoder[] = {
+    "shelfnet_coco/layers/decoder-bottom-conv1.bin",
+    "shelfnet_coco/layers/decoder-bottom-conv12.bin",
+    "shelfnet_coco/layers/decoder-up_conv_list-0-conv-conv.bin",
+    "shelfnet_coco/layers/decoder-up_conv_list-0-conv_atten.bin",
+    "shelfnet_coco/layers/decoder-up_dense_list-0-conv.bin",
+    "shelfnet_coco/layers/decoder-up_conv_list-1-conv-conv.bin",
+    "shelfnet_coco/layers/decoder-up_conv_list-1-conv_atten.bin",
+    "shelfnet_coco/layers/decoder-up_dense_list-1-conv.bin"
+    };
+
+    
+const char *ladder[] = {
+    "shelfnet_coco/layers/ladder-inconv-conv1.bin",
+    "shelfnet_coco/layers/ladder-inconv-conv12.bin",
+    "shelfnet_coco/layers/ladder-down_module_list-0-conv1.bin",
+    "shelfnet_coco/layers/ladder-down_module_list-0-conv12.bin",
+    "shelfnet_coco/layers/ladder-down_conv_list-0.bin",
+
+    "shelfnet_coco/layers/ladder-down_module_list-1-conv1.bin",
+    "shelfnet_coco/layers/ladder-down_module_list-1-conv12.bin",
+    "shelfnet_coco/layers/ladder-down_conv_list-1.bin",
+
+    "shelfnet_coco/layers/ladder-bottom-conv1.bin",
+    "shelfnet_coco/layers/ladder-bottom-conv12.bin",
+    
+    
+    
+    "shelfnet_coco/layers/ladder-up_conv_list-0-conv-conv.bin",
+    "shelfnet_coco/layers/ladder-up_conv_list-0-conv_atten.bin",
+    "shelfnet_coco/layers/ladder-up_dense_list-0-conv.bin",
+
+    
+    "shelfnet_coco/layers/ladder-up_conv_list-1-conv-conv.bin",
+    "shelfnet_coco/layers/ladder-up_conv_list-1-conv_atten.bin",
+    "shelfnet_coco/layers/ladder-up_dense_list-1-conv.bin"};
+
+const char *trans[] = {
+    "shelfnet_coco/layers/trans1-conv.bin",
+    "shelfnet_coco/layers/trans2-conv.bin",
+    "shelfnet_coco/layers/trans3-conv.bin"};
+int main()
+{
+
+    downloadWeightsifDoNotExist(input_bin, "shelfnet_coco", "https://cloud.hipert.unimore.it/s/KfQ9fGJQsgzNbiW/download");
+
+    int classes = 183;
+
+    // Network layout    
+    tk::dnn::dataDim_t dim(1, 3, 1024, 1024, 1);
+    tk::dnn::Network net(dim);
+
+    int bi = 0, di = 0, li = 0, ci = 0;
+    new tk::dnn::Conv2d(&net, 64, 7, 7, 2, 2, 3, 3, backbone[bi++], true);
+    new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+    tk::dnn::Layer* last = new tk::dnn::Pooling (&net, 3, 3, 2, 2, 1, 1, tk::dnn::POOLING_MAX);
+
+
+    
+    for(int i=0; i<2; ++i){
+        new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true);
+        new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, backbone[bi++], true);
+        new tk::dnn::Shortcut(&net, last);
+        last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+    }
+
+    std::vector<tk::dnn::Layer*> features;
+    for(int i=0;i<3;++i){
+        int out_channel = pow(2,7+i);
+        std::cout<<out_channel<<std::endl;
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 2, 2, 1, 1, backbone[bi++], true);
+        new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        tk::dnn::Layer* bn2 = new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, backbone[bi++], true);
+        new tk::dnn::Route(&net, &last, 1);
+        new tk::dnn::Conv2d (&net, out_channel, 1, 1, 2, 2, 0, 0, backbone[bi++], true);
+        new tk::dnn::Shortcut(&net, bn2);
+        last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, backbone[bi++], true);
+        new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, backbone[bi++], true);
+        
+        new tk::dnn::Shortcut(&net, last);
+        last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+        features.push_back(last);
+    }
+
+    for(int i=0; i<features.size(); ++i){
+        new tk::dnn::Route(&net, &features[i], 1);
+        int out_channel = pow(2,6+i);
+        new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, trans[i], true);
+        features[i] = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+    }
+
+    //DECODER
+    
+    last = features[2];
+    std::vector<tk::dnn::Layer*> up_out;
+    //bottom
+    new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, decoder[di++], true, false, 1, true);
+    new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+    new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, decoder[di++], true, false, 1, true);
+    new tk::dnn::Shortcut(&net, last);
+    last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+    up_out.push_back(last);
+
+    for(int i=0; i<2; ++i){
+        int out_channel = pow(2,7-i);
+        //up-conv
+        std::cout<<out_channel<<std::endl;
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, decoder[di++], true);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        
+        new tk::dnn::Pooling(&net, last->output_dim.w, last->output_dim.h, last->output_dim.w, last->output_dim.h, 0, 0, tk::dnn::POOLING_AVERAGE);
+        new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, decoder[di++], true);
+        
+        tk::dnn::Layer* act = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_SIGMOID);
+        new tk::dnn::Route(&net, &last, 1);
+        new tk::dnn::Shortcut(&net, act, true);
+
+        //interpolate
+        new tk::dnn::Resize(&net, 1,2,2);
+        new tk::dnn::Shortcut(&net, features[1-i]);
+
+        //up-dense
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, decoder[di++], true);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        up_out.push_back(last);
+    }
+
+    //LADDER
+
+    std::vector<tk::dnn::Layer*> down_out;
+    new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+    new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+    new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+    new tk::dnn::Shortcut(&net, last);
+    new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+    
+    for(int i=0; i<2;++i){
+        int out_channel = pow(2,6+i);
+        tk::dnn::Layer* l_last = new tk::dnn::Shortcut(&net, up_out[2-i]);
+    
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+        new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+        new tk::dnn::Shortcut(&net, l_last);
+        l_last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+        down_out.push_back(l_last);
+
+        new tk::dnn::Conv2d (&net, out_channel*2, 3, 3, 2, 2, 1, 1, ladder[li++], false);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.0f); //should be ReLU
+    }
+
+    new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+    new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+    new tk::dnn::Conv2d (&net, 256, 3, 3, 1, 1, 1, 1, ladder[li++], true, false, 1, true);
+    new tk::dnn::Shortcut(&net, last);
+    last = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_RELU);
+    up_out.clear();
+    up_out.push_back(last);
+
+    for(int i=0; i<2; ++i){
+        int out_channel = pow(2,7-i);
+        //up-conv
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        
+        new tk::dnn::Pooling(&net, last->output_dim.w, last->output_dim.h, last->output_dim.w, last->output_dim.h, 0, 0, tk::dnn::POOLING_AVERAGE);
+        new tk::dnn::Conv2d (&net, out_channel, 1, 1, 1, 1, 0, 0, ladder[li++], true);
+        
+        tk::dnn::Layer* act = new tk::dnn::Activation (&net, CUDNN_ACTIVATION_SIGMOID);
+        new tk::dnn::Route(&net, &last, 1);
+        new tk::dnn::Shortcut(&net, act, true);
+
+        //interpolate
+        new tk::dnn::Resize(&net, 1,2,2);
+        new tk::dnn::Shortcut(&net, down_out[1-i]);
+
+        // //up-dense
+        new tk::dnn::Conv2d (&net, out_channel, 3, 3, 1, 1, 1, 1, ladder[li++], true);
+        last = new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        up_out.push_back(last);
+    }
+
+
+    // for(int i=2;i>=0;--i){
+        // new tk::dnn::Route(&net, &up_out[i], 1);
+        new tk::dnn::Conv2d (&net, 64, 3, 3, 1, 1, 1, 1, conv_out[ci++], true);
+        new tk::dnn::Activation (&net, tk::dnn::ACTIVATION_LEAKY, 0.0f, 0.01);
+        new tk::dnn::Conv2d (&net, classes, 3, 3, 1, 1, 1, 1, conv_out[ci++], false);
+        /*up_out[i] =*/ new tk::dnn::Resize(&net, classes, net.input_dim.h, net.input_dim.w, true, tk::dnn::ResizeMode_t::LINEAR);
+    // // }
+
+    new tk::dnn::Softmax(&net);
+    
+    const char *output_bin = "shelfnet_coco/debug/softmax.bin";
+    
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+    std::cout<<"Input:"<<std::endl;
+
+    //print network model
+    net.print();
+
+    // // convert network to tensorRT
+    tk::dnn::NetworkRT netRT(&net, net.getNetworkRTName("shelfnet_coco"));
+
+    tk::dnn::dataDim_t dim1 = dim; //input dim
+    dnnType *cudnn_out = nullptr;
+    printCenteredTitle(" CUDNN inference ", '=', 30);
+    {
+        dim1.print();
+        TKDNN_TSTART
+        cudnn_out = net.infer(dim1, data);
+        TKDNN_TSTOP
+        dim1.print();
+    }
+
+    tk::dnn::dataDim_t dim2 = dim;
+    printCenteredTitle(" TENSORRT inference ", '=', 30);
+    {
+        dim2.print();
+        TKDNN_TSTART
+        netRT.infer(dim2, data);
+        TKDNN_TSTOP
+        dim2.print();
+    }
+
+    dnnType *rt_out1 = (dnnType *)netRT.buffersRT[1];
+
+    printCenteredTitle(std::string(" CHECK RESULTS ").c_str(), '=', 30);
+    dnnType *out1, *out1_h;
+    int odim1 = dim1.tot();
+    readBinaryFile(output_bin, odim1, &out1_h, &out1);
+
+    int ret_cudnn = 0, ret_tensorrt = 0, ret_cudnn_tensorrt = 0; 
+    // std::cout << "CUDNN vs correct" << std::endl;
+    // ret_cudnn |= checkResult(odim1, cudnn_out, out1, true, 20) == 0 ? 0 : ERROR_CUDNN;
+
+    std::cout << "TRT   vs correct" << std::endl;
+    ret_tensorrt |=checkResult(odim1, rt_out1, out1) == 0 ? 0 : ERROR_TENSORRT;
+
+    std::cout << "CUDNN vs TRT    " << std::endl;
+    ret_cudnn_tensorrt |= checkResult(odim1, cudnn_out, rt_out1) == 0 ? 0 : ERROR_CUDNNvsTENSORRT;
+    
+    cv::Mat viz = vizLayer2Mat(&net, net.num_layers-1);
+    cv::imwrite("test.png", viz);
+
+    return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt;
+}