Modified resize and boxes coordinates to float, to achieve same darknet accuracy

Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>
fix coords convert
2020-08-05 17:25:33 +02:00 · 2020-07-27 13:52:39 +02:00 · 2020-07-27 13:45:42 +02:00 · 2020-07-17 14:37:10 +02:00 · 2020-07-16 19:06:54 +02:00 · 2020-07-16 18:16:09 +02:00
44 changed files with 1219 additions and 452 deletions
@@ -13,4 +13,4 @@ build/
 *.pk
 *.table
 demo/COCO_val2017
-demo/BDD100k_val
+demo/BDD100K_val
@@ -10,6 +10,7 @@ if(DEBUG)
    add_definitions(-DDEBUG)
 endif()

+add_definitions(-DTKDNN_PATH="${CMAKE_CURRENT_SOURCE_DIR}")

 #-------------------------------------------------------------------------------
 # CUDA
@@ -20,6 +21,8 @@ SET(CUDA_SEPARABLE_COMPILATION ON)
 set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} --maxrregcount=32)

 find_package(CUDNN REQUIRED)
+include_directories(${CUDNN_INCLUDE_DIR})
+

 # compile
 file(GLOB tkdnn_CUSRC "src/kernels/*.cu" "src/sorting.cu")
@@ -1,10 +1,45 @@
 # tkDNN
-tkDNN is a Deep Neural Network library built with cuDNN and tensorRT primitives, specifically thought to work on NVIDIA Jetson Boards. It has been tested on TK1(branch cudnn2), TX1, TX2, AGX Xavier and several discrete GPU.
+tkDNN is a Deep Neural Network library built with cuDNN and tensorRT primitives, specifically thought to work on NVIDIA Jetson Boards. It has been tested on TK1(branch cudnn2), TX1, TX2, AGX Xavier, Nano and several discrete GPUs.
 The main goal of this project is to exploit NVIDIA boards as much as possible to obtain the best inference performance. It does not allow training. 

-Accepted paper @ IRC 2020, will soon been published.
+
+If you use tkDNN in your research, please cite one of the following papers. For use in commercial solutions, write at gattifrancesco@hotmail.it and micaela.verucchi@unimore.it or refer to https://hipert.unimore.it/ .
+
+```
+Accepted paper @ IRC 2020, will soon be published.
 M. Verucchi, L. Bartoli, F. Bagni, F. Gatti, P. Burgio and M. Bertogna, "Real-Time clustering and LiDAR-camera fusion on embedded platforms for self-driving cars",  in proceedings in IEEE Robotic Computing (2020)

+Accepted paper @ ETFA 2020, will soon be published.
+M. Verucchi, G. Brilli, D. Sapienza, M. Verasani, M. Arena, F. Gatti, A. Capotondi, R. Cavicchioli, M. Bertogna, M. Solieri
+"A Systematic Assessment of Embedded Neural Networks for Object Detection", in IEEE International Conference on Emerging Technologies and Factory Automation (2020)
+```
+
+## Results
+Inference FPS of yolov4 with tkDNN, average of 1200 images with the same dimesion as the input size, on 
+  * RTX 2080Ti (CUDA 10.2, TensorRT 7.0.0, Cudnn 7.6.5);
+  * Xavier AGX, Jetpack 4.3 (CUDA 10.0, CUDNN 7.6.3, tensorrt 6.0.1 );
+  * Tx2, Jetpack 4.2 (CUDA 10.0, CUDNN 7.3.1, tensorrt 5.0.6 );
+  * Jetson Nano, Jetpack 4.4  (CUDA 10.2, CUDNN 8.0.0, tensorrt 7.1.0 ). 
+
+| Platform   | Network    | FP32, B=1 | FP32, B=4	| FP16, B=1 |	FP16, B=4 |	INT8, B=1 |	INT8, B=4 | 
+| :------:   | :-----:    | :-----:   | :-----:   | :-----:   |	:-----:   |	:-----:   |	:-----:   | 
+| RTX 2080Ti | yolo4  320 | 118,59	  |237,31	    | 207,81	  | 443,32	  | 262,37	  | 530,93    | 
+| RTX 2080Ti | yolo4  416 | 104,81	  |162,86	    | 169,06	  | 293,78	  | 206,93	  | 353,26    | 
+| RTX 2080Ti | yolo4  512 | 92,98	    |132,43	    | 140,36	  | 215,17	  | 165,35	  | 254,96    | 
+| RTX 2080Ti | yolo4  608 | 63,77	    |81,53	    | 111,39	  | 152,89	  | 127,79	  | 184,72    | 
+| AGX Xavier | yolo4 320  |	26,78	    |32,05	    | 57,14	    | 79,05	    | 73,15	    | 97,56     |
+| AGX Xavier | yolo4 416  |	19,96	    |21,52	    | 41,01	    | 49,00	    | 50,81	    | 60,61     |
+| AGX Xavier | yolo4 512  |	16,58	    |16,98	    | 31,12	    | 33,84	    | 37,82	    | 41,28     |
+| AGX Xavier | yolo4 608  |	9,45 	    |10,13	    | 21,92	    | 23,36	    | 27,05	    | 28,93     |
+| Tx2        | yolo4 320	| 11,18	    | 12,07	    | 15,32	    | 16,31     | -         | -         |
+| Tx2        | yolo4 416	| 7,30	    | 7,58	    | 9,45	    | 9,90      | -         | -         |
+| Tx2        | yolo4 512	| 5,96	    | 5,95	    | 7,22	    | 7,23      | -         | -         |
+| Tx2        | yolo4 608	| 3,63	    | 3,65	    | 4,67	    | 4,70      | -         | -         |
+| Nano       | yolo4 320	| 4,23	    | 4,55	    | 6,14	    | 6,53      | -         | -         |
+| Nano       | yolo4 416	| 2,88	    | 3,00	    | 3,90	    | 4,04      | -         | -         |
+| Nano       | yolo4 512	| 2,32	    | 2,34	    | 3,02	    | 3,04      | -         | -         |
+| Nano       | yolo4 608	| 1,40	    | 1,41	    | 1,92	    | 1,93      | -         | -         |
+
 ## Index
 - [tkDNN](#tkdnn)
  - [Index](#index)
@@ -140,15 +175,25 @@ All models from darknet are now parsed directly from cfg, you still need to expo
  mish
 </details>

-## Run the demo
+## Run the demo 
+This is an example using yolov4.

-To run the an object detection demo follow these steps (example with yolov3):
+To run the an object detection first create the .rt file by running:
 ```
-rm yolo3_fp32.rt        # be sure to delete(or move) old tensorRT files
-./test_yolo3            # run the yolo test (is slow)
-./demo yolo3_fp32.rt ../demo/yolo_test.mp4 y
+rm yolo4_fp32.rt        # be sure to delete(or move) old tensorRT files
+./test_yolo4            # run the yolo test (is slow)
 ```
-In general the demo program takes 4 parameters:
+If you get problems in the creation, try to check the error activating the debug of TensorRT in this way:
+```
+cmake .. -DDEBUG=True
+make
+```
+
+Once you have succesfully created your rt file, run the demo: 
+```
+./demo yolo4_fp32.rt ../demo/yolo_test.mp4 y
+```
+In general the demo program takes 6 parameters:
 ```
 ./demo <network-rt-file> <path-to-video> <kind-of-network> <number-of-classes> <n-batches> <show-flag>
 ```
@@ -162,6 +207,7 @@ where

 N.b. By default it is used FP32 inference

+
 ![demo](https://user-images.githubusercontent.com/11562617/72547657-540e7800-388d-11ea-83c6-49dfea2a0607.gif)

 ### FP16 inference
@@ -177,24 +223,31 @@ N.b. Using FP16 inference will lead to some errors in the results (first or seco

 ### INT8 inference

-To run the an object detection demo with INT8 inference follow these steps (example with yolov3):
+To run the an object detection demo with INT8 inference three environment variables need to be set:
+  * ```export TKDNN_MODE=INT8```: set the 8-bit integer optimization
+  * ```export TKDNN_CALIB_IMG_PATH=/path/to/calibration/image_list.txt``` : image_list.txt has in each line the absolute path to a calibration image
+  * ```export TKDNN_CALIB_LABEL_PATH=/path/to/calibration/label_list.txt```: label_list.txt has in each line the absolute path to a calibration label
+  
+You should provide image_list.txt and label_list.txt, using training images. However, if you want to quickly test the INT8 inference you can run (from this repo root folder)
 ```
-export TKDNN_MODE=INT8  # set the 8-bit integer optimization
+bash scripts/download_validation.sh COCO
+```
+to automatically download COCO2017 validation (inside demo folder) and create those needed file. Use BDD insted of COCO to download BDD validation. 

-# image_list.txt contains the list of the absolute paths to the calibration images
-export TKDNN_CALIB_IMG_PATH=/path/to/calibration/image_list.txt
-
-# label_list.txt contains the list of the absolute paths to the calibration labels
-export TKDNN_CALIB_LABEL_PATH=/path/to/calibration/label_list.txt
+Then a complete example using yolo3 and COCO dataset would be:
+```
+export TKDNN_MODE=INT8
+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
 ```
-N.b. Using INT8 inference will lead to some errors in the results. 
-
-N.b. The test will be slower: this is due to the INT8 calibration, which may take some time to complete. 
-
-N.b. INT8 calibration requires TensorRT version greater than or equal to 6.0
+N.B. 
+ * Using INT8 inference will lead to some errors in the results. 
+ * The test will be slower: this is due to the INT8 calibration, which may take some time to complete. 
+ * INT8 calibration requires TensorRT version greater than or equal to 6.0
+ * Only 100 images are used to create the calibration table by default (set in the code).

 ### BatchSize bigger than 1
 ```
@@ -275,6 +328,8 @@ This demo also creates a json file named ```net_name_COCO_res.json``` containing
 | resnet101_cnet    | Centernet (Resnet101 backend)<sup>4</sup>     | [COCO 2017](http://cocodataset.org/)                          | 80        | 512x512       | [weights](https://cloud.hipert.unimore.it/s/5BTjHMWBcJk8g3i/download)     |
 | csresnext50-panet-spp    | Cross Stage Partial Network <sup>7</sup>     | [COCO 2014](http://cocodataset.org/)                          | 80        | 416x416       | [weights](https://cloud.hipert.unimore.it/s/Kcs4xBozwY4wFx8/download)     |
 | yolo4             | Yolov4 <sup>8</sup>                           | [COCO 2017](http://cocodataset.org/)                          | 80        | 416x416       | [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)     |
+| yolo4tiny             | Yolov4 tiny                           | [COCO 2017](http://cocodataset.org/)                          | 80        | 416x416       | [weights](https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/download)     |


 ## References
@@ -1,33 +1,66 @@
-# Find the header files
+# find the library
+if(CUDA_FOUND)
+  find_cuda_helper_libs(cudnn)
+  set(CUDNN_LIBRARY ${CUDA_cudnn_LIBRARY} CACHE FILEPATH "location of the cuDNN library")
+  unset(CUDA_cudnn_LIBRARY CACHE)

-find_path(CUDNN_INCLUDE_DIR
-        ${CMAKE_SYSROOT}/usr/local/include
-        ${CMAKE_SYSROOT}/usr/include
-        /usr/local/nvidia/tensorrt/include/
-        NO_DEFAULT_PATH
-)
+  find_cuda_helper_libs(nvinfer)
+  set(NVINFER_LIBRARY ${CUDA_nvinfer_LIBRARY} CACHE FILEPATH "location of the nvinfer library")
+  unset(CUDA_nvinfer_LIBRARY CACHE)
+endif()

-set(OLD_ROOT ${CMAKE_FIND_ROOT_PATH})
-list(APPEND CMAKE_FIND_ROOT_PATH /)
-list(APPEND CMAKE_FIND_LIBRARY_SUFFIXES .so.7)
-list(APPEND CMAKE_FIND_LIBRARY_SUFFIXES .so.5)
-find_library(CUDNN_LIB
-    NAMES cudnn
-    PATHS
-    /usr/local/driveworks/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib
-    /usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/
+# find the include
+if(CUDNN_LIBRARY)
+  find_path(CUDNN_INCLUDE_DIR
+    cudnn.h
+    PATHS ${CUDA_TOOLKIT_INCLUDE}
+    DOC "location of cudnn.h"
    NO_DEFAULT_PATH
-)
-find_library(CUDNN_NVLIB
-    NAMES "nvinfer"
-    PATHS
-    /usr/local/driveworks/targets/${CMAKE_SYSTEM_PROCESSOR}-Linux/lib
-    /usr/lib/${CMAKE_SYSTEM_PROCESSOR}-linux-gnu/
-    NO_DEFAULT_PATH
-)
-set(CMAKE_FIND_ROOT_PATH ${OLD_ROOT})
+  )

-set(CUDNN_LIBRARIES ${CUDNN_LIB} ${CUDNN_NVLIB})
-message("-- Found CUDNN: "  ${CUDNN_LIB})
-message("-- Found NVINFER: "  ${CUDNN_NVLIB})
-set(CUDNN_FOUND true)
+  if(NOT CUDNN_INCLUDE_DIR)
+    find_path(CUDNN_INCLUDE_DIR
+      cudnn.h
+      DOC "location of cudnn.h"
+    )
+  endif()
+
+  message("-- Found CUDNN: " ${CUDNN_LIBRARY})
+  message("-- Found CUDNN include: " ${CUDNN_INCLUDE_DIR})
+endif()
+
+if(NVINFER_LIBRARY)
+  find_path(NVINFER_INCLUDE_DIR
+    NvInfer.h
+    PATHS ${CUDA_TOOLKIT_INCLUDE}
+    DOC "location of NvInfer.h"
+    NO_DEFAULT_PATH
+  )
+
+  if(NOT NVINFER_INCLUDE_DIR)
+    find_path(NVINFER_INCLUDE_DIR
+        NvInfer.h
+        DOC "location of NvInfer.h"
+    )
+  endif()
+
+  message("-- Found NVINFER: " ${NVINFER_LIBRARY})
+  message("-- Found NVINFER include: " ${NVINFER_INCLUDE_DIR})
+endif()
+
+
+include(FindPackageHandleStandardArgs)
+find_package_handle_standard_args(CUDNN
+  FOUND_VAR CUDNN_FOUND
+  REQUIRED_VARS
+  CUDNN_LIBRARY
+  CUDNN_INCLUDE_DIR
+  VERSION_VAR CUDNN_VERSION
+)
+
+if(CUDNN_FOUND)
+  set(CUDNN_LIBRARIES ${CUDNN_LIBRARY} ${NVINFER_LIBRARY})
+  set(CUDNN_INCLUDE_DIRS ${CUDNN_INCLUDE_DIR} ${NVINFER_INCLUDE_DIR})
+endif()
+
+set(CUDNN_FOUND true)
@@ -153,7 +153,7 @@ int main(int argc, char *argv[])

        std::ofstream myfile;
        if(write_dets)
-            myfile.open ("det/"+f.lFilename.substr(f.lFilename.find("000")));
+            myfile.open ("det/"+f.lFilename.substr(f.lFilename.find("labels/") + 7));

        // save detections labels
        for(auto d:detected_bbox){
@@ -169,7 +169,7 @@ int main(int argc, char *argv[])
            f.det.push_back(b);

            if(write_dets)
-                myfile << d.cl << " "<< d.prob << " "<< d.x << " "<< d.y << " "<< d.w << " "<< d.h <<"\n";
+                myfile << d.cl << " "<< d.prob << " "<< b.x << " "<< b.y << " "<< b.w << " "<< b.h <<"\n";

 			if(show)// draw rectangle for detection
                cv::rectangle(batch_frames[0], cv::Point(d.x, d.y), cv::Point(d.x + d.w, d.y + d.h), cv::Scalar(0, 0, 255), 2);             
@@ -0,0 +1,7 @@
+FROM ceccocats/tkdnn:latest
+LABEL maintainer "Francesco Gatti"
+
+RUN cd && git clone https://github.com/ceccocats/tkDNN.git && cd tkDNN && mkdir build && cd build \ 
+    && cmake .. && make -j12
+
+
@@ -0,0 +1,57 @@
+FROM nvidia/cuda:10.2-cudnn7-devel-ubuntu18.04
+LABEL maintainer "Francesco Gatti"
+
+ADD nv-tensorrt-repo-ubuntu1804-cuda10.2-trt7.0.0.11-ga-20191216_1-1_amd64.deb /tmp/trt.deb
+RUN apt-get update && dpkg -i /tmp/trt.deb && rm /tmp/trt.deb && apt-get update
+RUN apt install -y libnvinfer7=7.0.0-1+cuda10.2 libnvinfer-dev=7.0.0-1+cuda10.2
+RUN DEBIAN_FRONTEND=noninteractive apt install -y git wget libeigen3-dev libyaml-cpp-dev
+RUN cd /tmp && \
+    wget https://github.com/Kitware/CMake/releases/download/v3.17.3/cmake-3.17.3-Linux-x86_64.sh && \
+    chmod +x cmake-3.17.3-Linux-x86_64.sh && \
+    ./cmake-3.17.3-Linux-x86_64.sh --prefix=/usr/local --exclude-subdir --skip-license && \
+    rm ./cmake-3.17.3-Linux-x86_64.sh
+
+RUN echo "INSTALL OPENCV"
+RUN apt-get install -y build-essential \
+    unzip \
+    pkg-config \
+    libjpeg-dev \
+    libpng-dev \
+    libtiff-dev \
+    libavcodec-dev \
+    libavformat-dev \
+    libswscale-dev \
+    libv4l-dev \
+    libxvidcore-dev \
+    libx264-dev \
+    libgtk-3-dev \
+    libatlas-base-dev \
+    gfortran \
+    libgstreamer1.0-dev \
+    libgstreamer-plugins-base1.0-dev \
+    libdc1394-22-dev \
+    libavresample-dev
+RUN cd && wget https://github.com/opencv/opencv/archive/4.3.0.tar.gz && tar -xf 4.3.0.tar.gz && rm *.tar.gz
+RUN cd && wget https://github.com/opencv/opencv_contrib/archive/4.3.0.tar.gz && tar -xf 4.3.0.tar.gz && rm *.tar.gz
+RUN cd && \
+    cd opencv-4.3.0 && mkdir build && cd build && \
+    cmake -D CMAKE_BUILD_TYPE=RELEASE \
+    -D CMAKE_INSTALL_PREFIX=/usr/local \
+    -D INSTALL_PYTHON_EXAMPLES=OFF \
+    -D INSTALL_C_EXAMPLES=OFF \
+    -D OPENCV_EXTRA_MODULES_PATH='~/opencv_contrib-4.3.0/modules' \
+    -D BUILD_EXAMPLES=OFF \
+    -D WITH_CUDA=ON \
+    -D CUDA_ARCH_BIN=7.2 \
+    -D CUDA_ARCH_PTX="" \
+    -D ENABLE_FAST_MATH=ON \
+    -D CUDA_FAST_MATH=ON \
+    -D WITH_CUBLAS=ON \
+    -D WITH_LIBV4L=ON \
+    -D WITH_GSTREAMER=ON \
+    -D WITH_GSTREAMER_0_10=OFF \
+    -D WITH_TBB=ON \
+    ../ && make -j12 && make install
+RUN apt clean
+
+
@@ -0,0 +1,21 @@
+# Use the prebuilt image
+```
+# build image
+docker build -t tkdnn:build -f Dockerfile . 
+```
+
+# Build Base Docker image
+```
+# make nvidia docker working
+# follow this guide: https://github.com/NVIDIA/nvidia-docker
+
+# dowload tensorrt
+# from: https://developer.nvidia.com/compute/machine-learning/tensorrt/secure/7.0/7.0.0.11/local_repo/nv-tensorrt-repo-ubuntu1804-cuda10.2-trt7.0.0.11-ga-20191216_1-1_amd64.deb
+
+# build image
+docker build -t ceccocats/tkdnn:latest -f Dockerfile.base .
+
+# run image
+docker run -ti --gpus all --rm ceccocats/tkdnn:latest bash
+```
+
@@ -1,6 +1,6 @@
 #pragma once
 #include <iostream>
-#include "tkdnn.h"
+#include "tkDNN/tkdnn.h"

 namespace tk { namespace dnn {

@@ -11,6 +11,7 @@ namespace tk { namespace dnn {
        int channels = 3;
        int batch_normalize=0;
        int groups = 1;
+        int group_id = 0;
        int filters=1;
        int size_x=1;
        int size_y=1;
@@ -27,267 +28,21 @@ namespace tk { namespace dnn {
        std::vector<int> layers;
        std::string activation = "linear";

+        friend std::ostream& operator<<(std::ostream& os, const darknetFields_t& f){
+            os << f.width << " " << f.height << " " << f.channels << " " << f.batch_normalize<< " " << f.filters << " "  << f.activation<< " " << f.scale_xy;
+            return os;
+        }
    };

-    std::ostream& operator<<(std::ostream& os, const darknetFields_t& f){
-        os << f.width << " " << f.height << " " << f.channels << " " << f.batch_normalize<< " " << f.filters << " "  << f.activation<< " " << f.scale_xy;
-        return os;
-    }
-
-    std::string darknetParseType(const std::string& line){
-            size_t start = line.find("[");
-            size_t end  = line.find("]");
-            if( start == std::string::npos || end == std::string::npos)
-                return "";
-            start++;
-            std::string type = line.substr(start, end-start);
-            return type;
-    }
-
-    bool divideNameAndValue(const std::string& line, std::string&name, std::string& value){
-        size_t sep = line.find("=");
-        if(sep == std::string::npos)
-            return false;
-        
-        name   = line.substr(0, sep);   
-        value  = line.substr(sep+1, line.size() - (sep+1));   
-        return true;
-    }
-
-    std::vector<int> fromStringToIntVec(const std::string& line, const char delimiter){
-        std::stringstream linestream(line);
-        std::string value;
-        std::vector<int> values;
-
-        while(getline(linestream,value,delimiter))
-            values.push_back(std::stoi(value));
-        return values;
-    }
-
-    bool darknetParseFields(const std::string& line, darknetFields_t& fields){
-
-        std::string name,value;
-        if(!divideNameAndValue(line, name, value))
-            return false;
-        if(name.find("width") !=  std::string::npos)
-            fields.width = std::stoi(value);
-        else if(name.find("height") !=  std::string::npos)
-            fields.height = std::stoi(value);
-        else if(name.find("channels") !=  std::string::npos)
-            fields.channels = std::stoi(value);
-        else if(name.find("batch_normalize") !=  std::string::npos)
-            fields.batch_normalize = std::stoi(value);
-        else if(name.find("filters") !=  std::string::npos)
-            fields.filters = std::stoi(value);
-        else if(name.find("activation") !=  std::string::npos)
-            fields.activation = value;
-        else if(name.find("size") !=  std::string::npos){
-            fields.size_x = std::stoi(value);
-            fields.size_y = std::stoi(value);
-        }
-        else if(name.find("size_x") !=  std::string::npos)
-            fields.size_x = std::stoi(value);
-        else if(name.find("size_y") !=  std::string::npos)
-            fields.size_y = std::stoi(value);
-        else if(name.find("stride") !=  std::string::npos){
-            fields.stride_x = std::stoi(value);
-            fields.stride_y = std::stoi(value);
-        }
-        else if(name.find("stride_x") !=  std::string::npos)
-            fields.stride_x = std::stoi(value);
-        else if(name.find("stride_y") !=  std::string::npos)
-            fields.stride_y = std::stoi(value);
-        else if(name.find("pad") !=  std::string::npos)
-            fields.pad = std::stoi(value);
-        else if(name.find("classes") !=  std::string::npos)
-            fields.classes = std::stoi(value);
-        else if(name.find("num") !=  std::string::npos)
-            fields.num = std::stoi(value);
-        else if(name.find("coords") !=  std::string::npos)
-            fields.coords = std::stoi(value);
-        else if(name.find("groups") !=  std::string::npos)
-            fields.groups = std::stoi(value);
-        else if(name.find("scale_x_y") !=  std::string::npos)
-            fields.scale_xy = std::stof(value);
-        else if(name.find("from") !=  std::string::npos)
-            fields.layers.push_back(std::stof(value));
-        else if(name.find("mask") !=  std::string::npos){
-            auto vec = fromStringToIntVec(value, ',');
-            fields.n_mask = vec.size();
-        }
-        else if(name.find("layers") !=  std::string::npos)
-            fields.layers = fromStringToIntVec(value, ',');
-
-        else
-            std::cout<<"Not supported field: "<<line<<std::endl;
-        return true;
-    }
-
-    tk::dnn::Network *darknetAddNet(darknetFields_t &fields) {
-        //std::cout<<"Add Net: "<<fields.type<<"\n";
-        dataDim_t dim(1, fields.channels, fields.height, fields.width);
-        return new tk::dnn::Network(dim);
-    }
-
-
-    void darknetAddLayer(tk::dnn::Network *net, darknetFields_t &f, std::string wgs_path, std::vector<tk::dnn::Layer*> &netLayers, const std::vector<std::string>& names) {
-        if(net == nullptr)
-            FatalError("Cant add a layer without a Net\n");
-
-        // padding compute
-        if(f.pad == 1) {
-            f.padding_x = f.padding_y = f.size_x /2;
-        }
-        //std::cout<<"Add layer: "<<f.type<<"\n";
-        if(f.type == "convolutional") {
-            std::string wgs = wgs_path + "/c" + std::to_string(netLayers.size()) + ".bin";
-            //printf("%d (%d,%d) (%d,%d) (%d,%d) %s %d %d\n", f.filters, f.size_x, f.size_y, f.stride_x, f.stride_y, f.padding_x, f.padding_y, wgs.c_str(), f.batch_normalize, f.groups);
-            tk::dnn::Conv2d *l= new tk::dnn::Conv2d(net, f.filters, f.size_x, f.size_y, f.stride_x, 
-                                f.stride_y, f.padding_x, f.padding_y, wgs, f.batch_normalize, false, f.groups);
-            netLayers.push_back(l);
-        } else if(f.type == "maxpool") {
-            if(f.stride_x == 1 && f.stride_y == 1)
-                netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y, 
-                    f.padding_x, f.padding_y, tk::dnn::POOLING_MAX_FIXEDSIZE));
-            else
-                netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y, 
-                    f.padding_x, f.padding_y, tk::dnn::POOLING_MAX));
-
-        } else if(f.type == "avgpool") {
-            netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y, 
-                f.padding_x, f.padding_y, tk::dnn::POOLING_AVERAGE));
-
-        } else if(f.type == "shortcut") {
-            if(f.layers.size() != 1) FatalError("no layers to shortcut\n");
-            int layerIdx = f.layers[0];
-            if(layerIdx < 0)
-                layerIdx = netLayers.size() + layerIdx; 
-            if(layerIdx < 0 || layerIdx >= netLayers.size()) FatalError("impossible to shortcut\n");
-            //std::cout<<"shortcut to "<<layerIdx<<" "<<netLayers[layerIdx]->getLayerName()<<"\n";
-            netLayers.push_back(new tk::dnn::Shortcut(net, netLayers[layerIdx]));
-
-        } else if(f.type == "upsample") {
-            netLayers.push_back(new tk::dnn::Upsample(net, f.stride_x));
-
-        } else if(f.type == "route") {
-            if(f.layers.size() == 0) FatalError("no layers to Route\n");
-            std::vector<tk::dnn::Layer*> layers;
-            for(int i=0; i<f.layers.size(); i++) {
-                int layerIdx = f.layers[i];
-                if(layerIdx < 0)
-                    layerIdx = netLayers.size() + layerIdx; 
-                if(layerIdx < 0 || layerIdx >= netLayers.size()) FatalError("impossible to route\n");
-                //std::cout<<"Route to "<<layerIdx<<" "<<netLayers[layerIdx]->getLayerName()<<"\n";
-                layers.push_back(netLayers[layerIdx]);
-            }
-            netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size()));
-
-        } else if(f.type == "reorg") {
-            netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x));
-
-        } else if(f.type == "region") {
-            netLayers.push_back(new tk::dnn::Region(net, f.classes, f.coords, f.num));
-
-        } else if(f.type == "yolo") {
-            std::string wgs = wgs_path + "/g" + std::to_string(netLayers.size()) + ".bin";
-            //printf("%d %d %s %d %f\n", f.classes, f.num/f.n_mask, wgs.c_str(), f.n_mask, f.scale_xy);
-            tk::dnn::Yolo *l = new tk::dnn::Yolo(net, f.classes, f.num/f.n_mask, wgs, f.n_mask, f.scale_xy);
-            if(names.size() != f.classes)
-                FatalError("Mismatch between number of classes and names");
-            l->classesNames = names;
-            netLayers.push_back(l);
-
-        } else{
-            FatalError("layer not supported: " + f.type);
-        }
-
-        // add activation
-        if(netLayers.size() > 0 && f.activation != "linear") {
-            tkdnnActivationMode_t act;
-            if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU);
-            else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY;
-            else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH;
-            else { FatalError("activation not supported: " + f.activation); }
-            netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act);
-        };
-    }
-
-    std::vector<std::string> darknetReadNames(const std::string& names_file){
-        std::ifstream if_names(names_file);
-        if(!if_names.is_open())
-            FatalError("cloud not open names file: " + names_file);       
-
-        std::vector<std::string> names;
-        std::string line;
-        while(std::getline(if_names, line))
-            if(line != "")
-                names.push_back(line);
-        
-        if_names.close();
-        return names;
-    }
-
-    tk::dnn::Network* darknetParser(const std::string& cfg_file, const std::string& wgs_path, const std::string& names_file) {
-
-        tk::dnn::Network *net = nullptr;
-        
-        // layers without activations to retrive correct id number
-        std::vector<tk::dnn::Layer*> netLayers;
-
-        std::ifstream if_cfg(cfg_file);
-        if(!if_cfg.is_open())
-            FatalError("cloud not open cfg file: " + cfg_file);
-
-        std::vector<std::string> names = darknetReadNames(names_file);
-
-        darknetFields_t fields; // will be filled with layers fields
-        std::string line;
-        while(std::getline(if_cfg, line)) {
-            // remove comments
-            std::size_t found = line.find("#");
-            if ( found != std::string::npos ) {
-                line = line.substr(0, found);
-            }
-
-            // skip empty lines
-            if(line.size() == 0)
-                continue;
-            
-            std::string type = darknetParseType(line);
-            if(type.size() > 0) {
-                // end of filled type
-                if(fields.type != "") {
-                    if(fields.type == "net")
-                        net = darknetAddNet(fields);
-                    else
-                        darknetAddLayer(net, fields, wgs_path, netLayers, names);
-                }
-
-                // new type
-                //std::cout<<"type: "<<type<<"\n";
-                fields = darknetFields_t(); // reset to default
-                fields.type = type;
-                continue;
-            }
-
-            if(darknetParseFields(line, fields)) {
-                // already parsed do nothing
-            } else {
-                FatalError("could not parse line: " + line);
-            }
-        }
-
-        // end of filled type
-        if(fields.type != "") {
-            darknetAddLayer(net, fields, wgs_path, netLayers, names);
-        }
-
-        if(net == nullptr) {
-            FatalError("net not found\n");
-        }
-        return net;
-    }
-        
+    std::string darknetParseType(const std::string& line);
+    bool divideNameAndValue(const std::string& line, std::string&name, std::string& value);
+    std::vector<int> fromStringToIntVec(const std::string& line, const char delimiter);
    
+    bool darknetParseFields(const std::string& line, darknetFields_t& fields);
+    tk::dnn::Network *darknetAddNet(darknetFields_t &fields);
+    void darknetAddLayer(tk::dnn::Network *net, darknetFields_t &f, std::string wgs_path, 
+                         std::vector<tk::dnn::Layer*> &netLayers, const std::vector<std::string>& names);
+    std::vector<std::string> darknetReadNames(const std::string& names_file);
+    tk::dnn::Network* darknetParser(const std::string& cfg_file, const std::string& wgs_path, const std::string& names_file);

 }}
@@ -273,8 +273,8 @@ public:
 protected:
    cudnnFilterDescriptor_t filterDesc;
    cudnnConvolutionDescriptor_t convDesc;
-    cudnnConvolutionFwdAlgo_t     algo;
-    cudnnConvolutionBwdDataAlgo_t bwAlgo;
+    cudnnConvolutionFwdAlgoPerf_t     algo;
+    cudnnConvolutionBwdDataAlgoPerf_t bwAlgo;
    cudnnTensorDescriptor_t biasTensorDesc;

    void initCUDNN(bool back = false);
@@ -509,7 +509,7 @@ public:
 class Route : public Layer {

 public:
-    Route(Network *net, Layer **layers, int layers_n); 
+    Route(Network *net, Layer **layers, int layers_n, int groups = 1, int group_id = 0); 
    virtual ~Route();
    virtual layerType_t getLayerType() { return LAYER_ROUTE; };

@@ -519,6 +519,8 @@ public:
    static const int MAX_LAYERS = 32;
    Layer *layers[MAX_LAYERS];  //ids of layers to be merged
    int layers_n; //number of layers
+    int groups;
+    int group_id;
 };


@@ -28,7 +28,7 @@ using namespace nvinfer1;
 #include "pluginsRT/ActivationMishRT.h"
 #include "pluginsRT/ReorgRT.h"
 #include "pluginsRT/RegionRT.h"
-//#include "pluginsRT/RouteRT.h"
+#include "pluginsRT/RouteRT.h"
 #include "pluginsRT/ShortcutRT.h"
 #include "pluginsRT/YoloRT.h"
 #include "pluginsRT/UpsampleRT.h"
@@ -0,0 +1,12 @@
+#pragma once
+#include <iostream>
+#include <opencv2/core/types.hpp>
+#include "tkdnn.h"
+
+namespace tk { namespace dnn {
+
+cv::Mat vizFloat2colorMap(cv::Mat map);
+cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim);
+cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim = 1000);
+    
+}}
@@ -13,6 +13,7 @@ private:
    int num = 0;
    int nMasks = 0;
    int nDets = 0;
+    bool letterbox = false;
    tk::dnn::Yolo::detection *dets = nullptr;
    tk::dnn::Yolo* yolo[3];

@@ -21,7 +22,7 @@ private:
    cv::Mat bgr_h;
    
 public:
-    Yolo3Detection() {};
+    Yolo3Detection(const bool letter_box=false) :letterbox(letter_box){}
    ~Yolo3Detection() {}; 

    bool init(const std::string& tensor_path, const int n_classes=80, const int n_batches=1);
@@ -8,7 +8,9 @@ class RouteRT : public IPlugin {
 	*/

 public:
-	RouteRT() {
+	RouteRT(int groups, int group_id) {
+		this->groups = groups;
+		this->group_id = group_id;
 	}

 	~RouteRT(){
@@ -22,7 +24,7 @@ public:
 	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
 		int out_c = 0;
 		for(int i=0; i<nbInputDims; i++) out_c += inputs[i].d[0];
-		return DimsCHW{out_c, inputs[0].d[1], inputs[0].d[2]};
+		return DimsCHW{out_c/groups, inputs[0].d[1], inputs[0].d[2]};
 	}

 	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
@@ -34,6 +36,7 @@ public:
 		}
 		h = inputDims[0].d[1];
 		w = inputDims[0].d[2];
+		c /= groups;
 	}

 	int initialize() override {
@@ -49,15 +52,18 @@ public:
 	}

 	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
-
+		
 		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);

-		int offset = 0;
-		for(int i=0; i<in; i++) {
-			dnnType *input = (dnnType*)reinterpret_cast<const dnnType*>(inputs[i]);
-			int in_dim = c_in[i]*h*w;
-			checkCuda( cudaMemcpyAsync(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) );
-			offset += in_dim;
+		for(int b=0; b<batchSize; b++) {
+			int offset = 0;
+			for(int i=0; i<in; i++) {
+				dnnType *input = (dnnType*)reinterpret_cast<const dnnType*>(inputs[i]);
+				int in_dim = c_in[i]*h*w;
+				int part_in_dim = in_dim / this->groups;
+				checkCuda( cudaMemcpyAsync(dstData + b*c*w*h + offset, input + b*c*w*h*groups + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream) );
+				offset += part_in_dim;
+			}
 		}

 		return 0;
@@ -65,11 +71,13 @@ public:


 	virtual size_t getSerializationSize() override {
-		return (4+MAX_INPUTS)*sizeof(int);
+		return (6+MAX_INPUTS)*sizeof(int);
 	}

 	virtual void serialize(void* buffer) override {
 		char *buf = reinterpret_cast<char*>(buffer);
+		tk::dnn::writeBUF(buf, groups);
+		tk::dnn::writeBUF(buf, group_id);
 		tk::dnn::writeBUF(buf, in);
 		for(int i=0; i<MAX_INPUTS; i++)
 			tk::dnn::writeBUF(buf, c_in[i]);
@@ -83,4 +91,5 @@ public:
 	int in;
 	int c_in[MAX_INPUTS];
 	int c, h, w;
+	int groups, group_id;
 };
@@ -1,7 +1,7 @@

 #include <tkdnn.h>
 int testInference(std::vector<std::string> input_bins, std::vector<std::string> output_bins, 
-  tk::dnn::Network *net, tk::dnn::NetworkRT *netRT = nullptr) {
+    tk::dnn::Network *net, tk::dnn::NetworkRT *netRT = nullptr) {

    std::vector<tk::dnn::Layer*> outputs;
    for(int i=0; i<net->num_layers; i++) {
@@ -67,7 +67,11 @@ int testInference(std::vector<std::string> input_bins, std::vector<std::string>
            std::cout<<"CUDNN vs TRT    "; 
            ret_cudnn_tensorrt |= checkResult(odim, cudnn_out[i], rt_out[i]) == 0 ? 0 : ERROR_CUDNNvsTENSORRT;
        }
-    }
-    return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt;

-  }
+        delete [] out_h;
+        checkCuda( cudaFree(out) );
+    }
+    delete [] input_h;
+    checkCuda( cudaFree(data) );
+    return ret_cudnn | ret_tensorrt | ret_cudnn_tensorrt;
+}
@@ -118,4 +118,8 @@ void matrixMulAdd(  cublasHandle_t handle, dnnType* srcData, dnnType* dstData,
 void getMemUsage(double& vm_usage_kb, double& resident_set_kb);
 void printCudaMemUsage();
 void removePathAndExtension(const std::string &full_string, std::string &name);
+static inline bool isCudaPointer(void *data) {
+  cudaPointerAttributes attr;
+  return cudaPointerGetAttributes(&attr, data) == 0;
+}
 #endif //UTILS_H
@@ -62,5 +62,5 @@ make -j4
 sudo make install
 sudo ldconfig

-cd '~/Downloads/opencv4/lib/python3.6/site-packages'
+cd ~/Downloads/opencv4/lib/python3.6/site-packages
 ln -s /usr/local/lib/python3.6/site-packages/cv2.cpython-36m-aarch64-linux-gnu.so cv2.so
@@ -74,6 +74,7 @@ do

    test_net yolo4
    test_net yolo4_berkeley
+    test_net yolo4tiny
    test_net yolo3
    test_net yolo3_berkeley
    test_net yolo3_coco4
@@ -62,25 +62,30 @@ void Conv2d::initCUDNN(bool back) {
    // init workspace
    workSpace = NULL;
    ws_sizeInBytes = 0;
+    int algo_count = 0;
    if(back) {
-        checkCUDNN( cudnnGetConvolutionBackwardDataAlgorithm(net->cudnnHandle,
-                                                        filterDesc, dstTensor, convDesc, srcTensor,
-                                                        CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST, 0, &bwAlgo) );
+        checkCUDNN( cudnnGetConvolutionBackwardDataAlgorithm_v7(net->cudnnHandle,
+                                                                filterDesc, dstTensor, convDesc, srcTensor, 1, &algo_count, &bwAlgo) );
        checkCUDNN(cudnnGetConvolutionBackwardDataWorkspaceSize(net->cudnnHandle,
-                                                           filterDesc, dstTensor, convDesc, srcTensor,
-                                                           bwAlgo, &ws_sizeInBytes));
+                                                                filterDesc, dstTensor, convDesc, srcTensor,
+                                                                bwAlgo.algo, &ws_sizeInBytes));
+        

        // invert tensors
        srcTensorDesc = dstTensor;
        dstTensorDesc = srcTensor;
    } else {
-        checkCUDNN( cudnnGetConvolutionForwardAlgorithm(net->cudnnHandle,
-                                                        srcTensor, filterDesc, convDesc, dstTensor,
-                                                        CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo) );
-        checkCUDNN(cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle,
-                                                           srcTensor, filterDesc, convDesc, dstTensor,
-                                                           algo, &ws_sizeInBytes));
+        
+        checkCUDNN( cudnnGetConvolutionForwardAlgorithm_v7(net->cudnnHandle,
+                                                         srcTensor, filterDesc, convDesc, dstTensor,
+                                                        1, &algo_count, &algo) );
+         checkCUDNN(cudnnGetConvolutionForwardWorkspaceSize(net->cudnnHandle,
+                                                            srcTensor, filterDesc, convDesc, dstTensor,
+                                                           algo.algo, &ws_sizeInBytes));
    }
+
+    if(algo_count < 1)
+        FatalError("Cannot retrieve convolutional algo");
 }

 void Conv2d::inferCUDNN(dnnType* srcData, bool back) {
@@ -91,12 +96,12 @@ void Conv2d::inferCUDNN(dnnType* srcData, bool back) {
        checkCUDNN(cudnnConvolutionBackwardData(net->cudnnHandle,
                                                &alpha, filterDesc, data_d,
                                                srcTensorDesc, srcData,
-                                                convDesc, bwAlgo, workSpace, ws_sizeInBytes,
+                                                convDesc, bwAlgo.algo, workSpace, ws_sizeInBytes,
                                                &beta, dstTensorDesc, dstData));
    } else {
        checkCUDNN(cudnnConvolutionForward(net->cudnnHandle,
                                           &alpha, srcTensorDesc, srcData, filterDesc,
-                                           data_d, convDesc, algo, workSpace, ws_sizeInBytes,
+                                           data_d, convDesc, algo.algo, workSpace, ws_sizeInBytes,
                                           &beta, dstTensorDesc, dstData));
    }

@@ -0,0 +1,263 @@
+#include "tkDNN/DarknetParser.h"
+
+namespace tk { namespace dnn {
+
+    std::string darknetParseType(const std::string& line){
+            size_t start = line.find("[");
+            size_t end  = line.find("]");
+            if( start == std::string::npos || end == std::string::npos)
+                return "";
+            start++;
+            std::string type = line.substr(start, end-start);
+            return type;
+    }
+
+    bool divideNameAndValue(const std::string& line, std::string&name, std::string& value){
+        size_t sep = line.find("=");
+        if(sep == std::string::npos)
+            return false;
+        
+        name   = line.substr(0, sep);   
+        value  = line.substr(sep+1, line.size() - (sep+1));   
+        return true;
+    }
+
+    std::vector<int> fromStringToIntVec(const std::string& line, const char delimiter){
+        std::stringstream linestream(line);
+        std::string value;
+        std::vector<int> values;
+
+        while(getline(linestream,value,delimiter))
+            values.push_back(std::stoi(value));
+        return values;
+    }
+
+    bool darknetParseFields(const std::string& line, darknetFields_t& fields){
+
+        std::string name,value;
+        if(!divideNameAndValue(line, name, value))
+            return false;
+        if(name.find("width") !=  std::string::npos)
+            fields.width = std::stoi(value);
+        else if(name.find("height") !=  std::string::npos)
+            fields.height = std::stoi(value);
+        else if(name.find("channels") !=  std::string::npos)
+            fields.channels = std::stoi(value);
+        else if(name.find("batch_normalize") !=  std::string::npos)
+            fields.batch_normalize = std::stoi(value);
+        else if(name.find("filters") !=  std::string::npos)
+            fields.filters = std::stoi(value);
+        else if(name.find("activation") !=  std::string::npos)
+            fields.activation = value;
+        else if(name.find("size") !=  std::string::npos){
+            fields.size_x = std::stoi(value);
+            fields.size_y = std::stoi(value);
+        }
+        else if(name.find("size_x") !=  std::string::npos)
+            fields.size_x = std::stoi(value);
+        else if(name.find("size_y") !=  std::string::npos)
+            fields.size_y = std::stoi(value);
+        else if(name.find("stride") !=  std::string::npos){
+            fields.stride_x = std::stoi(value);
+            fields.stride_y = std::stoi(value);
+        }
+        else if(name.find("stride_x") !=  std::string::npos)
+            fields.stride_x = std::stoi(value);
+        else if(name.find("stride_y") !=  std::string::npos)
+            fields.stride_y = std::stoi(value);
+        else if(name.find("pad") !=  std::string::npos)
+            fields.pad = std::stoi(value);
+        else if(name.find("classes") !=  std::string::npos)
+            fields.classes = std::stoi(value);
+        else if(name.find("num") !=  std::string::npos)
+            fields.num = std::stoi(value);
+        else if(name.find("coords") !=  std::string::npos)
+            fields.coords = std::stoi(value);
+        else if(name.find("groups") !=  std::string::npos)
+            fields.groups = std::stoi(value);
+        else if(name.find("group_id") !=  std::string::npos)
+            fields.group_id = std::stoi(value);
+        else if(name.find("scale_x_y") !=  std::string::npos)
+            fields.scale_xy = std::stof(value);
+        else if(name.find("from") !=  std::string::npos)
+            fields.layers.push_back(std::stof(value));
+        else if(name.find("mask") !=  std::string::npos){
+            auto vec = fromStringToIntVec(value, ',');
+            fields.n_mask = vec.size();
+        }
+        else if(name.find("layers") !=  std::string::npos)
+            fields.layers = fromStringToIntVec(value, ',');
+
+        else
+            std::cout<<"Not supported field: "<<line<<std::endl;
+        return true;
+    }
+
+    tk::dnn::Network *darknetAddNet(darknetFields_t &fields) {
+        //std::cout<<"Add Net: "<<fields.type<<"\n";
+        dataDim_t dim(1, fields.channels, fields.height, fields.width);
+        return new tk::dnn::Network(dim);
+    }
+
+
+    void darknetAddLayer(tk::dnn::Network *net, darknetFields_t &f, std::string wgs_path, std::vector<tk::dnn::Layer*> &netLayers, const std::vector<std::string>& names) {
+        if(net == nullptr)
+            FatalError("Cant add a layer without a Net\n");
+
+        // padding compute
+        if(f.pad == 1) {
+            f.padding_x = f.padding_y = f.size_x /2;
+        }
+        //std::cout<<"Add layer: "<<f.type<<"\n";
+        if(f.type == "convolutional") {
+            std::string wgs = wgs_path + "/c" + std::to_string(netLayers.size()) + ".bin";
+            //printf("%d (%d,%d) (%d,%d) (%d,%d) %s %d %d\n", f.filters, f.size_x, f.size_y, f.stride_x, f.stride_y, f.padding_x, f.padding_y, wgs.c_str(), f.batch_normalize, f.groups);
+            tk::dnn::Conv2d *l= new tk::dnn::Conv2d(net, f.filters, f.size_x, f.size_y, f.stride_x, 
+                                f.stride_y, f.padding_x, f.padding_y, wgs, f.batch_normalize, false, f.groups);
+            netLayers.push_back(l);
+        } else if(f.type == "maxpool") {
+            if(f.stride_x == 1 && f.stride_y == 1)
+                netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y, 
+                    f.padding_x, f.padding_y, tk::dnn::POOLING_MAX_FIXEDSIZE));
+            else
+                netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y, 
+                    f.padding_x, f.padding_y, tk::dnn::POOLING_MAX));
+
+        } else if(f.type == "avgpool") {
+            netLayers.push_back(new tk::dnn::Pooling(net, f.size_x, f.size_y, f.stride_x, f.stride_y, 
+                f.padding_x, f.padding_y, tk::dnn::POOLING_AVERAGE));
+
+        } else if(f.type == "shortcut") {
+            if(f.layers.size() != 1) FatalError("no layers to shortcut\n");
+            int layerIdx = f.layers[0];
+            if(layerIdx < 0)
+                layerIdx = netLayers.size() + layerIdx; 
+            if(layerIdx < 0 || layerIdx >= netLayers.size()) FatalError("impossible to shortcut\n");
+            //std::cout<<"shortcut to "<<layerIdx<<" "<<netLayers[layerIdx]->getLayerName()<<"\n";
+            netLayers.push_back(new tk::dnn::Shortcut(net, netLayers[layerIdx]));
+
+        } else if(f.type == "upsample") {
+            netLayers.push_back(new tk::dnn::Upsample(net, f.stride_x));
+
+        } else if(f.type == "route") {
+            if(f.layers.size() == 0) FatalError("no layers to Route\n");
+            std::vector<tk::dnn::Layer*> layers;
+            for(int i=0; i<f.layers.size(); i++) {
+                int layerIdx = f.layers[i];
+                if(layerIdx < 0)
+                    layerIdx = netLayers.size() + layerIdx; 
+                if(layerIdx < 0 || layerIdx >= netLayers.size()) FatalError("impossible to route\n");
+                //std::cout<<"Route to "<<layerIdx<<" "<<netLayers[layerIdx]->getLayerName()<<"\n";
+                layers.push_back(netLayers[layerIdx]);
+            }
+            netLayers.push_back(new tk::dnn::Route(net, layers.data(), layers.size(), f.groups, f.group_id));
+
+        } else if(f.type == "reorg") {
+            netLayers.push_back(new tk::dnn::Reorg(net, f.stride_x));
+
+        } else if(f.type == "region") {
+            netLayers.push_back(new tk::dnn::Region(net, f.classes, f.coords, f.num));
+
+        } else if(f.type == "yolo") {
+            std::string wgs = wgs_path + "/g" + std::to_string(netLayers.size()) + ".bin";
+            //printf("%d %d %s %d %f\n", f.classes, f.num/f.n_mask, wgs.c_str(), f.n_mask, f.scale_xy);
+            tk::dnn::Yolo *l = new tk::dnn::Yolo(net, f.classes, f.num/f.n_mask, wgs, f.n_mask, f.scale_xy);
+            if(names.size() != f.classes)
+                FatalError("Mismatch between number of classes and names");
+            l->classesNames = names;
+            netLayers.push_back(l);
+
+        } else{
+            FatalError("layer not supported: " + f.type);
+        }
+
+        // add activation
+        if(netLayers.size() > 0 && f.activation != "linear") {
+            tkdnnActivationMode_t act;
+            if(f.activation == "relu") act = tkdnnActivationMode_t(CUDNN_ACTIVATION_RELU);
+            else if(f.activation == "leaky") act = tk::dnn::ACTIVATION_LEAKY;
+            else if(f.activation == "mish") act = tk::dnn::ACTIVATION_MISH;
+            else { FatalError("activation not supported: " + f.activation); }
+            netLayers[netLayers.size()-1] = new tk::dnn::Activation(net, act);
+        };
+    }
+
+    std::vector<std::string> darknetReadNames(const std::string& names_file){
+        std::ifstream if_names(names_file);
+        if(!if_names.is_open())
+            FatalError("cloud not open names file: " + names_file);       
+
+        std::vector<std::string> names;
+        std::string line;
+        while(std::getline(if_names, line))
+            if(line != "")
+                names.push_back(line);
+        
+        if_names.close();
+        return names;
+    }
+
+    tk::dnn::Network* darknetParser(const std::string& cfg_file, const std::string& wgs_path, const std::string& names_file) {
+
+        tk::dnn::Network *net = nullptr;
+        
+        // layers without activations to retrive correct id number
+        std::vector<tk::dnn::Layer*> netLayers;
+
+        std::ifstream if_cfg(cfg_file);
+        if(!if_cfg.is_open())
+            FatalError("cloud not open cfg file: " + cfg_file);
+
+        std::vector<std::string> names = darknetReadNames(names_file);
+
+        darknetFields_t fields; // will be filled with layers fields
+        std::string line;
+        while(std::getline(if_cfg, line)) {
+            // remove comments
+            std::size_t found = line.find("#");
+            if ( found != std::string::npos ) {
+                line = line.substr(0, found);
+            }
+
+            // skip empty lines
+            if(line.size() == 0)
+                continue;
+            
+            std::string type = darknetParseType(line);
+            if(type.size() > 0) {
+                // end of filled type
+                if(fields.type != "") {
+                    if(fields.type == "net")
+                        net = darknetAddNet(fields);
+                    else
+                        darknetAddLayer(net, fields, wgs_path, netLayers, names);
+                }
+
+                // new type
+                //std::cout<<"type: "<<type<<"\n";
+                fields = darknetFields_t(); // reset to default
+                fields.type = type;
+                continue;
+            }
+
+            if(darknetParseFields(line, fields)) {
+                // already parsed do nothing
+            } else {
+                FatalError("could not parse line: " + line);
+            }
+        }
+
+        // end of filled type
+        if(fields.type != "") {
+            darknetAddLayer(net, fields, wgs_path, netLayers, names);
+        }
+
+        if(net == nullptr) {
+            FatalError("net not found\n");
+        }
+        return net;
+    }
+        
+    
+
+}}
@@ -132,21 +132,14 @@ void BatchStream::readCVimage(std::string inputFileName, std::vector<float>& res

 void BatchStream::readLabels(std::string inputFileName, std::vector<float>& ris) {
    std::ifstream is(inputFileName.c_str());
-    //read only the first number: the image sub-portion class
-    while (true) {
+    
+    std::string line;
+    while (std::getline(is, line))
+    {
+        std::istringstream iss(line);
        float val;
-		is >> val;
-        if (!is) {
-            break;
-        }
-        // insert the first number and skip all others
+        if(!(iss >> val)) { break; } // error
        ris.push_back(val);
-        while( true ) {
-            char c;
-            is >> c;
-            if (is.peek() == '\n') //detect "\n"
-                break;
-        }   
    }
 }

@@ -86,7 +86,11 @@ LSTM::LSTM( Network *net, int hiddensize, bool returnSeq, std::string fname_weig
    // RNN descriptors
    checkCUDNN(cudnnCreateRNNDescriptor(&rnnDesc));

-    checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle, 
+#if CUDNN_MAJOR > 7
+    checkCUDNN(cudnnSetRNNDescriptor_v6(net->cudnnHandle,
+#else
+    checkCUDNN(cudnnSetRNNDescriptor(net->cudnnHandle,
+#endif
        rnnDesc, stateSize, numLayers, dropoutDesc,
        cudnnRNNInputMode_t::CUDNN_LINEAR_INPUT,
        //(bidirectional ? cudnnDirectionMode_t::CUDNN_BIDIRECTIONAL : cudnnDirectionMode_t::CUDNN_UNIDIRECTIONAL),
@@ -449,12 +449,15 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
        // }
        // std::cout<<"\n";
    }
-    
-    IConcatenationLayer *lRT = networkRT->addConcatenation(tens, l->layers_n);
-    //IPlugin *plugin = new RouteRT();
-    //IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin);
-    checkNULL(lRT);

+    if(l->groups > 1){
+        IPlugin *plugin = new RouteRT(l->groups, l->group_id);
+        IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin);
+        checkNULL(lRT);
+        return lRT;
+    }
+    IConcatenationLayer *lRT = networkRT->addConcatenation(tens, l->layers_n);
+    checkNULL(lRT);
    return lRT;
 }

@@ -595,7 +598,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) {

 bool NetworkRT::serialize(const char *filename) {

-    std::ofstream p(filename);
+    std::ofstream p(filename, std::ios::binary);
    if (!p) {
        FatalError("could not open plan output file");
        return false;
@@ -766,9 +769,9 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
        r->w = readBUF<int>(buf);
        return r;
    } 
-/*
+
    if(name.find("Route") == 0) {
-        RouteRT *r = new RouteRT();
+        RouteRT *r = new RouteRT(readBUF<int>(buf),readBUF<int>(buf));
        r->in = readBUF<int>(buf);
        for(int i=0; i<RouteRT::MAX_INPUTS; i++)
            r->c_in[i] = readBUF<int>(buf);
@@ -777,7 +780,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa
        r->w = readBUF<int>(buf);
        return r;
    } 
-*/
+
    if(name.find("Deformable") == 0) {
        DeformableConvRT *r = new DeformableConvRT(readBUF<int>(buf), readBUF<int>(buf), readBUF<int>(buf),
                                                    readBUF<int>(buf), readBUF<int>(buf), readBUF<int>(buf),
@@ -0,0 +1,69 @@
+#include <opencv2/core/core.hpp>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/videoio.hpp>
+#include <opencv2/imgproc/imgproc.hpp>
+#include "tkDNN/NetworkViz.h"
+
+namespace tk { namespace dnn {
+
+cv::Mat vizFloat2colorMap(cv::Mat map) {
+
+    double min;
+    double max;
+    cv::minMaxIdx(map, &min, &max);
+    cv::Mat adjMap;
+    // expand your range to 0..255. Similar to histEq();
+    map.convertTo(adjMap,CV_8UC1, 255 / (max-min), -min); 
+    //return adjMap;
+
+    
+    cv::Mat falseColorsMap;
+    applyColorMap(adjMap, falseColorsMap, cv::COLORMAP_HOT);
+    return falseColorsMap;
+}
+
+cv::Mat vizData2Mat(dnnType *dataInput, tk::dnn::dataDim_t dim, int imgdim) {
+    dnnType *data = nullptr;
+
+    // copy to CPU
+    if(isCudaPointer(dataInput)) {
+        data = new dnnType[dim.tot()];
+        checkCuda( cudaMemcpy(data, dataInput, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost) );
+    } else {
+        data = dataInput;
+    }
+
+    int gridDim = ceil(sqrt(dim.c));
+    cv::Size gridSize(dim.w*gridDim, dim.h*gridDim);
+    cv::Mat grid = cv::Mat(gridSize, CV_8UC3, cv::Scalar(0));    
+
+    for(int i=0; i<dim.c;i++) {
+        cv::Mat raw = vizFloat2colorMap(cv::Mat(cv::Size(dim.w, dim.h),CV_32FC1, data + dim.w*dim.h*i));
+        int r = i / gridDim;
+        int c = i - r * gridDim;
+        raw.copyTo(grid.rowRange(r*dim.h, r*dim.h + dim.h).colRange(c*dim.w, c*dim.w + dim.w));
+    }
+
+    float ar = float(dim.w)/dim.h;
+    cv::Size vdim(ar*imgdim, imgdim);
+    cv::Mat viz;
+    cv::resize(grid, viz, vdim, 0, 0, 0);
+    
+    // free memory
+    if(isCudaPointer(dataInput)) {
+        delete [] data;
+    }
+    return viz;
+}
+
+cv::Mat vizLayer2Mat(tk::dnn::Network *net, int layer, int imgdim) {
+    if(layer >= net->num_layers)
+        FatalError("Could not viz layer\n");
+    return vizData2Mat(net->layers[layer]->dstData, net->layers[layer]->output_dim, imgdim);
+
+    //cv::imwrite("viz/layer" + std::to_string(layer) + ".png", viz);
+    //cv::imshow("layer", viz);
+    //cv::waitKey(0);
+}
+
+}}
@@ -5,7 +5,7 @@

 namespace tk { namespace dnn {

-Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) {
+Route::Route(Network *net, Layer **layers, int layers_n, int groups, int group_id) : Layer(net) {

    // copy input layers
    if(layers_n > MAX_LAYERS) {
@@ -15,6 +15,8 @@ Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) {
        this->layers[i] = layers[i];
    }
    this->layers_n = layers_n;
+    this->groups = groups;
+    this->group_id = group_id;
    
    //get dims
    output_dim.l = 1;
@@ -32,6 +34,7 @@ Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) {
        output_dim.c += layers[i]->output_dim.c;
    }

+    output_dim.c /= this->groups;
    input_dim = output_dim;

    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
@@ -49,8 +52,9 @@ dnnType* Route::infer(dataDim_t &dim, dnnType* srcData) {
    for(int i=0; i<layers_n; i++) {
        dnnType *input = layers[i]->dstData;
        int in_dim = layers[i]->output_dim.tot();
-        checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
-        offset += in_dim;
+        int part_in_dim = in_dim / this->groups;
+        checkCuda( cudaMemcpy(dstData + offset, input + this->group_id*part_in_dim, part_in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+        offset += part_in_dim;
    }

    //update data dimensions    
@@ -52,28 +52,80 @@ bool Yolo3Detection::init(const std::string& tensor_path, const int n_classes, c
    return true;
 } 

-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));
+cv::Mat resize_image(cv::Mat im, int w, int h)
+{
+    cv::Mat resized = cv::Mat(cv::Size(w,h), CV_32FC3, cv::Scalar(0) );
+    cv::Mat part = cv::Mat(cv::Size(w,im.rows), CV_32FC3, cv::Scalar(0) );
+    int r, c, k;
+    float w_scale = (float)(im.cols - 1) / (w - 1);
+    float h_scale = (float)(im.rows - 1) / (h - 1);
+    
+    for(k = 0; k < im.channels(); ++k){
+        for(r = 0; r < im.rows; ++r){
+            for(c = 0; c < w; ++c){
+                float val = 0;
+                if(c == w-1 || im.cols == 1){
+                    val = im.at<cv::Vec3f>(r, im.cols-1)[k];
+                } else {
+                    float sx = c*w_scale;
+                    int ix = (int) sx;
+                    float dx = sx - ix;
+                    val = (1 - dx) * im.at<cv::Vec3f>(r, ix)[k] + dx * im.at<cv::Vec3f>(r,ix+1)[k];
+                }
+                part.at<cv::Vec3f>(r,c)[k] = val;
+            }
+        }
    }
-#else
-    cv::resize(frame, frame, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
+    for(k = 0; k < im.channels(); ++k){
+        for(r = 0; r < h; ++r){
+            float sy = r*h_scale;
+            int iy = (int) sy;
+            float dy = sy - iy;
+            for(c = 0; c < w; ++c){
+                float val = (1-dy) * part.at<cv::Vec3f>(iy, c)[k];
+                resized.at<cv::Vec3f>(r, c)[k] = val;
+            }
+            if(r == h-1 || im.rows == 1) continue;
+            for(c = 0; c < w; ++c){
+                float val = dy * part.at<cv::Vec3f>(iy+1, c)[k];
+                resized.at<cv::Vec3f>(r,c)[k] += val;
+            }
+        }
+    }
+
+    return resized;
+}
+
+void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
    frame.convertTo(imagePreproc, CV_32FC3, 1/255.0); 

+    if(letterbox){
+        int im_w = frame.cols;
+        int im_h = frame.rows;
+        int net_w = netRT->input_dim.w;
+        int net_h = netRT->input_dim.h;
+        if(net_w == net_h && letterbox){
+            float ratio = ( im_w > im_h ) ? float(im_w)/float(net_w) : float(im_h)/float(net_h);
+
+            int new_h = im_h/ratio;
+            int new_w = im_w/ratio;
+
+            imagePreproc = resize_image(imagePreproc, new_w, new_h);
+
+            cv::Mat borders;
+            int top = (net_h - new_h)/2;
+            int bottom = (net_h - new_h) - top;
+            int left = (net_w - new_w)/2;
+            int right = (net_w - new_w) - left;
+
+            cv::copyMakeBorder(imagePreproc,imagePreproc, top, bottom, left, right, cv::BORDER_CONSTANT, cv::Scalar(0.5,0.5,0.5));
+        }
+        else
+            FatalError("letterbox not spported with h!=w");
+    }
+    else
+        imagePreproc = resize_image(imagePreproc, netRT->input_dim.w, netRT->input_dim.h);
+
    //split channels
    cv::split(imagePreproc,bgr);//split source

@@ -84,7 +136,6 @@ void Yolo3Detection::preprocess(cv::Mat &frame, const int bi){
        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
 }

 void Yolo3Detection::postprocess(const int bi, const bool mAP){
@@ -105,41 +156,67 @@ void Yolo3Detection::postprocess(const int bi, const bool mAP){
    }
    tk::dnn::Yolo::mergeDetections(dets, nDets, classes);

+    int im_w = originalSize[bi].width;
+    int im_h = originalSize[bi].height;
+    int net_w = netRT->input_dim.w;
+    int net_h = netRT->input_dim.h;
+    int new_h, new_w;
+
+    int top = 0, left = 0;
+
+    if(letterbox){
+        float ratio = ( im_w > im_h ) ? float(im_w)/float(net_w) : float(im_h)/float(net_h);
+        x_ratio = ratio;
+        y_ratio = ratio;
+        std::cout<<ratio<<std::endl;
+
+        int new_h = im_h/ratio;
+        int new_w = im_w/ratio;
+
+        top = (net_h - new_h)/2;
+        left = (net_w - new_w)/2;
+    }
+    else{
+        new_h = net_h;
+        new_w = net_w;
+    }
+    
+    float deltaw = net_w - new_w;
+    float deltah = net_h - new_h;
+    float ratiow = (float)new_w / net_w;
+    float ratioh = (float)new_h / net_h;
+
    // fill detected
    detected.clear();
    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 obj_class = -1;
-        float prob = 0;
+
+        float x0   = (b.x - left - b.w/2.);
+        float x1   = (b.x - left + b.w/2.);
+        float y0   = (b.y - top - b.h/2.);
+        float y1   = (b.y - top + b.h/2.);
+
+        // convert to image coords
+        x0 = x_ratio*x0;
+        x1 = x_ratio*x1;
+        y0 = y_ratio*y0;
+        y1 = y_ratio*y1;
+        
        for(int c=0; c<classes; c++) {
            if(dets[j].prob[c] >= confThreshold) {
-                obj_class = c;
-                prob = dets[j].prob[c];
-            }
-        }
+                int obj_class = c;
+                float prob = dets[j].prob[c];

-        if(obj_class >= 0) {
-            // convert to image coords
-            x0 = x_ratio*x0;
-            x1 = x_ratio*x1;
-            y0 = y_ratio*y0;
-            y1 = y_ratio*y1;
-              
-            tk::dnn::box res;
-            res.cl = obj_class;
-            res.prob = prob;
-            res.x = x0;
-            res.y = y0;
-            res.w = x1 - x0;
-            res.h = y1 - y0;
-            if(mAP)
-                for(int c=0; c<classes; c++) 
-                    res.probs.push_back(dets[j].prob[c]);
-            detected.push_back(res);
+                tk::dnn::box res;
+                res.cl = obj_class;
+                res.prob = prob;
+                res.x = x0;
+                res.y = y0;
+                res.w = x1 - x0;
+                res.h = y1 - y0;
+
+                detected.push_back(res);
+            }
        }
    }
    batchDetected.push_back(detected);
@@ -342,14 +342,16 @@ void printJsonCOCOFormat(std::ofstream *out_file, const std::string image_path,
                //min threshold confidence is set in DetectionNN.h
                if (bbox[i].probs[j] > 0) {

-                    *out_file <<    "{\"image_id\":" << image_id <<
+                    *out_file << std::fixed  << std::setprecision(6) <<
+                                    "{\"image_id\":" << image_id <<
                                    ", \"category_id\":" << coco_ids[j] << 
                                    ", \"bbox\":[" << bx << ", " << by << ", " << bw << ", " << bh << 
                                    "], \"score\":" << bbox[i].probs[j] << "},\n";
                }
            }
        else
-            *out_file <<    "{\"image_id\":" << image_id <<
+            *out_file <<     std::fixed  << std::setprecision(6) <<
+                            "{\"image_id\":" << image_id <<
                            ", \"category_id\":" << coco_ids[bbox[i].cl] << 
                            ", \"bbox\":[" << bx << ", " << by << ", " << bw << ", " << bh << 
                            "], \"score\":" << bbox[i].prob << "},\n";
@@ -0,0 +1,281 @@
+[net]
+# Testing
+#batch=1
+#subdivisions=1
+# Training
+batch=64
+subdivisions=1
+width=416
+height=416
+channels=3
+momentum=0.9
+decay=0.0005
+angle=0
+saturation = 1.5
+exposure = 1.5
+hue=.1
+
+learning_rate=0.00261
+burn_in=1000
+max_batches = 500200
+policy=steps
+steps=400000,450000
+scales=.1,.1
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers=-1
+groups=2
+group_id=1
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -1,-2
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -6,-1
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers=-1
+groups=2
+group_id=1
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -1,-2
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -6,-1
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers=-1
+groups=2
+group_id=1
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -1,-2
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[route]
+layers = -6,-1
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+##################################
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=255
+activation=linear
+
+
+
+[yolo]
+mask = 3,4,5
+anchors = 10,14,  23,27,  37,58,  81,82,  135,169,  344,319
+classes=80
+num=6
+jitter=.3
+scale_x_y = 1.05
+cls_normalizer=1.0
+iou_normalizer=0.07
+iou_loss=ciou
+ignore_thresh = .7
+truth_thresh = 1
+random=0
+resize=1.5
+nms_kind=greedynms
+beta_nms=0.6
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 23
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=255
+activation=linear
+
+[yolo]
+mask = 1,2,3
+anchors = 10,14,  23,27,  37,58,  81,82,  135,169,  344,319
+classes=80
+num=6
+jitter=.3
+scale_x_y = 1.05
+cls_normalizer=1.0
+iou_normalizer=0.07
+iou_loss=ciou
+ignore_thresh = .7
+truth_thresh = 1
+random=0
+resize=1.5
+nms_kind=greedynms
+beta_nms=0.6
@@ -15,8 +15,8 @@ int main() {
        bin_path + "/debug/layer137_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/csresnext50-panet-spp.cfg";
-    std::string name_path = "../tests/darknet/names/coco.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/csresnext50-panet-spp.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/Kcs4xBozwY4wFx8/download");

    // parse darknet network
@@ -15,10 +15,9 @@ int main() {
        bin_path + "/debug/layer137_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/csresnext50-panet-spp_berkeley.cfg";
-    std::string name_path = "../tests/darknet/names/berkeley.names";
-    // FIXME: wrong weights
-    // downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s//download");
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/csresnext50-panet-spp_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/q82qHAtqpoaFYo5/download");

    // parse darknet network
    tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
@@ -0,0 +1,70 @@
+#include<iostream>
+#include<vector>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/imgproc/imgproc.hpp>
+
+#include "tkdnn.h"
+#include "test.h"
+#include "DarknetParser.h"
+#include "NetworkViz.h"
+
+int main(int argc, char *argv[]) {
+    if(argc <2)
+        FatalError("you must provide an input image");
+    std::string input_image = argv[1];
+    std::string bin_path  = "yolo3";
+    std::string wgs_path  = bin_path + "/layers";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
+    downloadWeightsifDoNotExist(wgs_path, bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download");
+
+    // parse darknet network
+    tk::dnn::Network *net = tk::dnn::darknetParser(cfg_path, wgs_path, name_path);
+    net->print();
+
+    // input data
+    dnnType *input_d;
+    checkCuda( cudaMalloc(&input_d, sizeof(dnnType)*net->input_dim.tot()));
+
+    // load image
+    cv::Mat frame, frameFloat;
+    frame = cv::imread(input_image);
+    cv::resize(frame, frame, cv::Size(net->input_dim.w, net->input_dim.h));
+    frame.convertTo(frameFloat, CV_32FC3, 1/255.0); 
+
+    //split channels
+    cv::Mat bgr[3];
+    cv::split(frameFloat,bgr);//split source
+
+    //write channels
+    for(int i=0; i<net->input_dim.c; i++) {
+        int idx = i*frameFloat.rows*frameFloat.cols;
+        int ch = net->input_dim.c-1 -i;
+        checkCuda( cudaMemcpy(input_d + idx, (void*)bgr[ch].data, frameFloat.rows*frameFloat.cols*sizeof(dnnType), cudaMemcpyHostToDevice));     
+    }
+    
+    tk::dnn::dataDim_t dim =  net->input_dim;     
+    dim.print();
+    std::cout<<"infer\n";
+    net->infer(dim, input_d);    
+
+    // output directory
+    std::string output_viz = "viz/";
+    system( (std::string("mkdir -p ") + output_viz).c_str() );
+
+    for(int i=0; i<net->num_layers; i++) {
+        std::string output_png = output_viz + "/layer" + std::to_string(i) + ".png";
+        std::cout<<"saving "<<output_png<<"\n";
+        cv::Mat viz = vizLayer2Mat(net, i);
+        cv::imwrite(output_png, viz);
+        //cv::imshow("layer", viz);
+        //cv::waitKey(0);
+    }
+
+    checkCuda(cudaFree(input_d));
+    net->releaseLayers();
+    delete net;
+    return 0;
+}
+
+ 
@@ -13,8 +13,8 @@ int main() {
        bin_path + "/layers/output.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo2.cfg";
-    std::string name_path = "../tests/darknet/names/coco.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo2.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/nf4PJ3k8bxBETwL/download");

    // parse darknet network
@@ -13,8 +13,8 @@ int main() {
        bin_path + "/layers/output.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo2_voc.cfg";
-    std::string name_path = "../tests/darknet/names/voc.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo2_voc.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/voc.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/DJC5Fi2pEjfNDP9/download");

    // parse darknet network
@@ -13,8 +13,8 @@ int main() {
        bin_path + "/layers/output.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo2tiny.cfg";
-    std::string name_path = "../tests/darknet/names/coco.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo2tiny.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
    // FIXME: wrong weights
    //downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s//download");

@@ -15,8 +15,8 @@ int main() {
        bin_path + "/debug/layer106_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo3.cfg";
-    std::string name_path = "../tests/darknet/names/coco.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/jPXmHyptpLoNdNR/download");

    // parse darknet network
@@ -31,4 +31,4 @@ int main() {
    delete net;
    delete netRT;
    return ret;
-}
+}
@@ -15,8 +15,8 @@ int main() {
        bin_path + "/debug/layer106_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo3_512.cfg";
-    std::string name_path = "../tests/darknet/names/coco.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3_512.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/RGecMeGLD4cXEWL/download");

    // parse darknet network
@@ -15,8 +15,8 @@ int main() {
        bin_path + "/debug/layer106_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo3_berkeley.cfg";
-    std::string name_path = "../tests/darknet/names/berkeley.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3_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/o5cHa4AjTKS64oD/download");

    // parse darknet network
@@ -15,8 +15,8 @@ int main() {
        bin_path + "/debug/layer106_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo3_coco4.cfg";
-    std::string name_path = "../tests/darknet/names/coco4.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3_coco4.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco4.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/o27NDzSAartbyc4/download");

    // parse darknet network
@@ -15,8 +15,8 @@ int main() {
        bin_path + "/debug/layer106_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo3_flir.cfg";
-    std::string name_path = "../tests/darknet/names/flir.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3_flir.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/flir.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/62DECncmF6bMMiH/download");

    // parse darknet network
@@ -14,8 +14,8 @@ int main() {
        bin_path + "/debug/layer23_out.bin",
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo3tiny.cfg";
-    std::string name_path = "../tests/darknet/names/coco.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3tiny.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/LMcSHtWaLeps8yN/download");

    // parse darknet network
@@ -14,8 +14,8 @@ int main() {
        bin_path + "/debug/layer23_out.bin",
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo3tiny_512.cfg";
-    std::string name_path = "../tests/darknet/names/coco.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo3tiny_512.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/8Zt6bHwHADqP4JC/download");

    // parse darknet network
@@ -15,8 +15,8 @@ int main() {
        bin_path + "/debug/layer161_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo4.cfg";
-    std::string name_path = "../tests/darknet/names/coco.names";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/d97CFzYqCPCp5Hg/download");

    // parse darknet network
@@ -15,8 +15,8 @@ int main() {
        bin_path + "/debug/layer161_out.bin"
    };
    std::string wgs_path  = bin_path + "/layers";
-    std::string cfg_path  = "../tests/darknet/cfg/yolo4_berkeley.cfg";
-    std::string name_path = "../tests/darknet/names/berkeley.names";
+    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/nkWFa5fgb4NTdnB/download");

    // parse darknet network
@@ -0,0 +1,33 @@
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+#include "test.h"
+#include "DarknetParser.h"
+
+int main() {
+    std::string bin_path  = "yolo4tiny";
+    std::vector<std::string> input_bins = { 
+        bin_path + "/layers/input.bin"
+    };
+    std::vector<std::string> output_bins = {
+        bin_path + "/debug/layer30_out.bin",
+        bin_path + "/debug/layer37_out.bin"
+    };
+    std::string wgs_path  = bin_path + "/layers";
+    std::string cfg_path  = std::string(TKDNN_PATH) + "/tests/darknet/cfg/yolo4tiny.cfg";
+    std::string name_path = std::string(TKDNN_PATH) + "/tests/darknet/names/coco.names";
+    downloadWeightsifDoNotExist(input_bins[0], bin_path, "https://cloud.hipert.unimore.it/s/iRnc4pSqmx78gJs/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;
+}
Author	SHA1	Message	Date
Micaela Verucchi	8ff7cad2e1	Modified resize and boxes coordinates to float, to achieve same darknet accuracy Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>	2020-08-05 17:25:33 +02:00
Francesco Gatti	a5d2d4792a	fix coords convert	2020-07-27 13:52:39 +02:00
Francesco Gatti	3a0802d70c	Resolve detection objects pick by prob threshold. Before this it will only pick the last object with prob > thresh wich is absolutely wrong Now it picks all the objects with prob > thesh. fixes #94	2020-07-27 13:45:42 +02:00
Micaela Verucchi	f4970d1e6f	Update README.md	2020-07-17 14:37:10 +02:00
Micaela Verucchi	6a68f19b2c	Fix patch from @ahmedius2 , tkDNN now supports CUDNN 8.0.1 (Fix #74 ) Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com> Francesco Gatti <gattifrancesco@hotmail.it>	2020-07-16 19:06:54 +02:00
tk	c4aad7fe95	Patch for CUDNN 8.0.1 Signed-off-by: tk <micaelaverucchi@gmail.com>	2020-07-16 18:16:09 +02:00
Francesco Gatti	b2df9fc110	Update README.md	2020-07-13 19:51:09 +02:00
Francesco Gatti	b12cf0d7c2	docker	2020-07-13 19:50:00 +02:00
Micaela Verucchi	7c2155decf	Add weights download link for csresnext50-panet-spp_berkeley ( fix #63 ) Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>	2020-07-01 11:56:46 +02:00
Francesco Gatti	04602f3952	Yolo4-tiny batched fix #59	2020-06-30 19:37:16 +02:00
Micaela Verucchi	fe2e4eae92	yolo4tiny works on tensorRT 🐬 🐬 🐬 🐬 Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>	2020-06-30 15:52:58 +02:00
Micaela Verucchi	61aa24c6b7	yolov4tiny works on CUDNN Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>	2020-06-30 15:19:03 +02:00
Francesco Gatti	6d9beb1ec5	Merge pull request #53 from omaralvarez/master Fix error when parsing label files with unexpected chars	2020-06-25 10:07:41 +02:00
Omar Alvarez	2817ade782	Fix parsing label files with unexpected chars	2020-06-22 13:40:15 +02:00
Francesco Gatti	1dfc69ba89	viz yolo3 preprocess	2020-06-16 13:19:30 +02:00
Francesco Gatti	07193fc343	dealloc in test.h fix #36	2020-06-16 12:44:12 +02:00
Francesco Gatti	1b8f45703f	darknet parser cpp	2020-06-16 12:41:48 +02:00
Francesco Gatti	3d3a2427c9	viz yolo3	2020-06-16 12:36:45 +02:00
Micaela Verucchi	285c77d6dd	Imptove relative paths Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com> Francesco Gatti <gattifrancesco@hotmail.it>	2020-06-15 17:15:37 +02:00
Francesco Gatti	6dff675db7	fix wrong commit	2020-06-14 13:03:48 +02:00
Francesco Gatti	cbfc8ea4f2	serialize fix	2020-06-14 13:01:45 +02:00
Francesco Gatti	567dc0f75d	cmake cudnn fix	2020-06-14 12:57:29 +02:00
Micaela Verucchi	ab6d2d1766	Update README Signed-off-by: Micaela Verucchi <micaela.verucchi@unimore.it>	2020-06-12 11:39:46 +02:00
Micaela Verucchi	be9e327aef	Fix minor, update readme Signed-off-by: Micaela Verucchi <micaelaverucchi@gmail.com>	2020-06-09 21:01:19 +02:00
Francesco Gatti	211eff8ad0	Merge branch 'master' of https://github.com/ceccocats/tkDNN	2020-06-09 20:53:17 +02:00
Francesco Gatti	20303ac32e	cudnn8 compile	2020-06-09 20:53:05 +02:00
Micaela Verucchi	3ea23815a4	Update README.md	2020-06-03 11:01:17 +02:00