add DLA, plugin for shortcut and leaky. new verison 0.4

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

.gitignore

@@ -8,3 +8,4 @@ build/
 *.h5
 *.tar.gz
 *.weights
+.idea/

CMakeLists.txt

@@ -1,47 +1,46 @@
-cmake_minimum_required(VERSION 2.8)
+cmake_minimum_required(VERSION 3.5)
+
 project (tkDNN)
+set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} ${CMAKE_CURRENT_SOURCE_DIR}/cmake)
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -fPIC")
+include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include/tkDNN)
 
-set(BUILD_DEPS true CACHE BOOL "If true download deps")
-
-if( ${BUILD_DEPS} )
-	message("Launching pre-build dependency installer script...")
-
-    execute_process (COMMAND bash -c "bash build_models.sh download"
-        WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tests)
-
-	set(BUILD_DEPS false CACHE BOOL "If true download deps" FORCE)
-	message("Finished dowloading test weights")
-endif()
-
+# project specific flags
 if(DEBUG)
     add_definitions(-DDEBUG)
 endif()
 
-find_package(CUDA QUIET REQUIRED)
-find_package(OpenCV QUIET)
-if(OPENCV)
-    message("Compiling with openCV support")
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV")
-else()
-    message(WARNING "compiling without OpenCV")
-endif()
 
-cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS})
-cuda_add_library(kernels SHARED src/kernels/activation_elu.cu 
-                                src/kernels/activation_leaky.cu
-                                src/kernels/activation_logistic.cu
-                                src/kernels/reorg.cu
-                                src/kernels/softmax.cu
-                                src/kernels/convert.cu)
+#-------------------------------------------------------------------------------
+# CUDA
+#-------------------------------------------------------------------------------
+find_package(CUDA 9.0 REQUIRED)
+SET(CUDA_SEPARABLE_COMPILATION ON)
+#set(CUDA_NVCC_FLAGS "${CUDA_NVCC_FLAGS} -arch=sm_30 --compiler-options '-fPIC'")
 
+find_package(CUDNN REQUIRED)
+
+# compile
+file(GLOB tkdnn_CUSRC "src/kernels/*.cu")
+cuda_include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${CUDNN_INCLUDE_DIRS})
+cuda_add_library(kernels SHARED ${tkdnn_CUSRC})
+
+
+#-------------------------------------------------------------------------------
+# External Libraries
+#-------------------------------------------------------------------------------
+find_package(OpenCV REQUIRED)
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DOPENCV")
+
+
+#-------------------------------------------------------------------------------
+# Build Libraries
+#-------------------------------------------------------------------------------
 file(GLOB tkdnn_SRC "src/*.cpp")
-set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer ${OpenCV_LIBS})
+set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} ${CUDNN_LIBRARIES} ${OpenCV_LIBS})
 
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11")
-if(NOT OPENCV) 
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -D_GLIBCXX_USE_CXX11_ABI=0") 
-endif()
-include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS})
+include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS} ${OPENCV_INCLUDE_DIRS} ${NVINFER_INCLUDES})
 add_library(tkDNN SHARED ${tkdnn_SRC})
 target_link_libraries(tkDNN ${tkdnn_LIBS})
 
@@ -62,6 +61,9 @@ target_link_libraries(test_mnistRT tkDNN)
 add_executable(test_yolo tests/yolo/yolo.cpp)
 target_link_libraries(test_yolo tkDNN)
 
+add_executable(test_yolo_voc tests/yolo_voc/yolo_voc.cpp)
+target_link_libraries(test_yolo_voc tkDNN)
+
 add_executable(test_yolo_tiny tests/yolo_tiny/yolo_tiny.cpp)
 target_link_libraries(test_yolo_tiny tkDNN)
 
@@ -71,21 +73,54 @@ target_link_libraries(test_yolo_relu tkDNN)
 
 add_executable(test_yolo_224 tests/yolo_224/yolo_224.cpp)
 target_link_libraries(test_yolo_224 tkDNN)
+
+add_executable(test_yolo_berkeley tests/yolo_berkeley/yolo_berkeley.cpp)
+target_link_libraries(test_yolo_berkeley tkDNN)
+
+add_executable(test_yolo3_coco4 tests/yolo3_coco4/yolo3_coco4.cpp)
+target_link_libraries(test_yolo3_coco4 tkDNN)
+
+add_executable(test_yolo3_berkeley tests/yolo3_berkeley/yolo3_berkeley.cpp)
+target_link_libraries(test_yolo3_berkeley tkDNN)
+
+add_executable(test_yolo3_flir tests/yolo3_flir/yolo3_flir.cpp)
+target_link_libraries(test_yolo3_flir tkDNN)
 ################################################################################
 
 
 add_executable(test_rtinference tests/test_rtinference/rtinference.cpp)
 target_link_libraries(test_rtinference tkDNN)
 
-add_executable(detection demo/detection/detection.cpp)
-target_link_libraries(detection tkDNN)
+add_executable(yolo3_demo demo/demo/demo.cpp)
+target_link_libraries(yolo3_demo tkDNN)
 
-#install
-if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
-    set (CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}/install" 
-         CACHE PATH "default install path" FORCE)
-endif()
+
+#-------------------------------------------------------------------------------
+# Install
+#-------------------------------------------------------------------------------
+#if (CMAKE_INSTALL_PREFIX_INITIALIZED_TO_DEFAULT)
+#    set (CMAKE_INSTALL_PREFIX "${CMAKE_BINARY_DIR}/install" 
+#         CACHE PATH "default install path" FORCE)
+#endif()
 message("install dir:" ${CMAKE_INSTALL_PREFIX})
-install(DIRECTORY include/ DESTINATION include/${CMAKE_PROJECT_NAME}
-FILES_MATCHING PATTERN "*.h")
+install(DIRECTORY include/ DESTINATION include/)
 install(TARGETS tkDNN kernels DESTINATION lib)
+install(DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/cmake/" # source directory
+        DESTINATION "share/tkDNN/cmake/" # target directory
+)
+
+
+#-------------------------------------------------------------------------------
+# Prepare for test
+#-------------------------------------------------------------------------------
+set(TEST_DATA true CACHE BOOL "If true download deps")
+if( ${TEST_DATA} )
+	message("Launching pre-build dependency installer script...")
+
+    execute_process (COMMAND bash -c "bash build_models.sh download"
+        WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tests)
+
+	set(TEST_DATA false CACHE BOOL "If true download deps" FORCE)
+	message("Finished dowloading test weights")
+endif()
+

README.md

@@ -1,11 +1,12 @@
 # tkDNN
-tkDNN is a Deep Neural Network library built with cuDNN primitives specifically thought to work on NVIDIA TK1 board.<br>
+tkDNN is a Deep Neural Network library built with cuDNN primitives specifically thought to work on NVIDIA TK1(and all successive) board.<br>
 The main scope is to do high performance inference on already trained models.
 
-this branch is actually work on every NVIDIA GPU that support the dependencies:
-* CUDA 8
-* CUDNN 6
-* TENSORRT 2
+this branch actually work on every NVIDIA GPU that support the dependencies:
+* CUDA 10.0
+* CUDNN 7.603
+* TENSORRT 6.01
+* OPENCV 4.1
 
 ## Workflow
 The recommended workflow follow these step:
@@ -20,6 +21,7 @@ Build with cmake
 mkdir build
 cd build
 cmake ..
+# use -DTEST_DATA=False to skip dataset download
 make
 ```
 during the cmake configuration it will be dowloaded the weights needed for running
@@ -32,4 +34,18 @@ Assumiung you have correctly builded the library these are the test ready to exe
 * test_mnistRT: the mnist network hardcoded in using tensorRT apis (TENSORRT only)
 * test_yolo: YOLO detection network (CUDNN and TENSORRT)
 * test_yolo_tiny: smaller version of YOLO (CUDNN and TENSRRT)
+* test_yolo3_berkeley: our yolo3 version trained with BDD100K dateset 
 
+## yolo3 berkeley demo detection
+For the live detection you need to precompile the tensorRT file by luncing the desidered network test, this is the recommended process:
+```
+export TKDNN_MODE=FP16   # set the half floating point optimization
+rm yolo3_berkeley.rt		 # be sure to delete(or move) old tensorRT files
+./test_yolo3_berkeley              # run the yolo test (is slow)
+# with f16 inference the result will be a bit incorrect
+```
+this will genereate a yolo3_berkeley.rt file that can be used for live detection:
+```
+./yolo3_demo                               # launch detection on a demo video
+./yolo3_demo yolo3_berkeley.rt /dev/video0 # launch detection on device 0
+```

cmake/FindCUDNN.cmake

@@ -0,0 +1,33 @@
+# Find the header files
+
+find_path(CUDNN_INCLUDE_DIR
+        ${CMAKE_SYSROOT}/usr/local/include
+        ${CMAKE_SYSROOT}/usr/include
+        /usr/local/nvidia/tensorrt/include/
+        NO_DEFAULT_PATH
+)
+
+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/
+    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)

cmake/tkDNNConfig.cmake

@@ -0,0 +1,24 @@
+message("-- Found tkDNN")
+set(CMAKE_MODULE_PATH ${CMAKE_CURRENT_LIST_DIR})
+set(CMAKE_CXX_FLAGS  "${CMAKE_CXX_FLAGS} --std=c++11 -fPIC")
+
+find_package(CUDA REQUIRED)
+find_package(OpenCV REQUIRED)
+find_package(CUDNN REQUIRED)
+
+set(tkDNN_INCLUDE_DIRS 
+	${CUDA_INCLUDE_DIRS} 
+	${OPENCV_INCLUDE_DIRS} 
+    ${CUDNN_INCLUDE_DIRS}
+)
+
+set(tkDNN_LIBRARIES 
+    tkDNN 
+    kernels 
+    ${CUDA_LIBRARIES} 
+    ${CUDA_CUBLAS_LIBRARIES}
+	${CUDNN_LIBRARIES}
+	${OpenCV_LIBS}
+)
+
+set(tkDNN_FOUND true)

demo/demo/demo.cpp

@@ -0,0 +1,109 @@
+#include <iostream>
+#include <signal.h>
+#include <stdlib.h>     /* srand, rand */
+#include <unistd.h>
+#include <mutex>
+#include "utils.h"
+
+#include <opencv2/core/core.hpp>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/videoio.hpp>
+#include <opencv2/imgproc/imgproc.hpp>
+
+#include "Yolo3Detection.h"
+
+bool 			gRun;
+bool SAVE_RESULT = false;
+
+void sig_handler(int signo) {
+    std::cout<<"request gateway stop\n";
+    gRun = false;
+}
+
+int main(int argc, char *argv[]) {
+
+    std::cout<<"detection\n";
+    signal(SIGINT, sig_handler);
+
+
+    char *net = "yolo3_berkeley.rt";
+    if(argc > 1)
+        net = argv[1]; 
+    char *input = "../demo/yolo_test.mp4";
+    if(argc > 2)
+        input = argv[2]; 
+
+    tk::dnn::Yolo3Detection yolo;
+    yolo.init(net);
+
+    gRun = true;
+
+    cv::VideoCapture cap(input);
+    if(!cap.isOpened())
+        gRun = false; 
+    else
+        std::cout<<"camera started\n";
+
+
+    cv::VideoWriter resultVideo;
+    if(SAVE_RESULT) {
+        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;
+    cv::Mat dnn_input;
+    cv::namedWindow("detection", cv::WINDOW_NORMAL);
+    
+    while(gRun) {
+        cap >> frame; 
+        if(!frame.data) {
+            break;
+        }  
+ 
+        // this will be resized to the net format
+        dnn_input = frame.clone();
+        // TODO: async infer
+        yolo.update(dnn_input);
+
+        // draw dets
+        for(int i=0; i<yolo.detected.size(); i++) {
+            tk::dnn::box b = yolo.detected[i];
+            int x0   				= b.x;
+            int x1   				= b.x + b.w;
+            int y0   				= b.y;
+            int y1   				= b.y + b.h;
+            std::string det_class 	= yolo.getYoloLayer()->classesNames[b.cl];
+            float prob 				= b.prob;
+
+            std::cout<<det_class<<" ("<<prob<<"): "<<x0<<" "<<y0<<" "<<x1<<" "<<y1<<"\n";
+			// draw rectangle
+            cv::rectangle(frame, cv::Point(x0, y0), cv::Point(x1, y1), yolo.colors[b.cl], 2); 
+
+	        // draw label
+            int baseline = 0;
+            float fontScale = 0.5;
+            int thickness = 2;
+            cv::Size textSize = getTextSize(det_class, cv::FONT_HERSHEY_SIMPLEX, fontScale, thickness, &baseline);
+            cv::rectangle(frame, cv::Point(x0, y0), cv::Point((x0 + textSize.width - 2), (y0 - textSize.height - 2)), yolo.colors[b.cl], -1);                      
+            cv::putText(frame, det_class, cv::Point(x0, (y0 - (baseline / 2))), cv::FONT_HERSHEY_SIMPLEX, fontScale, cv::Scalar(255, 255, 255), thickness);
+        }
+    
+        cv::imshow("detection", frame);
+        cv::waitKey(1);
+        if(SAVE_RESULT)
+            resultVideo << frame;
+    }
+
+    std::cout<<"detection end\n";   
+
+
+    std::cout<<COL_GREENB<<"\n\nTime stats:\n";
+    std::cout<<"Min: "<<*std::min_element(yolo.stats.begin(), yolo.stats.end())<<" ms\n";    
+    std::cout<<"Max: "<<*std::max_element(yolo.stats.begin(), yolo.stats.end())<<" ms\n";    
+    double mean = 0; for(int i=0; i<yolo.stats.size(); i++) mean += yolo.stats[i]; mean /= yolo.stats.size();
+    std::cout<<"Avg: "<<mean<<" ms\n"<<COL_END;    
+    return 0;
+}
+

demo/detection/detection.cpp

@@ -1,249 +0,0 @@
-#include<iostream>
-#include "tkdnn.h"
-#include <stdlib.h>     /* srand, rand */
-#include <unistd.h>
-
-#include <opencv2/core/core.hpp>
-#include <opencv2/highgui/highgui.hpp>
-#include <opencv2/imgproc/imgproc.hpp>
-
-const char *reg_bias = "../tests/yolo/layers/g31.bin";
-
-int prob_sort(const void *pa, const void *pb) {
-    tkDNN::box a = *(tkDNN::box *)pa;
-    tkDNN::box b = *(tkDNN::box *)pb;
-    float diff = a.prob - b.prob;
-    if(diff < 0) return 1;
-    else if(diff > 0) return -1;
-    return 0;
-}
-
-cv::Mat GetSquareImage(const cv::Mat& img, int target_width) {
-    int width = img.cols, height = img.rows;
-
-    cv::Mat square = cv::Mat::zeros( target_width, target_width, img.type() );
-
-    int max_dim = ( width >= height ) ? width : height;
-    float scale = ( ( float ) target_width ) / max_dim;
-    cv::Rect roi;
-    if ( width >= height )
-    {
-        roi.width = target_width;
-        roi.x = 0;
-        roi.height = height * scale;
-        roi.y = ( target_width - roi.height ) / 2;
-    }
-    else
-    {
-        roi.y = 0;
-        roi.height = target_width;
-        roi.width = width * scale;
-        roi.x = ( target_width - roi.width ) / 2;
-    }
-
-    cv::resize( img, square( roi ), roi.size() );
-
-    return square;
-}
-
-//return inference time
-double compute_image( cv::Mat imageORIG, 
-                    tkDNN::NetworkRT *netRT, tkDNN::RegionInterpret *rI,
-                    dnnType *input, dnnType *output) {
-
-    //Resize with padding and convert to float
-    cv::Mat image = GetSquareImage(imageORIG, netRT->input_dim.w);
-    cv::Mat imageF;
-    image.convertTo(imageF, CV_32FC3, 1/255.0); 
-
-    //split channels
-    cv::Mat bgr[3];   //destination array
-    cv::split(imageF,bgr);//split source
-    
-    //write channels
-    int idx = 0;
-    memcpy((void*)&input[idx], (void*)bgr[2].data, imageF.rows*imageF.cols*sizeof(dnnType));
-    idx = imageF.rows*imageF.cols;
-    memcpy((void*)&input[idx], (void*)bgr[1].data, imageF.rows*imageF.cols*sizeof(dnnType));
-    idx *= 2;    
-    memcpy((void*)&input[idx], (void*)bgr[0].data, imageF.rows*imageF.cols*sizeof(dnnType));
-
-    //DO INFERENCE
-    printCenteredTitle(" TENSORRT inference ", '=', 30); 
-    TIMER_START
-    checkCuda(  cudaMemcpyAsync(netRT->buffersRT[netRT->buf_input_idx], input, 
-                                netRT->input_dim.tot()*sizeof(float), 
-                                cudaMemcpyHostToDevice, netRT->stream));
-    netRT->enqueue();
-    checkCuda(  cudaMemcpyAsync(output, netRT->buffersRT[netRT->buf_output_idx], 
-                                netRT->output_dim.tot()*sizeof(float), 
-                                cudaMemcpyDeviceToHost, netRT->stream));
-    cudaStreamSynchronize(netRT->stream);
-    TIMER_STOP
-    
-
-    rI->interpretData(output, imageORIG.cols, imageORIG.rows);
-
-    return t_ns;
-}
-
-int print_usage() {
-    std::cout<<"usage: ./detection net.rt validation_list.txt"
-             <<" [-t <thresh>] [-s] [-i <iterations>]\n"
-             <<" -t: set thresh value\n -s: show images as compute\n"
-             <<" -i: images to compute\n\n" 
-             <<"> validation_list.txt format: \n"
-             <<"    path/to/image.jpg path/to/label.txt\n"
-             <<"> label.txt format: \n"  
-             <<"    <object-class> <x> <y> <width> <height>\n"
-             <<"    x and y are the box center, "   
-             <<"all values are relative to the image size\n\n";
-    return 1;
-}
-
-
-int main(int argc, char *argv[]) {
-
-    //params
-    char *tensor_path = NULL;
-    char *imageset_path = NULL; 
-    float thresh = 0.3f;
-    bool  show = false;
-    int iterations = INT_MAX;
-
-    //parse params
-    int c;
-    while ((c = getopt (argc, argv, "t:si:")) != -1) {
-        switch(c) {
-        case 't':   thresh = atof(optarg);       break;
-        case 's':   show = true;                 break;
-        case 'i':   iterations = atoi(optarg);   break;
-        case '?':   
-            return print_usage();
-        default:    return print_usage();
-        }
-    }
-
-    if(argc - optind == 2) {
-        tensor_path   = argv[optind];
-        imageset_path = argv[optind+1];
-    } else {
-        std::cout<<"not enough arguments.\n";
-        return print_usage();
-    }
-    //end parsing
-
-    if(!fileExist(tensor_path))
-        FatalError("unable to read serialRT file");
-    //convert network to tensorRT
-    tkDNN::NetworkRT netRT(NULL, tensor_path);
-    tkDNN::RegionInterpret rI(netRT.input_dim, netRT.output_dim, 80, 4, 5, thresh, reg_bias);
-
-    dnnType *input = new float[netRT.input_dim.tot()];
-    dnnType *output = new float[netRT.output_dim.tot()];
-
-    std::string line;
-    std::ifstream imageset(imageset_path);
-    if(!imageset.is_open())
-        FatalError("could not read imageset");
-    
-    double mTime = 0;
-    float mAP = 0;
-    int processed_images;
-
-    for(processed_images=1; 
-        processed_images-1 < iterations && getline(imageset, line); 
-        processed_images++) {
-        
-        std::string image_path = line.substr(0, line.find(" "));
-        std::string label_path = line.substr(line.find(" ")+1, line.size());
-        std::cout<<image_path<<"\n"<<label_path<<"\n";
-
-        //LOAD IMAGE
-        cv::Mat img = cv::imread(image_path.c_str(), CV_LOAD_IMAGE_COLOR);
-        if(!img.data)                              
-            FatalError("Could not open image");
-        std::cout<<"Image size: ("<<img.cols<<"x"<<img.rows<<")\n";
-
-        mTime += compute_image(img, &netRT, &rI, input, output); 
-
-        std::ifstream labels(label_path.c_str());
-        if(!labels.is_open())
-            FatalError("could not read labels");
-
-
-        qsort(rI.res_boxes, rI.res_boxes_n, sizeof(tkDNN::box), prob_sort);
-        for(int i=0; i<rI.res_boxes_n; i++) {
-            tkDNN::box bx = rI.res_boxes[i];
-            std::cout<<" ("<<int(bx.prob*100)<<"%) "<<bx.cl
-                     <<": "<<bx.x<<" "<<bx.y<<" "<<bx.w<<" "<<bx.h<<"\n";
-
-            cv::rectangle(img, cv::Point(bx.x - bx.w/2, bx.y - bx.h/2), 
-                               cv::Point(bx.x + bx.w/2, bx.y + bx.h/2),
-            cv::Scalar( 0, 0, 255), 2);
-        }
-
-        std::cout<<"GROUND TRUTH\n";
-        tkDNN::box gt[256];
-        int gt_n = 0; 
-        int cl;
-        float x, y, w, h;
-        while(labels>>cl) {
-            labels>>x>>y>>w>>h;
-            w *= img.cols; x *= img.cols; 
-            h *= img.rows; y *= img.rows; 
-            std::cout<<cl<<": "<<x<<" "<<y<<" "<<w<<" "<<h<<"\n";
-            gt[gt_n].x = x;
-            gt[gt_n].y = y;
-            gt[gt_n].w = w;
-            gt[gt_n].h = h;
-            gt[gt_n].cl = cl;
-            gt_n++;            
-
-            cv::rectangle(img, cv::Point(x -w/2, y -h/2), 
-                               cv::Point(x +w/2, y +h/2),
-            cv::Scalar( 255, 0, 0), 2);
-        }
-
-        //AP calculation
-        float AP = 0;
-        for(int i=rI.res_boxes_n; i>=1; i--) {  //for each detected evaluate sub group
-            
-            int prec = 0;
-            for(int j=0; j<i; j++) {            //for each detected in sub group
-                for(int z=0; z<gt_n; z++) {     //control each ground truth
-                    float iou = tkDNN::RegionInterpret::box_iou(rI.res_boxes[j], gt[z]);
-                    if(iou > 0.6f && rI.res_boxes[j].cl == gt[z].cl) {
-                        prec++;
-                        break;
-                    }
-                }
-            }
-
-            AP += float(prec)/i;
-        }
-        AP = AP/gt_n;
-        std::cout<<"AP: "<<AP<<"\n";        
-        
-        mAP += AP;
-        std::cout<<"#### processed: "<<processed_images
-                 <<", mAP: "<<mAP/processed_images<<"\n";        
-
-        //show results
-        if(show) {
-            cv::namedWindow("result");
-            cv::imshow("result", img);
-            cv::waitKey(10);
-        }        
-    }
-
-    //print results to file
-    processed_images -= 1;
-    std::ofstream res("results.txt", std::ios::app);
-    res<<"#### "<<tensor_path<<"\n";
-    res<<"processed images: "<<processed_images<<"\n";
-    res<<"mean inference time: "<<mTime/processed_images<<"\n";
-    res<<"mean AP: "<<mAP/processed_images<<"\n";
-    res<<"thesh used: "<<thresh<<"\n\n";
-    return 0;
-}

include/tkDNN/Layer.h

@@ -2,10 +2,11 @@
 #define LAYER_H
 
 #include<iostream>
+#include<vector>
 #include "utils.h"
 #include "Network.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 enum layerType_t {
     LAYER_DENSE,
@@ -17,9 +18,14 @@ enum layerType_t {
     LAYER_SOFTMAX,
     LAYER_ROUTE,
     LAYER_REORG,
-    LAYER_REGION
+    LAYER_SHORTCUT,
+    LAYER_UPSAMPLE,
+    LAYER_REGION,
+    LAYER_YOLO
 };
 
+#define TKDNN_BN_MIN_EPSILON 1e-5
+
 /** 
     Simple layer Father class
 */
@@ -50,7 +56,10 @@ public:
             case LAYER_SOFTMAX:     return "Softmax";
             case LAYER_ROUTE:       return "Route";            
             case LAYER_REORG:       return "Reorg";
+            case LAYER_SHORTCUT:    return "Shortcut";
+            case LAYER_UPSAMPLE:    return "Upsample";
             case LAYER_REGION:      return "Region";
+            case LAYER_YOLO:        return "Yolo";
             default:                return "unknown";
         }
     }
@@ -69,7 +78,7 @@ class LayerWgs : public Layer {
 
 public:
     LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kt, 
-             const char* fname_weights, bool batchnorm = false); 
+             std::string fname_weights, bool batchnorm = false); 
     virtual ~LayerWgs();
 
     int inputs, outputs;
@@ -102,7 +111,7 @@ public:
 class Dense : public LayerWgs {
 
 public:
-    Dense(Network *net, int out_ch, const char* fname_weights); 
+    Dense(Network *net, int out_ch, std::string fname_weights); 
     virtual ~Dense();
     virtual layerType_t getLayerType() { return LAYER_DENSE; };
 
@@ -145,7 +154,7 @@ class Conv2d : public LayerWgs {
 public:
     Conv2d( Network *net, int out_ch, int kernelH, int kernelW, 
                 int strideH, int strideW, int paddingH, int paddingW,
-                const char* fname_weights, bool batchnorm = false); 
+                std::string fname_weights, bool batchnorm = false); 
     virtual ~Conv2d();
     virtual layerType_t getLayerType() { return LAYER_CONV2D; };
 
@@ -282,6 +291,39 @@ public:
     int stride;
 };
 
+/**
+    Shortcut layer
+    sum with stride another layer
+*/
+class Shortcut : public Layer {
+
+public:
+    Shortcut(Network *net, Layer *backLayer); 
+    virtual ~Shortcut();
+    virtual layerType_t getLayerType() { return LAYER_SHORTCUT; };
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+
+public:
+    Layer *backLayer;
+};
+
+/**
+    Upsample layer
+    Mantain same dimension but change C*H*W distribution
+*/
+class Upsample : public Layer {
+
+public:
+    Upsample(Network *net, int stride);
+    virtual ~Upsample();
+    virtual layerType_t getLayerType() { return LAYER_UPSAMPLE; };
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+
+    int stride;
+    bool reverse;
+};
 
 struct box {
     int cl;
@@ -294,9 +336,46 @@ struct sortable_bbox {
     float **probs;
 };
 
+/**
+    Yolo3 layer
+*/
+class Yolo : public Layer {
+
+public:
+    struct box {
+        float x, y, w, h;
+    };
+
+    struct detection{
+        Yolo::box bbox;
+        int classes;
+        float *prob;
+        float *mask;
+        float objectness;
+        int sort_class;
+    };
+
+    Yolo(Network *net, int classes, int num, std::string fname_weights);
+    virtual ~Yolo();
+    virtual layerType_t getLayerType() { return LAYER_YOLO; };
+
+    int classes, num;
+    dnnType *mask_h, *mask_d; //anchors
+    dnnType *bias_h, *bias_d; //anchors
+    std::vector<std::string> classesNames;
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+    int computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh);
+
+    dnnType *predictions;
+
+    static const int MAX_DETECTIONS = 256;
+    static Yolo::detection *allocateDetections(int nboxes, int classes);
+    static void             mergeDetections(Yolo::detection *dets, int ndets, int classes);
+};
+
 /**
     Region layer
-    Mantain same dimension but change C*H*W distribution
 */
 class Region : public Layer {
 
@@ -314,7 +393,7 @@ class RegionInterpret {
 
 public:
     RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, 
-                    int classes, int coords, int num, float thresh, const char* fname_weights);
+                    int classes, int coords, int num, float thresh, std::string fname_weights);
     ~RegionInterpret();
 
     dataDim_t input_dim, output_dim;
@@ -340,5 +419,5 @@ public:
     static float box_iou(box a, box b);
 };
 
-}
+}}
 #endif //LAYER_H

include/tkDNN/Network.h

@@ -3,7 +3,7 @@
 
 #include "utils.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 /**
     Data rapresentation beetween layers
@@ -59,8 +59,8 @@ public:
     dataDim_t input_dim;
     dataDim_t getOutputDim();
 
-    bool fp16;
+    bool fp16, dla;
 };
 
-}
+}}
 #endif //NETWORK_H

include/tkDNN/NetworkRT.h

@@ -1,12 +1,47 @@
 #ifndef NETWORKRT_H
 #define NETWORKRT_H
 
+#include <string.h> // memcpy
 #include "utils.h"
 #include "Network.h"
 #include "Layer.h"
 #include "NvInfer.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
+
+template<typename T> void writeBUF(char*& buffer, const T& val)
+{
+    *reinterpret_cast<T*>(buffer) = val;
+    buffer += sizeof(T);
+}
+
+template<typename T> T readBUF(const char*& buffer)
+{
+    T val = *reinterpret_cast<const T*>(buffer);
+    buffer += sizeof(T);
+    return val;
+}
+
+using namespace nvinfer1;
+#include "pluginsRT/ActivationLeakyRT.h"
+#include "pluginsRT/ReorgRT.h"
+#include "pluginsRT/RegionRT.h"
+//#include "pluginsRT/RouteRT.h"
+#include "pluginsRT/ShortcutRT.h"
+#include "pluginsRT/YoloRT.h"
+#include "pluginsRT/UpsampleRT.h"
+//#include "pluginsRT/Int8Calibrator.h"
+
+class PluginFactory : IPluginFactory
+{
+public:
+    YoloRT *yolos[16];
+    int n_yolos;
+
+	virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength);
+};
+
+
 
 class NetworkRT {
 
@@ -18,13 +53,17 @@ public:
     
     nvinfer1::ICudaEngine *engineRT;
     nvinfer1::IExecutionContext *contextRT;
-    void* buffersRT[2];
+
+    const static int MAX_BUFFERS_RT = 10;
+    void* buffersRT[MAX_BUFFERS_RT];
     int buf_input_idx, buf_output_idx;
 
     dataDim_t input_dim, output_dim;
     dnnType *output;
     cudaStream_t stream;
 
+    PluginFactory *pluginFactory;
+
     NetworkRT(Network *net, const char *name);
     virtual ~NetworkRT();
 
@@ -43,24 +82,13 @@ public:
     nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Route *l);
     nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Reorg *l);
     nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Region *l);
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Shortcut *l);
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Yolo *l);
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Upsample *l);
 
     bool serialize(const char *filename);
     bool deserialize(const char *filename);
 };
 
-
-template<typename T> void writeBUF(char*& buffer, const T& val)
-{
-    *reinterpret_cast<T*>(buffer) = val;
-    buffer += sizeof(T);
-}
-
-template<typename T> T readBUF(const char*& buffer)
-{
-    T val = *reinterpret_cast<const T*>(buffer);
-    buffer += sizeof(T);
-    return val;
-}
-
-}
+}}
 #endif //NETWORKRT_H

include/tkDNN/Yolo3Detection.h

@@ -0,0 +1,67 @@
+#include <iostream>
+#include <signal.h>
+#include <stdlib.h>     /* srand, rand */
+#include <unistd.h>
+#include <mutex>
+#include "utils.h"
+
+#include <opencv2/core/core.hpp>
+#include <opencv2/highgui/highgui.hpp>
+#include <opencv2/imgproc/imgproc.hpp>
+
+#include "tkdnn.h"
+
+namespace tk { namespace dnn {
+
+/**
+ * 
+ * @author Francesco Gatti
+ */
+class Yolo3Detection {
+
+    private:
+        tk::dnn::NetworkRT *netRT = nullptr;
+        tk::dnn::Yolo* yolo[3];
+        dnnType *input, *input_d;
+
+        int ndets = 0;
+        tk::dnn::Yolo::detection *dets = nullptr;
+
+        cv::Mat imageF;
+        cv::Mat bgr[3]; 
+
+    public:
+        int classes = 0;
+        int num = 0;
+        float thresh = 0.3;
+        cv::Scalar colors[256];
+
+        // this is filled with results
+        std::vector<tk::dnn::box> detected;
+
+        // keep track of inference times (ms)
+        std::vector<double> stats;
+
+        Yolo3Detection() {}
+
+        virtual ~Yolo3Detection() {}
+
+        /**
+         * Method used for inizialize the class
+         * 
+         * @return Success of the initialization
+         */
+        bool init(std::string tensor_path);
+
+        void update(cv::Mat &frame);
+
+        tk::dnn::Yolo* getYoloLayer(int n=0) {
+            if(n<3)
+                return yolo[n];
+            else 
+                return nullptr;
+        }
+
+};
+
+}}

include/tkDNN/kernels.h

@@ -7,11 +7,21 @@ void activationELUForward(dnnType* srcData, dnnType* dstData, int size, cudaStre
 void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
 void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
 
+void fill(dnnType* data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0));
+
 void reorgForward(  dnnType* srcData, dnnType* dstData, 
                     int n, int c, int h, int w, int stride, cudaStream_t stream = cudaStream_t(0));
 void softmaxForward(float *input, int n, int batch, int batch_offset, 
                     int groups, int group_offset, int stride, float temp, float *output, cudaStream_t stream = cudaStream_t(0));
 
 
+void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int w1, int s1,
+                                                         int n2, int c2, int h2, int w2, int s2, 
+                     cudaStream_t stream = cudaStream_t(0));
+
+void upsampleForward(dnnType* srcData, dnnType* dstData, 
+                     int n, int c, int h, int w, int s, int forward, float scale,                                   
+                     cudaStream_t stream = cudaStream_t(0));
+
 void float2half(float* srcData, __half* dstData, int size, const cudaStream_t stream = cudaStream_t(0));
 #endif //KERNELS_H

include/tkDNN/models/Yolo3.h

@@ -0,0 +1,289 @@
+int preYoloFilters = (classes+5)*3; 
+
+std::string input_bin = bin_path + "/layers/input.bin";
+std::vector<std::string> output_bins = {
+    bin_path + "/debug/layer82_out.bin",
+    bin_path + "/debug/layer94_out.bin",
+    bin_path + "/debug/layer106_out.bin"
+};
+std::string c0_bin    = bin_path + "/layers/c0.bin";
+std::string c1_bin    = bin_path + "/layers/c1.bin";
+std::string c2_bin    = bin_path + "/layers/c2.bin";
+std::string c3_bin    = bin_path + "/layers/c3.bin";
+std::string c5_bin    = bin_path + "/layers/c5.bin";
+std::string c6_bin    = bin_path + "/layers/c6.bin";
+std::string c7_bin    = bin_path + "/layers/c7.bin";
+std::string c9_bin    = bin_path + "/layers/c9.bin";
+std::string c10_bin   = bin_path + "/layers/c10.bin";
+std::string c12_bin   = bin_path + "/layers/c12.bin";
+std::string c13_bin   = bin_path + "/layers/c13.bin";
+std::string c14_bin   = bin_path + "/layers/c14.bin";
+std::string c16_bin   = bin_path + "/layers/c16.bin";
+std::string c17_bin   = bin_path + "/layers/c17.bin";
+std::string c19_bin   = bin_path + "/layers/c19.bin";
+std::string c20_bin   = bin_path + "/layers/c20.bin";
+std::string c22_bin   = bin_path + "/layers/c22.bin";
+std::string c23_bin   = bin_path + "/layers/c23.bin";
+std::string c25_bin   = bin_path + "/layers/c25.bin";
+std::string c26_bin   = bin_path + "/layers/c26.bin";
+std::string c28_bin   = bin_path + "/layers/c28.bin";
+std::string c29_bin   = bin_path + "/layers/c29.bin";
+std::string c31_bin   = bin_path + "/layers/c31.bin";
+std::string c32_bin   = bin_path + "/layers/c32.bin";
+std::string c34_bin   = bin_path + "/layers/c34.bin";
+std::string c35_bin   = bin_path + "/layers/c35.bin";
+std::string c37_bin   = bin_path + "/layers/c37.bin";
+std::string c38_bin   = bin_path + "/layers/c38.bin";
+std::string c39_bin   = bin_path + "/layers/c39.bin";
+std::string c41_bin   = bin_path + "/layers/c41.bin";
+std::string c42_bin   = bin_path + "/layers/c42.bin";
+std::string c44_bin   = bin_path + "/layers/c44.bin";
+std::string c45_bin   = bin_path + "/layers/c45.bin";
+std::string c47_bin   = bin_path + "/layers/c47.bin";
+std::string c48_bin   = bin_path + "/layers/c48.bin";
+std::string c50_bin   = bin_path + "/layers/c50.bin";
+std::string c51_bin   = bin_path + "/layers/c51.bin";
+std::string c53_bin   = bin_path + "/layers/c53.bin";
+std::string c54_bin   = bin_path + "/layers/c54.bin";
+std::string c56_bin   = bin_path + "/layers/c56.bin";
+std::string c57_bin   = bin_path + "/layers/c57.bin";
+std::string c59_bin   = bin_path + "/layers/c59.bin";
+std::string c60_bin   = bin_path + "/layers/c60.bin";
+std::string c62_bin   = bin_path + "/layers/c62.bin";
+std::string c63_bin   = bin_path + "/layers/c63.bin";
+std::string c64_bin   = bin_path + "/layers/c64.bin";
+std::string c66_bin   = bin_path + "/layers/c66.bin";
+std::string c67_bin   = bin_path + "/layers/c67.bin";
+std::string c69_bin   = bin_path + "/layers/c69.bin";
+std::string c70_bin   = bin_path + "/layers/c70.bin";
+std::string c72_bin   = bin_path + "/layers/c72.bin";
+std::string c73_bin   = bin_path + "/layers/c73.bin";
+std::string c75_bin   = bin_path + "/layers/c75.bin";
+std::string c76_bin   = bin_path + "/layers/c76.bin";
+std::string c77_bin   = bin_path + "/layers/c77.bin";
+std::string c78_bin   = bin_path + "/layers/c78.bin";
+std::string c79_bin   = bin_path + "/layers/c79.bin";
+std::string c80_bin   = bin_path + "/layers/c80.bin";
+std::string c81_bin   = bin_path + "/layers/c81.bin";
+std::string g82_bin   = bin_path + "/layers/g82.bin";
+std::string c84_bin   = bin_path + "/layers/c84.bin";
+std::string c87_bin   = bin_path + "/layers/c87.bin";
+std::string c88_bin   = bin_path + "/layers/c88.bin";
+std::string c89_bin   = bin_path + "/layers/c89.bin";
+std::string c90_bin   = bin_path + "/layers/c90.bin";
+std::string c91_bin   = bin_path + "/layers/c91.bin";
+std::string c92_bin   = bin_path + "/layers/c92.bin";
+std::string c93_bin   = bin_path + "/layers/c93.bin";
+std::string g94_bin   = bin_path + "/layers/g94.bin";
+std::string c96_bin   = bin_path + "/layers/c96.bin";
+std::string c99_bin   = bin_path + "/layers/c99.bin";
+std::string c100_bin  = bin_path + "/layers/c100.bin";
+std::string c101_bin  = bin_path + "/layers/c101.bin";
+std::string c102_bin  = bin_path + "/layers/c102.bin";
+std::string c103_bin  = bin_path + "/layers/c103.bin";
+std::string c104_bin  = bin_path + "/layers/c104.bin";
+std::string c105_bin  = bin_path + "/layers/c105.bin";
+std::string g106_bin  = bin_path + "/layers/g106.bin";
+
+tk::dnn::Conv2d     c0   (&net,  32, 3, 3, 1, 1, 1, 1,  c0_bin, true);
+tk::dnn::Activation a0   (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c1   (&net,  64, 3, 3, 2, 2, 1, 1,  c1_bin, true);
+tk::dnn::Activation a1   (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c2   (&net,  32, 1, 1, 1, 1, 0, 0,  c2_bin, true);
+tk::dnn::Activation a2   (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c3   (&net,  64, 3, 3, 1, 1, 1, 1,  c3_bin, true);
+tk::dnn::Activation a3   (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s4   (&net, &a1);
+tk::dnn::Conv2d     c5   (&net, 128, 3, 3, 2, 2, 1, 1,  c5_bin, true);
+tk::dnn::Activation a5   (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c6   (&net,  64, 1, 1, 1, 1, 0, 0,  c6_bin, true);
+tk::dnn::Activation a6   (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c7   (&net, 128, 3, 3, 1, 1, 1, 1,  c7_bin, true);
+tk::dnn::Activation a7   (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s8   (&net, &a5);
+tk::dnn::Conv2d     c9   (&net,  64, 1, 1, 1, 1, 0, 0,  c9_bin, true);
+tk::dnn::Activation a9   (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c10  (&net, 128, 3, 3, 1, 1, 1, 1, c10_bin, true);
+tk::dnn::Activation a10  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s11  (&net, &s8);
+
+tk::dnn::Conv2d     c12  (&net, 256, 3, 3, 2, 2, 1, 1, c12_bin, true);
+tk::dnn::Activation a12  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c13  (&net, 128, 1, 1, 1, 1, 0, 0, c13_bin, true);
+tk::dnn::Activation a13  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c14  (&net, 256, 3, 3, 1, 1, 1, 1, c14_bin, true);
+tk::dnn::Activation a14  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s15  (&net, &a12);
+
+tk::dnn::Conv2d     c16  (&net, 128, 1, 1, 1, 1, 0, 0, c16_bin, true);
+tk::dnn::Activation a16  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c17  (&net, 256, 3, 3, 1, 1, 1, 1, c17_bin, true);
+tk::dnn::Activation a17  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s18  (&net, &s15);
+tk::dnn::Conv2d     c19  (&net, 128, 1, 1, 1, 1, 0, 0, c19_bin, true);
+tk::dnn::Activation a19  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c20  (&net, 256, 3, 3, 1, 1, 1, 1, c20_bin, true);
+tk::dnn::Activation a20  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s21  (&net, &s18);
+tk::dnn::Conv2d     c22  (&net, 128, 1, 1, 1, 1, 0, 0, c22_bin, true);
+tk::dnn::Activation a22  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c23  (&net, 256, 3, 3, 1, 1, 1, 1, c23_bin, true);
+tk::dnn::Activation a23  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s24  (&net, &s21);
+tk::dnn::Conv2d     c25  (&net, 128, 1, 1, 1, 1, 0, 0, c25_bin, true);
+tk::dnn::Activation a25  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c26  (&net, 256, 3, 3, 1, 1, 1, 1, c26_bin, true);
+tk::dnn::Activation a26  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s27  (&net, &s24);
+tk::dnn::Conv2d     c28  (&net, 128, 1, 1, 1, 1, 0, 0, c28_bin, true);
+tk::dnn::Activation a28  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c29  (&net, 256, 3, 3, 1, 1, 1, 1, c29_bin, true);
+tk::dnn::Activation a29  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s30  (&net, &s27);
+tk::dnn::Conv2d     c31  (&net, 128, 1, 1, 1, 1, 0, 0, c31_bin, true);
+tk::dnn::Activation a31  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c32  (&net, 256, 3, 3, 1, 1, 1, 1, c32_bin, true);
+tk::dnn::Activation a32  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s33  (&net, &s30);
+tk::dnn::Conv2d     c34  (&net, 128, 1, 1, 1, 1, 0, 0, c34_bin, true);
+tk::dnn::Activation a34  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c35  (&net, 256, 3, 3, 1, 1, 1, 1, c35_bin, true);
+tk::dnn::Activation a35  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s36  (&net, &s33);
+
+tk::dnn::Conv2d     c37  (&net, 512, 3, 3, 2, 2, 1, 1, c37_bin, true);
+tk::dnn::Activation a37  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c38  (&net, 256, 1, 1, 1, 1, 0, 0, c38_bin, true);
+tk::dnn::Activation a38  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c39  (&net, 512, 3, 3, 1, 1, 1, 1, c39_bin, true);
+tk::dnn::Activation a39  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s40  (&net, &a37);
+
+tk::dnn::Conv2d     c41  (&net, 256, 1, 1, 1, 1, 0, 0, c41_bin, true);
+tk::dnn::Activation a41  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c42  (&net, 512, 3, 3, 1, 1, 1, 1, c42_bin, true);
+tk::dnn::Activation a42  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s43  (&net, &s40);
+tk::dnn::Conv2d     c44  (&net, 256, 1, 1, 1, 1, 0, 0, c44_bin, true);
+tk::dnn::Activation a44  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c45  (&net, 512, 3, 3, 1, 1, 1, 1, c45_bin, true);
+tk::dnn::Activation a45  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s46  (&net, &s43);
+tk::dnn::Conv2d     c47  (&net, 256, 1, 1, 1, 1, 0, 0, c47_bin, true);
+tk::dnn::Activation a47  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c48  (&net, 512, 3, 3, 1, 1, 1, 1, c48_bin, true);
+tk::dnn::Activation a48  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s49  (&net, &s46);
+tk::dnn::Conv2d     c50  (&net, 256, 1, 1, 1, 1, 0, 0, c50_bin, true);
+tk::dnn::Activation a50  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c51  (&net, 512, 3, 3, 1, 1, 1, 1, c51_bin, true);
+tk::dnn::Activation a51  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s52  (&net, &s49);
+tk::dnn::Conv2d     c53  (&net, 256, 1, 1, 1, 1, 0, 0, c53_bin, true);
+tk::dnn::Activation a53  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c54  (&net, 512, 3, 3, 1, 1, 1, 1, c54_bin, true);
+tk::dnn::Activation a54  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s55  (&net, &s52);
+tk::dnn::Conv2d     c56  (&net, 256, 1, 1, 1, 1, 0, 0, c56_bin, true);
+tk::dnn::Activation a56  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c57  (&net, 512, 3, 3, 1, 1, 1, 1, c57_bin, true);
+tk::dnn::Activation a57  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s58  (&net, &s55);
+tk::dnn::Conv2d     c59  (&net, 256, 1, 1, 1, 1, 0, 0, c59_bin, true);
+tk::dnn::Activation a59  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c60  (&net, 512, 3, 3, 1, 1, 1, 1, c60_bin, true);
+tk::dnn::Activation a60  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s61  (&net, &s58);
+
+tk::dnn::Conv2d     c62  (&net,1024, 3, 3, 2, 2, 1, 1, c62_bin, true);
+tk::dnn::Activation a62  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c63  (&net, 512, 1, 1, 1, 1, 0, 0, c63_bin, true);
+tk::dnn::Activation a63  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c64  (&net,1024, 3, 3, 1, 1, 1, 1, c64_bin, true);
+tk::dnn::Activation a64  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s65  (&net, &a62);
+
+tk::dnn::Conv2d     c66  (&net, 512, 1, 1, 1, 1, 0, 0, c66_bin, true);
+tk::dnn::Activation a66  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c67  (&net,1024, 3, 3, 1, 1, 1, 1, c67_bin, true);
+tk::dnn::Activation a67  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s68  (&net, &s65);
+
+tk::dnn::Conv2d     c69  (&net, 512, 1, 1, 1, 1, 0, 0, c69_bin, true);
+tk::dnn::Activation a69  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c70  (&net,1024, 3, 3, 1, 1, 1, 1, c70_bin, true);
+tk::dnn::Activation a70  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s71  (&net, &s68);
+
+tk::dnn::Conv2d     c72  (&net, 512, 1, 1, 1, 1, 0, 0, c72_bin, true);
+tk::dnn::Activation a72  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c73  (&net,1024, 3, 3, 1, 1, 1, 1, c73_bin, true);
+tk::dnn::Activation a73  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Shortcut   s74  (&net, &s71);
+
+tk::dnn::Conv2d     c75  (&net, 512, 1, 1, 1, 1, 0, 0, c75_bin, true);
+tk::dnn::Activation a75  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c76  (&net,1024, 3, 3, 1, 1, 1, 1, c76_bin, true);
+tk::dnn::Activation a76  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c77  (&net, 512, 1, 1, 1, 1, 0, 0, c77_bin, true);
+tk::dnn::Activation a77  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c78  (&net,1024, 3, 3, 1, 1, 1, 1, c78_bin, true);
+tk::dnn::Activation a78  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c79  (&net, 512, 1, 1, 1, 1, 0, 0, c79_bin, true);
+tk::dnn::Activation a79  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c80  (&net,1024, 3, 3, 1, 1, 1, 1, c80_bin, true);
+tk::dnn::Activation a80  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c81  (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c81_bin, false);
+tk::dnn::Yolo     yolo0  (&net, classes, 3, g82_bin);
+
+tk::dnn::Layer *m83_layers[1] = { &a79 };
+tk::dnn::Route      m83  (&net, m83_layers, 1);
+tk::dnn::Conv2d     c84  (&net, 256, 1, 1, 1, 1, 0, 0, c84_bin, true);
+tk::dnn::Activation a84  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Upsample   u85  (&net, 2);
+
+tk::dnn::Layer *m86_layers[2] = { &u85, &s61 };
+tk::dnn::Route      m86  (&net, m86_layers, 2);
+tk::dnn::Conv2d     c87  (&net, 256, 1, 1, 1, 1, 0, 0, c87_bin, true);
+tk::dnn::Activation a87  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c88  (&net, 512, 3, 3, 1, 1, 1, 1, c88_bin, true);
+tk::dnn::Activation a88  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c89  (&net, 256, 1, 1, 1, 1, 0, 0, c89_bin, true);
+tk::dnn::Activation a89  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c90  (&net, 512, 3, 3, 1, 1, 1, 1, c90_bin, true);
+tk::dnn::Activation a90  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c91  (&net, 256, 1, 1, 1, 1, 0, 0, c91_bin, true);
+tk::dnn::Activation a91  (&net, tk::dnn::ACTIVATION_LEAKY);
+
+tk::dnn::Conv2d     c92  (&net, 512, 3, 3, 1, 1, 1, 1, c92_bin, true);
+tk::dnn::Activation a92  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c93  (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c93_bin, false);
+tk::dnn::Yolo     yolo1  (&net, classes, 3, g94_bin);
+
+tk::dnn::Layer *m95_layers[1] = { &a91 };
+tk::dnn::Route      m95  (&net, m95_layers, 1);
+tk::dnn::Conv2d     c96  (&net, 128, 1, 1, 1, 1, 0, 0, c96_bin, true);
+tk::dnn::Activation a96  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Upsample   u97  (&net, 2);
+
+tk::dnn::Layer *m98_layers[2] = { &u97, &s36 };
+tk::dnn::Route      m98  (&net, m98_layers, 2);
+tk::dnn::Conv2d     c99  (&net, 128, 1, 1, 1, 1, 0, 0, c99_bin, true);
+tk::dnn::Activation a99  (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c100 (&net, 256, 3, 3, 1, 1, 1, 1, c100_bin, true);
+tk::dnn::Activation a100 (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c101 (&net, 128, 1, 1, 1, 1, 0, 0, c101_bin, true);
+tk::dnn::Activation a101 (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c102 (&net, 256, 3, 3, 1, 1, 1, 1, c102_bin, true);
+tk::dnn::Activation a102 (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c103 (&net, 128, 1, 1, 1, 1, 0, 0, c103_bin, true);
+tk::dnn::Activation a103 (&net, tk::dnn::ACTIVATION_LEAKY);
+
+tk::dnn::Conv2d     c104 (&net, 256, 3, 3, 1, 1, 1, 1, c104_bin, true);
+tk::dnn::Activation a104 (&net, tk::dnn::ACTIVATION_LEAKY);
+tk::dnn::Conv2d     c105 (&net, preYoloFilters, 1, 1, 1, 1, 0, 0, c105_bin, false);
+tk::dnn::Yolo      yolo2 (&net, classes, 3, g106_bin);
+
+yolo[0] = &yolo0;
+yolo[1] = &yolo1;
+yolo[2] = &yolo2;

include/tkDNN/pluginsRT/ActivationLeakyRT.h

@@ -1,5 +1,5 @@
 #include<cassert>
-#include "kernels.h"
+#include "../kernels.h"
 
 class ActivationLeakyRT : public IPlugin {
 
@@ -53,7 +53,7 @@ public:
 
 	virtual void serialize(void* buffer) override {
 		char *buf = reinterpret_cast<char*>(buffer);
-		tkDNN::writeBUF(buf, size);
+		tk::dnn::writeBUF(buf, size);
 	}
 
 	int size;

include/tkDNN/pluginsRT/Int8Calibrator.h

@@ -8,7 +8,7 @@
 class BatchStream
 {
 public:
-	BatchStream(tkDNN::dataDim_t dim, int batchSize, int maxBatches)
+	BatchStream(tk::dnn::dataDim_t dim, int batchSize, int maxBatches)
 	{
         mBatchSize = batchSize;
         mMaxBatches = maxBatches;

include/tkDNN/pluginsRT/RegionRT.h

@@ -1,5 +1,5 @@
 #include<cassert>
-#include "kernels.h"
+#include "../kernels.h"
 
 class RegionRT : public IPlugin {
 
@@ -69,17 +69,17 @@ public:
 
 
 	virtual size_t getSerializationSize() override {
-		return 6*sizeof(int) + 1*sizeof(float);
+		return 6*sizeof(int);
 	}
 
 	virtual void serialize(void* buffer) override {
 		char *buf = reinterpret_cast<char*>(buffer);
-		tkDNN::writeBUF(buf, classes);
-		tkDNN::writeBUF(buf, coords);
-		tkDNN::writeBUF(buf, num);
-		tkDNN::writeBUF(buf, c);
-		tkDNN::writeBUF(buf, h);
-		tkDNN::writeBUF(buf, w);
+		tk::dnn::writeBUF(buf, classes);
+		tk::dnn::writeBUF(buf, coords);
+		tk::dnn::writeBUF(buf, num);
+		tk::dnn::writeBUF(buf, c);
+		tk::dnn::writeBUF(buf, h);
+		tk::dnn::writeBUF(buf, w);
 	}
 
 	int c, h, w;

include/tkDNN/pluginsRT/ReorgRT.h

@@ -1,5 +1,5 @@
 #include<cassert>
-#include "kernels.h"
+#include "../kernels.h"
 
 class ReorgRT : public IPlugin {
 
@@ -53,10 +53,10 @@ public:
 
 	virtual void serialize(void* buffer) override {
 		char *buf = reinterpret_cast<char*>(buffer);
-		tkDNN::writeBUF(buf, stride);
-		tkDNN::writeBUF(buf, c);
-		tkDNN::writeBUF(buf, h);
-		tkDNN::writeBUF(buf, w);
+		tk::dnn::writeBUF(buf, stride);
+		tk::dnn::writeBUF(buf, c);
+		tk::dnn::writeBUF(buf, h);
+		tk::dnn::writeBUF(buf, w);
 	}
 
 	int c, h, w, stride;

include/tkDNN/pluginsRT/RouteRT.h

@@ -0,0 +1,82 @@
+#include<cassert>
+#include "../kernels.h"
+
+class RouteRT : public IPlugin {
+
+public:
+	RouteRT() {
+	}
+
+	~RouteRT(){
+
+	}
+
+	int getNbOutputs() const override {
+		return 1;
+	}
+
+	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]};
+	}
+
+	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+		in = nbInputs;
+		c = 0;
+		for(int i=0; i<nbInputs; i++) {
+			c_in[i] = inputDims[i].d[0]; 
+			c += inputDims[i].d[0];
+		}
+		h = inputDims[0].d[1];
+		w = inputDims[0].d[2];
+	}
+
+	int initialize() override {
+
+		return 0;
+	}
+
+	virtual void terminate() override {
+	}
+
+	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+		return 0;
+	}
+
+	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;
+		}
+
+		return 0;
+	}
+
+
+	virtual size_t getSerializationSize() override {
+		return (4+MAX_INPUTS)*sizeof(int);
+	}
+
+	virtual void serialize(void* buffer) override {
+		char *buf = reinterpret_cast<char*>(buffer);
+		tk::dnn::writeBUF(buf, in);
+		for(int i=0; i<MAX_INPUTS; i++)
+			tk::dnn::writeBUF(buf, c_in[i]);
+
+		tk::dnn::writeBUF(buf, c);
+		tk::dnn::writeBUF(buf, h);
+		tk::dnn::writeBUF(buf, w);
+	}
+
+	static const int MAX_INPUTS = 4;
+	int in;
+	int c_in[MAX_INPUTS];
+	int c, h, w;
+};

include/tkDNN/pluginsRT/ShortcutRT.h

@@ -0,0 +1,65 @@
+#include<cassert>
+#include "../kernels.h"
+
+class ShortcutRT : public IPlugin {
+
+public:
+	ShortcutRT() {
+	}
+
+	~ShortcutRT(){
+
+	}
+
+	int getNbOutputs() const override {
+		return 1;
+	}
+
+	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+		return DimsCHW{inputs[0].d[0], inputs[0].d[1], inputs[0].d[2]};
+	}
+
+	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+		c = inputDims[0].d[0];
+		h = inputDims[0].d[1];
+		w = inputDims[0].d[2];
+	}
+
+	int initialize() override {
+
+		return 0;
+	}
+
+	virtual void terminate() override {
+	}
+
+	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+		return 0;
+	}
+
+	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+
+		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+		dnnType *srcDataBack = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
+		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+
+		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
+		shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, c, h, w, 1, stream);
+
+		return 0;
+	}
+
+
+	virtual size_t getSerializationSize() override {
+		return 3*sizeof(int);
+	}
+
+	virtual void serialize(void* buffer) override {
+		char *buf = reinterpret_cast<char*>(buffer);
+		tk::dnn::writeBUF(buf, c);
+		tk::dnn::writeBUF(buf, h);
+		tk::dnn::writeBUF(buf, w);
+	}
+
+	int c, h, w;
+};

include/tkDNN/pluginsRT/UpsampleRT.h

@@ -0,0 +1,65 @@
+#include<cassert>
+#include "../kernels.h"
+
+class UpsampleRT : public IPlugin {
+
+public:
+	UpsampleRT(int stride) {
+		this->stride = stride;
+	}
+
+	~UpsampleRT(){
+
+	}
+
+	int getNbOutputs() const override {
+		return 1;
+	}
+
+	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+		return DimsCHW(inputs[0].d[0], inputs[0].d[1]*stride, inputs[0].d[2]*stride);
+	}
+
+	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+		c = inputDims[0].d[0];
+		h = inputDims[0].d[1];
+		w = inputDims[0].d[2];
+	}
+
+	int initialize() override {
+
+		return 0;
+	}
+
+	virtual void terminate() override {
+	}
+
+	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+		return 0;
+	}
+
+	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+
+		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+	    
+		fill(dstData, batchSize*c*h*w*stride*stride, 0.0, stream);
+    	upsampleForward(srcData, dstData, batchSize, c, h, w, stride, 1, 1, stream);
+		return 0;
+	}
+
+
+	virtual size_t getSerializationSize() override {
+		return 4*sizeof(int);
+	}
+
+	virtual void serialize(void* buffer) override {
+		char *buf = reinterpret_cast<char*>(buffer);
+		tk::dnn::writeBUF(buf, stride);
+		tk::dnn::writeBUF(buf, c);
+		tk::dnn::writeBUF(buf, h);
+		tk::dnn::writeBUF(buf, w);
+	}
+
+	int c, h, w, stride;
+};

include/tkDNN/pluginsRT/YoloRT.h

@@ -0,0 +1,116 @@
+#include<cassert>
+#include "../kernels.h"
+
+#define YOLORT_CLASSNAME_W 256
+
+class YoloRT : public IPlugin {
+
+
+
+public:
+	YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr) {
+
+		this->classes = classes;
+		this->num = num;
+
+        mask = new dnnType[num];
+        bias = new dnnType[num*3*2];
+        if(yolo != nullptr) {
+            memcpy(mask, yolo->mask_h, sizeof(dnnType)*num);
+            memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*3*2);
+			classesNames = yolo->classesNames;
+        }
+	}
+
+	~YoloRT(){
+
+	}
+
+	int getNbOutputs() const override {
+		return 1;
+	}
+
+	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+		return inputs[0];
+	}
+
+	void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override {
+		c = inputDims[0].d[0];
+		h = inputDims[0].d[1];
+		w = inputDims[0].d[2];
+	}
+
+	int initialize() override {
+
+		return 0;
+	}
+
+	virtual void terminate() override {
+	}
+
+	virtual size_t getWorkspaceSize(int maxBatchSize) const override {
+		return 0;
+	}
+
+	virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override {
+
+		dnnType *srcData = (dnnType*)reinterpret_cast<const dnnType*>(inputs[0]);
+		dnnType *dstData = reinterpret_cast<dnnType*>(outputs[0]);
+
+		checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
+
+		for (int b = 0; b < batchSize; ++b){
+			for(int n = 0; n < num; ++n){
+				int index = entry_index(b, n*w*h, 0, batchSize);
+				activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream);
+				
+				index = entry_index(b, n*w*h, 4, batchSize);
+				activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream);
+			}
+		}
+
+		//std::cout<<"YOLO END\n";
+		return 0;
+	}
+
+
+	virtual size_t getSerializationSize() override {
+		return 5*sizeof(int) + num*sizeof(dnnType) + num*3*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char);
+	}
+
+	virtual void serialize(void* buffer) override {
+		char *buf = reinterpret_cast<char*>(buffer);
+		tk::dnn::writeBUF(buf, classes);
+		tk::dnn::writeBUF(buf, num);
+		tk::dnn::writeBUF(buf, c);
+		tk::dnn::writeBUF(buf, h);
+		tk::dnn::writeBUF(buf, w);
+        for(int i=0; i<num; i++)
+    		tk::dnn::writeBUF(buf, mask[i]);
+        for(int i=0; i<3*2*num; i++)
+    		tk::dnn::writeBUF(buf, bias[i]);
+
+		// save classes names
+		for(int i=0; i<classes; i++) {
+			char tmp[YOLORT_CLASSNAME_W];
+			strcpy(tmp, classesNames[i].c_str());
+			for(int j=0; j<YOLORT_CLASSNAME_W; j++) {
+				tk::dnn::writeBUF(buf, tmp[j]);
+			}
+		}
+	}
+
+	int c, h, w;
+    int classes, num;
+	std::vector<std::string> classesNames;
+
+    dnnType *mask;
+    dnnType *bias;
+
+	int entry_index(int batch, int location, int entry, int batchSize) {
+		int n =   location / (w*h);
+		int loc = location % (w*h);
+		return batch*c*h*w*batchSize + n*w*h*(4+classes+1) + entry*w*h + loc;
+	}
+
+};

include/tkDNN/tkdnn.h

@@ -5,4 +5,4 @@
 #include "Layer.h"
 #include "NetworkRT.h"
 
-#define TKDNN_VERSION 200
+#define TKDNN_VERSION 400

include/tkDNN/utils.h

@@ -90,7 +90,7 @@
 
 void printCenteredTitle(const char *title, char fill, int dim);
 bool fileExist(const char *fname);
-void readBinaryFile(const char* fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0);
+void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek = 0);
 int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device = true);
 void printDeviceVector(int size, dnnType* vec_d, bool device = true);
 void resize(int size, dnnType **data);

src/Activation.cpp

@@ -3,7 +3,7 @@
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Activation::Activation(Network *net, int act_mode) : 
     Layer(net) {
@@ -47,7 +47,7 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
 
     if(act_mode == ACTIVATION_LEAKY) {
         activationLEAKYForward(srcData, dstData, dim.tot());
-    
+
     } else {
         dnnType alpha = dnnType(1);
         dnnType beta  = dnnType(0);
@@ -63,4 +63,4 @@ dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Conv2d.cpp

@@ -2,11 +2,11 @@
 
 #include "Layer.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW, 
                 int strideH, int strideW, int paddingH, int paddingW,
-                const char* fname_weights, bool batchnorm) : 
+                std::string fname_weights, bool batchnorm) : 
     
     LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, 
              fname_weights, batchnorm) {
@@ -117,7 +117,7 @@ dnnType* Conv2d::infer(dataDim_t &dim, dnnType* srcData) {
             dstTensorDesc, dstData, dstTensorDesc, 
             dstData, biasTensorDesc, //same tensor descriptor as bias 
             scales_d, bias_d, mean_d, variance_d, 
-            CUDNN_BN_MIN_EPSILON);
+            TKDNN_BN_MIN_EPSILON);
     }
     //update data dimensions    
     dim = output_dim;
@@ -125,4 +125,4 @@ dnnType* Conv2d::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Dense.cpp

@@ -2,9 +2,9 @@
 
 #include "Layer.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
-Dense::Dense(Network *net, int out_ch, const char* fname_weights) : 
+Dense::Dense(Network *net, int out_ch, std::string fname_weights) : 
     LayerWgs(net, net->getOutputDim().tot(), out_ch, 1, 1, 1, fname_weights) {
 
     output_dim.n = 1;
@@ -55,4 +55,4 @@ dnnType* Dense::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Flatten.cpp

@@ -3,7 +3,7 @@
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Flatten::Flatten(Network *net) : Layer(net) {
 
@@ -33,4 +33,4 @@ dnnType* Flatten::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Layer.cpp

@@ -2,19 +2,22 @@
 
 #include "Layer.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Layer::Layer(Network *net) {
 
     this->net = net;
-    this->input_dim = net->getOutputDim();
-    this->output_dim = input_dim;
-    
-    checkCUDNN( cudnnCreateTensorDescriptor(&srcTensorDesc) );
-    checkCUDNN( cudnnCreateTensorDescriptor(&dstTensorDesc) );
 
-    if(!net->addLayer(this))
-        FatalError("Net reached max number of layers");    
+    if(net != nullptr) {
+        this->input_dim = net->getOutputDim();
+        this->output_dim = input_dim;
+        
+        checkCUDNN( cudnnCreateTensorDescriptor(&srcTensorDesc) );
+        checkCUDNN( cudnnCreateTensorDescriptor(&dstTensorDesc) );
+
+        if(!net->addLayer(this))
+            FatalError("Net reached max number of layers");    
+    }
 }
 
 Layer::~Layer() {
@@ -23,4 +26,4 @@ Layer::~Layer() {
     checkCUDNN( cudnnDestroyTensorDescriptor(dstTensorDesc) );
 }
 
-}
+}}

src/LayerWgs.cpp

@@ -4,11 +4,11 @@
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 LayerWgs::LayerWgs(Network *net, int inputs, int outputs, 
                    int kh, int kw, int kl, 
-                   const char* fname_weights, bool batchnorm) : Layer(net) {
+                   std::string fname_weights, bool batchnorm) : Layer(net) {
 
     this->inputs  = inputs;
     this->outputs = outputs;    
@@ -29,7 +29,7 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs,
         seek += outputs;
         readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek);
 
-        float eps = CUDNN_BN_MIN_EPSILON;
+        float eps = TKDNN_BN_MIN_EPSILON;
 
         power_h = new dnnType[outputs];
         for(int i=0; i<outputs; i++) power_h[i] = 1.0f;
@@ -114,4 +114,4 @@ LayerWgs::~LayerWgs() {
     }
 }
 
-}
+}}

src/MulAdd.cpp

@@ -3,7 +3,7 @@
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 MulAdd::MulAdd(Network *net, dnnType mul, dnnType add) : Layer(net) {
 
@@ -41,4 +41,4 @@ dnnType* MulAdd::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Network.cpp

@@ -5,7 +5,7 @@
 #include "Network.h"
 #include "Layer.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Network::Network(dataDim_t input_dim) {
     this->input_dim = input_dim;
@@ -17,19 +17,28 @@ Network::Network(dataDim_t input_dim) {
              <<", CUDNN v"<<cu_ver<<")\n";
     dataType = CUDNN_DATA_FLOAT;
     tensorFormat = CUDNN_TENSOR_NCHW;
+    num_layers = 0;
+
+    fp16 = false;
+    dla = false;
+    if(const char* env_p = std::getenv("TKDNN_MODE")) {
+        if(strcmp(env_p, "FP16") == 0)
+            fp16 = true;
+	else if(strcmp(env_p, "DLA") == 0) {
+    	    dla = true;
+	    fp16 = true;	
+	}
+    }
+   
+    if(fp16)
+        std::cout<<COL_REDB<<"!! FP16 INERENCE ENABLED !!"<<COL_END<<"\n";
+    if(dla)
+        std::cout<<COL_GREENB<<"!! DLA INERENCE ENABLED !!"<<COL_END<<"\n";
+
 
     checkCUDNN( cudnnCreate(&cudnnHandle) );
     checkERROR( cublasCreate(&cublasHandle) );
 
-    num_layers = 0;
-
-    fp16 = false;
-    if(const char* env_p = std::getenv("TKDNN_MODE"))
-        if(strcmp(env_p, "FP16") == 0)
-            fp16 = true;
-   
-    if(fp16)
-        std::cout<<COL_REDB<<"!! FP16 INERENCE ENABLED !!"<<COL_END<<"\n";
 }
 
 Network::~Network() {
@@ -99,4 +108,4 @@ void Network::print() {
 }
 
 
-}
+}}

src/NetworkRT.cpp

@@ -11,10 +11,6 @@
 #include "NetworkRT.h"
 
 using namespace nvinfer1;
-#include "pluginsRT/ActivationLeakyRT.cpp"
-#include "pluginsRT/ReorgRT.cpp"
-#include "pluginsRT/RegionRT.cpp"
-#include "pluginsRT/Int8Calibrator.cpp"
 
 // Logger for info/warning/errors
 class Logger : public ILogger {
@@ -25,7 +21,7 @@ class Logger : public ILogger {
     }
 } loggerRT;
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 std::map<Layer*, nvinfer1::ITensor*>tensors; 
 
@@ -38,7 +34,8 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
   
     builderRT = createInferBuilder(loggerRT);
     std::cout<<"Float16 support: "<<builderRT->platformHasFastFp16()<<"\n";
-	std::cout<<"Int8 support: "<<builderRT->platformHasFastInt8()<<"\n";
+    std::cout<<"Int8 support: "<<builderRT->platformHasFastInt8()<<"\n";
+    std::cout<<"DLAs: "<<builderRT->getNbDLACores()<<"\n";
     networkRT = builderRT->createNetwork();
 
     if(!fileExist(name)) {
@@ -50,18 +47,18 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
         builderRT->setMaxBatchSize(1);
         builderRT->setMaxWorkspaceSize(1 << 30);
 
-        //change datatype based on system specs
-        if(builderRT->platformHasFastInt8()) {
-            BatchStream bstream({32,dim.c, dim.h, dim.w}, 32, 1);
-            Int8EntropyCalibrator calib(bstream, 0, false);
-            builderRT->setInt8Mode(true);
-            builderRT->setInt8Calibrator(&calib);
-
-        } else if(net->fp16 && builderRT->platformHasFastFp16()) {
+        if(net->fp16 && builderRT->platformHasFastFp16()) {
             dtRT = DataType::kHALF;
             builderRT->setHalf2Mode(true);
         }
-
+        if(net->dla && builderRT->getNbDLACores() > 0) {
+            dtRT = DataType::kHALF;
+            builderRT->setFp16Mode(true);
+            builderRT->allowGPUFallback(true);
+            builderRT->setDefaultDeviceType(DeviceType::kDLA);
+            builderRT->setDLACore(0);
+        }
+       
         //add input layer
         ITensor *input = networkRT->addInput("data", DataType::kFLOAT, 
                         DimsCHW{ dim.c, dim.h, dim.w});
@@ -74,6 +71,10 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
             Ilay->setName( (l->getLayerName() + std::to_string(i)).c_str() );
             
             input = Ilay->getOutput(0);
+            input->setName( (l->getLayerName() + std::to_string(i) + "_out").c_str() );
+            
+            if(l->getLayerType() == LAYER_YOLO)
+                networkRT->markOutput(*input);
             tensors[l] = input;
         }
         if(input == NULL)
@@ -96,9 +97,9 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
 	contextRT = engineRT->createExecutionContext();
 
 	// input and output buffer pointers that we pass to the engine - the engine requires exactly IEngine::getNbBindings(),
-	// of these, but in this case we know that there is exactly one input and one output.
-	if(engineRT->getNbBindings() != 2)
-        FatalError("Incorrect buffers number");
+	std::cout<<"Input/outputs numbers: "<<engineRT->getNbBindings()<<"\n";
+    if(engineRT->getNbBindings() > MAX_BUFFERS_RT)
+        FatalError("over RT buffer array size");
 
 	// In order to bind the buffers, we need to know the names of the input and output tensors.
 	// note that indices are guaranteed to be less than IEngine::getNbBindings()
@@ -119,10 +120,13 @@ NetworkRT::NetworkRT(Network *net, const char *name) {
     output_dim.c = oDim.d[0];
     output_dim.h = oDim.d[1];
     output_dim.w = oDim.d[2];
+    output_dim.print();
 	
     // create GPU buffers and a stream
-    checkCuda(cudaMalloc(&buffersRT[buf_input_idx],  input_dim.tot()*sizeof(dnnType)));
-    checkCuda(cudaMalloc(&buffersRT[buf_output_idx], output_dim.tot()*sizeof(dnnType)));
+    for(int i=0; i<engineRT->getNbBindings(); i++) {
+        Dims dim = engineRT->getBindingDimensions(i);
+        checkCuda(cudaMalloc(&buffersRT[i], dim.d[0]*dim.d[1]*dim.d[2]*sizeof(dnnType)));
+    }
     checkCuda(cudaMalloc(&output, output_dim.tot()*sizeof(dnnType)));
 	checkCuda(cudaStreamCreate(&stream));
 }
@@ -133,9 +137,9 @@ NetworkRT::~NetworkRT() {
 
 dnnType* NetworkRT::infer(dataDim_t &dim, dnnType* data) {
 
-    checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, input_dim.tot()*sizeof(float), cudaMemcpyDeviceToDevice, stream));
+    checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, input_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
     contextRT->enqueue(1, buffersRT, stream, nullptr);
-    checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], output_dim.tot()*sizeof(float), cudaMemcpyDeviceToDevice, stream));
+    checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream));
     cudaStreamSynchronize(stream);
 
     dim = output_dim;
@@ -167,6 +171,12 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
         return convert_layer(input, (Reorg*) l);
     if(type == LAYER_REGION)
         return convert_layer(input, (Region*) l);
+    if(type == LAYER_SHORTCUT)
+        return convert_layer(input, (Shortcut*) l);
+    if(type == LAYER_YOLO)
+        return convert_layer(input, (Yolo*) l);
+    if(type == LAYER_UPSAMPLE)
+        return convert_layer(input, (Upsample*) l);
 
     FatalError("Layer not implemented in tensorRT");
     return NULL;
@@ -264,10 +274,19 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
 
     if(l->act_mode == ACTIVATION_LEAKY) {
         //std::cout<<"New plugin LEAKY\n";
+        
+        /*
+        // plugin version
         IPlugin *plugin = new ActivationLeakyRT();
         IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
         checkNULL(lRT);
         return lRT;
+        */
+
+        IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kLEAKY_RELU);
+        lRT->setAlpha(0.1);
+        checkNULL(lRT);
+        return lRT;
 
     } else if(l->act_mode == CUDNN_ACTIVATION_RELU) {
         IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU);
@@ -292,10 +311,13 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Softmax *l) {
 ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) {
     //std::cout<<"convert route\n";
 
-    ITensor *tens[256];
-    for(int i=0; i<l->layers_n; i++)
+    ITensor **tens = new ITensor*[l->layers_n];
+    for(int i=0; i<l->layers_n; i++) {
         tens[i] = tensors[l->layers[i]];
+    }
     IConcatenationLayer *lRT = networkRT->addConcatenation(tens, l->layers_n);
+    //IPlugin *plugin = new RouteRT();
+    //IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin);
     checkNULL(lRT);
 
     return lRT;
@@ -321,6 +343,48 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Region *l) {
     return lRT;
 }
 
+ILayer* NetworkRT::convert_layer(ITensor *input, Shortcut *l) {
+    //std::cout<<"convert Shortcut\n";
+
+    //std::cout<<"New plugin Shortcut\n";
+    
+    ITensor *back_tens = tensors[l->backLayer];
+    /*
+    // plugin version
+    IPlugin *plugin = new ShortcutRT();
+    ITensor **inputs = new ITensor*[2];
+    inputs[0] = input;
+    inputs[1] = back_tens; 
+    IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin);
+    checkNULL(lRT);
+    */
+
+    IElementWiseLayer *lRT = networkRT->addElementWise(*input, *back_tens, ElementWiseOperation::kSUM);
+    checkNULL(lRT);
+    
+    return lRT;
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Yolo *l) {
+    //std::cout<<"convert Yolo\n";
+
+    //std::cout<<"New plugin YOLO\n";
+    IPlugin *plugin = new YoloRT(l->classes, l->num, l);
+    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    checkNULL(lRT);
+    return lRT;
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Upsample *l) {
+    //std::cout<<"convert Upsample\n";
+
+    //std::cout<<"New plugin UPSAMPLE\n";
+    IPlugin *plugin = new UpsampleRT(l->stride);
+    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    checkNULL(lRT);
+    return lRT;
+}
+
 bool NetworkRT::serialize(const char *filename) {
 
     std::ofstream p(filename);
@@ -338,44 +402,6 @@ bool NetworkRT::serialize(const char *filename) {
     return true;
 }
 
-class PluginFactory : IPluginFactory
-{
-public:
-	virtual IPlugin* createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
-        const char * buf = reinterpret_cast<const char*>(serialData);
-        
-        std::string name(layerName);
-
-        if(name.find("Activation") == 0) {
-            ActivationLeakyRT *a = new ActivationLeakyRT();
-            a->size = readBUF<int>(buf);
-            return a;
-        }
-
-        if(name.find("Region") == 0) {
-            RegionRT *r = new RegionRT(readBUF<int>(buf),    //classes
-                                       readBUF<int>(buf),    //coords
-                                       readBUF<int>(buf));   //num
-
-        	r->c = readBUF<int>(buf);
-		    r->h = readBUF<int>(buf);
-		    r->w = readBUF<int>(buf);
-            return r;
-        } 
-
-        if(name.find("Reorg") == 0) {
-            ReorgRT *r = new ReorgRT(readBUF<int>(buf)); //stride
-        	r->c = readBUF<int>(buf);
-		    r->h = readBUF<int>(buf);
-		    r->w = readBUF<int>(buf);
-            return r;
-        } 
-
-        FatalError("Cant deserialize Plugin");
-        return NULL;
-    }
-};
-
 bool NetworkRT::deserialize(const char *filename) {
 
     char *gieModelStream{nullptr};
@@ -390,13 +416,99 @@ bool NetworkRT::deserialize(const char *filename) {
         file.close();
     }
 
-    PluginFactory plfact;
-
+    pluginFactory = new PluginFactory();
     runtimeRT = createInferRuntime(loggerRT);
-    engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) &plfact);
+    engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) pluginFactory);
     //if (gieModelStream) delete [] gieModelStream;
 
     return true;
 }
 
+
+
+IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialData, size_t serialLength) {
+    const char * buf = reinterpret_cast<const char*>(serialData);
+    
+    std::string name(layerName);
+
+    if(name.find("Activation") == 0) {
+        ActivationLeakyRT *a = new ActivationLeakyRT();
+        a->size = readBUF<int>(buf);
+        return a;
+    }
+
+    if(name.find("Region") == 0) {
+        RegionRT *r = new RegionRT(readBUF<int>(buf),    //classes
+                                    readBUF<int>(buf),    //coords
+                                    readBUF<int>(buf));   //num
+
+        r->c = readBUF<int>(buf);
+        r->h = readBUF<int>(buf);
+        r->w = readBUF<int>(buf);
+        return r;
+    } 
+
+    if(name.find("Reorg") == 0) {
+        ReorgRT *r = new ReorgRT(readBUF<int>(buf)); //stride
+        r->c = readBUF<int>(buf);
+        r->h = readBUF<int>(buf);
+        r->w = readBUF<int>(buf);
+        return r;
+    } 
+
+    if(name.find("Shortcut") == 0) {
+        ShortcutRT *r = new ShortcutRT();
+        r->c = readBUF<int>(buf);
+        r->h = readBUF<int>(buf);
+        r->w = readBUF<int>(buf);
+        return r;
+    } 
+
+    if(name.find("Yolo") == 0) {
+        YoloRT *r = new YoloRT(readBUF<int>(buf),    //classes
+                                readBUF<int>(buf));   //num
+        r->c = readBUF<int>(buf);
+        r->h = readBUF<int>(buf);
+        r->w = readBUF<int>(buf);
+        for(int i=0; i<r->num; i++)
+            r->mask[i] = readBUF<dnnType>(buf);
+        for(int i=0; i<3*2*r->num; i++)
+            r->bias[i] = readBUF<dnnType>(buf);
+
+		// save classes names
+        r->classesNames.resize(r->classes);
+		for(int i=0; i<r->classes; i++) {
+            char tmp[YOLORT_CLASSNAME_W];
+			for(int j=0; j<YOLORT_CLASSNAME_W; j++)
+				tmp[j] = readBUF<char>(buf);
+            r->classesNames[i] = std::string(tmp);
+		}
+
+        yolos[n_yolos++] = r;
+        return r;
+    } 
+
+    if(name.find("Upsample") == 0) {
+        UpsampleRT *r = new UpsampleRT(readBUF<int>(buf)); //stride
+        r->c = readBUF<int>(buf);
+        r->h = readBUF<int>(buf);
+        r->w = readBUF<int>(buf);
+        return r;
+    } 
+/*
+    if(name.find("Route") == 0) {
+        RouteRT *r = new RouteRT();
+        r->in = readBUF<int>(buf);
+        for(int i=0; i<RouteRT::MAX_INPUTS; i++)
+            r->c_in[i] = readBUF<int>(buf);
+        r->c = readBUF<int>(buf);
+        r->h = readBUF<int>(buf);
+        r->w = readBUF<int>(buf);
+        return r;
+    } 
+*/
+    FatalError("Cant deserialize Plugin");
+    return NULL;
 }
+
+}}

src/Pooling.cpp

@@ -3,7 +3,7 @@
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
                   tkdnnPoolingMode_t pool_mode) : 
@@ -107,4 +107,4 @@ dnnType* Pooling::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Region.cpp

@@ -3,12 +3,13 @@
 #ifdef OPENCV
     #include <opencv2/core/core.hpp>
     #include <opencv2/highgui/highgui.hpp>
+    #include <opencv2/imgproc/imgproc.hpp>
 #endif
 
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Region::Region(Network *net, int classes, int coords, int num) : 
     Layer(net) {
@@ -24,7 +25,7 @@ Region::Region(Network *net, int classes, int coords, int num) :
     output_dim.w = input_dim.w;
     output_dim.l = input_dim.l;
 
-    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
 }
 
 Region::~Region() {
@@ -65,7 +66,7 @@ dnnType* Region::infer(dataDim_t &dim, dnnType* srcData) {
 
 /* Intepret class */
 RegionInterpret::RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, 
-    int classes, int coords, int num, float thresh, const char* fname_weights) {
+    int classes, int coords, int num, float thresh, std::string fname_weights) {
 
     this->input_dim = input_dim;
     this->output_dim = output_dim;
@@ -338,4 +339,4 @@ void RegionInterpret::showImageResult(dnnType *input_h) {
 #endif
 }
 
-}
+}}

src/Reorg.cpp

@@ -3,7 +3,7 @@
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Reorg::Reorg(Network *net, int stride) : Layer(net) {
 
@@ -15,7 +15,7 @@ Reorg::Reorg(Network *net, int stride) : Layer(net) {
     output_dim.w = input_dim.w/stride;
     output_dim.l = input_dim.l;
     
-    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
 }
 
 Reorg::~Reorg() {
@@ -31,4 +31,4 @@ dnnType* Reorg::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Route.cpp

@@ -3,7 +3,7 @@
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) {
 
@@ -42,7 +42,7 @@ dnnType* Route::infer(dataDim_t &dim, dnnType* srcData) {
     int offset = 0;
     for(int i=0; i<layers_n; i++) {
         dnnType *input = layers[i]->dstData;
-        int in_dim = layers[i]->input_dim.tot();
+        int in_dim = layers[i]->output_dim.tot();
         checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
         offset += in_dim;
     }
@@ -53,4 +53,4 @@ dnnType* Route::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Shortcut.cpp

@@ -0,0 +1,37 @@
+#include <iostream>
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tk { namespace dnn {
+
+Shortcut::Shortcut(Network *net, Layer *backLayer) : Layer(net) {
+
+    this->backLayer = backLayer;
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
+
+    if( backLayer->output_dim.c != input_dim.c ||
+        backLayer->output_dim.w != input_dim.w ||
+        backLayer->output_dim.h != input_dim.h   )
+        FatalError("Shortcut dim missmatch");
+}
+
+Shortcut::~Shortcut() {
+
+    checkCuda( cudaFree(dstData) );
+}
+
+dnnType* Shortcut::infer(dataDim_t &dim, dnnType* srcData) {
+
+    dataDim_t bdim = this->backLayer->output_dim;
+
+    checkCuda(cudaMemcpy(dstData, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+    shortcutForward(this->backLayer->dstData, dstData, dim.n, dim.c, dim.h, dim.w, 1, bdim.n, bdim.c, bdim.h, bdim.w, 1);
+
+    //update data dimensions    
+    dim = output_dim;
+
+    return dstData;
+}
+
+}}

src/Softmax.cpp

@@ -3,7 +3,7 @@
 #include "Layer.h"
 #include "kernels.h"
 
-namespace tkDNN {
+namespace tk { namespace dnn {
 
 Softmax::Softmax(Network *net) : Layer(net) {
 
@@ -44,4 +44,4 @@ dnnType* Softmax::infer(dataDim_t &dim, dnnType* srcData) {
     return dstData;
 }
 
-}
+}}

src/Upsample.cpp

@@ -0,0 +1,35 @@
+#include <iostream>
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tk { namespace dnn {
+
+Upsample::Upsample(Network *net, int stride) : Layer(net) {
+
+    this->stride = stride;
+
+    output_dim.n = input_dim.n;
+    output_dim.c = input_dim.c;
+    output_dim.h = input_dim.h*stride;
+    output_dim.w = input_dim.w*stride;
+    output_dim.l = input_dim.l;
+    
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
+}
+
+Upsample::~Upsample() {
+
+    checkCuda( cudaFree(dstData) );
+}
+
+dnnType* Upsample::infer(dataDim_t &dim, dnnType* srcData) {
+
+    fill(dstData, output_dim.tot(), 0.0);
+    upsampleForward(srcData, dstData, input_dim.n, input_dim.c, input_dim.h, input_dim.w, stride, 1, 1);
+
+    dim = output_dim;
+    return dstData;
+}
+
+}}

src/Yolo.cpp

@@ -0,0 +1,251 @@
+#include <iostream>
+
+#ifdef OPENCV
+    #include <opencv2/core/core.hpp>
+    #include <opencv2/highgui/highgui.hpp>
+    #include <opencv2/imgproc/imgproc.hpp>
+#endif
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tk { namespace dnn {
+
+Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights) : 
+    Layer(net) {
+    
+    this->classes = classes;
+    this->num = num;
+
+    // load anchors
+    if(fname_weights != "") {
+        int seek = 0;
+        readBinaryFile(fname_weights, num, &mask_h, &mask_d, seek);
+        seek += num;
+        readBinaryFile(fname_weights, 3*num*2, &bias_h, &bias_d, seek);
+    }
+
+    // init default classes name
+    classesNames.clear();
+    for(int i=0; i<classes; i++) {
+        classesNames.push_back(std::to_string(i));
+    }
+
+    // same
+    output_dim.n = input_dim.n;
+    output_dim.c = input_dim.c;
+    output_dim.h = input_dim.h;
+    output_dim.w = input_dim.w;
+    output_dim.l = input_dim.l;
+
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
+    predictions = nullptr;
+}
+
+Yolo::~Yolo() {
+    checkCuda( cudaFree(dstData) );
+}
+
+int entry_index(int batch, int location, int entry, 
+            int classes, dataDim_t &input_dim, dataDim_t &output_dim) {
+    int n =   location / (input_dim.w*input_dim.h);
+    int loc = location % (input_dim.w*input_dim.h);
+    return batch*output_dim.tot() + n*input_dim.w*input_dim.h*(4+classes+1) +
+           entry*input_dim.w*input_dim.h + loc;
+}
+
+Yolo::box get_yolo_box(float *x, float *biases, int n, int index, int i, int j, int lw, int lh, int w, int h, int stride) {
+    Yolo::box b;
+    b.x = (i + x[index + 0*stride]) / lw;
+    b.y = (j + x[index + 1*stride]) / lh;
+    b.w = exp(x[index + 2*stride]) * biases[2*n]   / w;
+    b.h = exp(x[index + 3*stride]) * biases[2*n+1] / h;
+    return b;
+}
+
+dnnType* Yolo::infer(dataDim_t &dim, dnnType* srcData) {
+
+    checkCuda( cudaMemcpy(dstData, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+
+    for (int b = 0; b < dim.n; ++b){
+        for(int n = 0; n < num; ++n){
+            int index = entry_index(b, n*dim.w*dim.h, 0, classes, input_dim, output_dim);
+            activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
+            
+            index = entry_index(b, n*dim.w*dim.h, 4, classes, input_dim, output_dim);
+            activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*dim.w*dim.h);
+        }
+    }
+
+    dim = output_dim;
+    return dstData;
+}
+
+void correct_yolo_boxes(Yolo::detection *dets, int n, int w, int h, int netw, int neth, int relative)
+{
+    int i;
+    int new_w=0;
+    int new_h=0;
+    if (((float)netw/w) < ((float)neth/h)) {
+        new_w = netw;
+        new_h = (h * netw)/w;
+    } else {
+        new_h = neth;
+        new_w = (w * neth)/h;
+    }
+    for (i = 0; i < n; ++i){
+        Yolo::box b = dets[i].bbox;
+        b.x =  (b.x - (netw - new_w)/2./netw) / ((float)new_w/netw); 
+        b.y =  (b.y - (neth - new_h)/2./neth) / ((float)new_h/neth); 
+        b.w *= (float)netw/new_w;
+        b.h *= (float)neth/new_h;
+        if(!relative){
+            b.x *= w;
+            b.w *= w;
+            b.y *= h;
+            b.h *= h;
+        }
+        dets[i].bbox = b;
+    }
+}
+
+int Yolo::computeDetections(Yolo::detection *dets, int &ndets, int netw, int neth, float thresh) {
+
+    if(predictions == nullptr)
+        predictions = new dnnType[output_dim.tot()];
+    checkCuda( cudaMemcpy(predictions, dstData, output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost));
+
+    int lw = output_dim.w;
+    int lh = output_dim.h;
+
+    if (output_dim.n == 2) {
+        FatalError("BATCH of 2 not supported"); 
+        //avg_flipped_yolo(l);
+    }
+    int i,j,n;
+    int count = ndets;
+    for (i = 0; i < lw*lh; ++i){
+        int row = i / lw;
+        int col = i % lw;
+        for(n = 0; n < num; ++n){
+            int obj_index  = entry_index(0, n*lw*lh + i, 4, classes, input_dim, output_dim);
+            float objectness = predictions[obj_index];
+            if(objectness <= thresh) continue;
+            int box_index  = entry_index(0, n*lw*lh + i, 0, classes, input_dim, output_dim);
+            
+            dets[count].bbox = get_yolo_box(predictions, bias_h, mask_h[n], box_index, col, row, lw, lh, netw, neth, lw*lh);
+            dets[count].objectness = objectness;
+            dets[count].classes = classes;
+            for(j = 0; j < classes; ++j){
+                int class_index = entry_index(0, n*lw*lh + i, 4 + 1 + j, classes, input_dim, output_dim);
+                float prob = objectness*predictions[class_index];
+                dets[count].prob[j] = (prob > thresh) ? prob : 0;
+            }
+            
+            ++count;
+            if(count >= MAX_DETECTIONS)
+                FatalError("reach max boxes");
+        }
+    }
+
+    correct_yolo_boxes(dets + ndets, count, netw, neth, netw, neth, 0);
+    ndets = count;
+    return count;
+}
+
+//////////////////////////////////////////////////////////////////
+float yolo_overlap(float x1, float w1, float x2, float w2)
+{
+    float l1 = x1 - w1/2;
+    float l2 = x2 - w2/2;
+    float left = l1 > l2 ? l1 : l2;
+    float r1 = x1 + w1/2;
+    float r2 = x2 + w2/2;
+    float right = r1 < r2 ? r1 : r2;
+    return right - left;
+}
+
+float yolo_box_intersection(Yolo::box a, Yolo::box b)
+{
+    float w = yolo_overlap(a.x, a.w, b.x, b.w);
+    float h = yolo_overlap(a.y, a.h, b.y, b.h);
+    if(w < 0 || h < 0) return 0;
+    float area = w*h;
+    return area;
+}
+
+float yolo_box_union(Yolo::box a, Yolo::box b)
+{
+    float i = yolo_box_intersection(a, b);
+    float u = a.w*a.h + b.w*b.h - i;
+    return u;
+}
+
+float yolo_box_iou(Yolo::box a, Yolo::box b)
+{
+    return yolo_box_intersection(a, b)/yolo_box_union(a, b);
+}
+
+int yolo_nms_comparator(const void *pa, const void *pb)
+{
+    Yolo::detection a = *(Yolo::detection *)pa;
+    Yolo::detection b = *(Yolo::detection *)pb;
+    float diff = 0;
+    if(b.sort_class >= 0){
+        diff = a.prob[b.sort_class] - b.prob[b.sort_class];
+    } else {
+        diff = a.objectness - b.objectness;
+    }
+    if(diff < 0) return 1;
+    else if(diff > 0) return -1;
+    return 0;
+}
+//////////////////////////////////////////////////////////////////7
+
+Yolo::detection *Yolo::allocateDetections(int nboxes, int classes) {
+    
+    int i;
+    Yolo::detection *dets = (Yolo::detection*) calloc(nboxes, sizeof(Yolo::detection));
+    for(i = 0; i < nboxes; ++i){
+        dets[i].prob = (float*) calloc(classes, sizeof(float));
+    }
+    return dets;
+}
+
+void Yolo::mergeDetections(Yolo::detection *dets, int ndets, int classes) {
+    double nms_thresh = 0.45;
+    int total = ndets;
+
+    int i, j, k;
+    k = total-1;
+    for(i = 0; i <= k; ++i){
+        if(dets[i].objectness == 0){
+            detection swap = dets[i];
+            dets[i] = dets[k];
+            dets[k] = swap;
+            --k;
+            --i;
+        }
+    }
+    total = k+1;
+
+    for(k = 0; k < classes; ++k){
+        for(i = 0; i < total; ++i){
+            dets[i].sort_class = k;
+        }
+        qsort(dets, total, sizeof(detection), yolo_nms_comparator);
+        for(i = 0; i < total; ++i){
+            if(dets[i].prob[k] == 0) continue;
+            box a = dets[i].bbox;
+            for(j = i+1; j < total; ++j){
+                box b = dets[j].bbox;
+                if (yolo_box_iou(a, b) > nms_thresh){
+                    dets[j].prob[k] = 0;
+                }
+            }
+        }
+    }
+
+}
+
+}}

src/Yolo3Detection.cpp

@@ -0,0 +1,152 @@
+#include "Yolo3Detection.h"
+
+namespace tk { namespace dnn {
+
+float _colors[6][3] = { {1,0,1}, {0,0,1},{0,1,1},{0,1,0},{1,1,0},{1,0,0} };
+float get_color(int c, int x, int max)
+{
+    float ratio = ((float)x/max)*5;
+    int i = floor(ratio);
+    int j = ceil(ratio);
+    ratio -= i;
+    float r = (1-ratio) * _colors[i % 6][c % 3] + ratio*_colors[j % 6][c % 3];
+    //printf("%f\n", r);
+    return r;
+}
+
+bool Yolo3Detection::init(std::string tensor_path) {
+
+    //const char *tensor_path = "../data/yolo3/yolo3_berkeley.rt";
+
+    //convert network to tensorRT
+    std::cout<<(tensor_path).c_str()<<"\n";
+    netRT = new tk::dnn::NetworkRT(NULL, (tensor_path).c_str() );
+
+    if(netRT->pluginFactory->n_yolos != 3) {
+        FatalError("this is not yolo3");
+    }
+
+    for(int i=0; i<netRT->pluginFactory->n_yolos; i++) {
+        YoloRT *yRT = netRT->pluginFactory->yolos[i];
+        classes = yRT->classes;
+        num = yRT->num;
+
+        // make a yolo layer for interpret predictions
+        yolo[i] = new tk::dnn::Yolo(nullptr, classes, num, ""); // yolo without input and bias
+        yolo[i]->mask_h = new dnnType[num];
+        yolo[i]->bias_h = new dnnType[num*3*2];
+        memcpy(yolo[i]->mask_h, yRT->mask, sizeof(dnnType)*num);
+        memcpy(yolo[i]->bias_h, yRT->bias, sizeof(dnnType)*num*3*2);
+        yolo[i]->input_dim = yolo[i]->output_dim = tk::dnn::dataDim_t(1, yRT->c, yRT->h, yRT->w);
+        yolo[i]->classesNames = yRT->classesNames;
+    }
+
+    dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
+    
+    checkCuda(cudaMallocHost(&input, sizeof(dnnType)*netRT->input_dim.tot()));
+    checkCuda(cudaMalloc(&input_d, sizeof(dnnType)*netRT->input_dim.tot()));
+
+
+    // class colors precompute    
+    for(int c=0; c<classes; c++) {
+        int offset = c*123457 % classes;
+        float r = get_color(2, offset, classes);
+        float g = get_color(1, offset, classes);
+        float b = get_color(0, offset, classes);
+        colors[c] = cv::Scalar(int(255.0*b), int(255.0*g), int(255.0*r));
+    }
+    return true;
+} 
+
+
+void Yolo3Detection::update(cv::Mat &imageORIG) {
+
+    if(!imageORIG.data) {
+        std::cout<<"YOLO: NO IMAGE DATA\n";
+        return;
+    }     
+    float xRatio =  float(imageORIG.cols) / float(netRT->input_dim.w);
+    float yRatio =  float(imageORIG.rows) / float(netRT->input_dim.h);
+
+    resize(imageORIG, imageORIG, cv::Size(netRT->input_dim.w, netRT->input_dim.h));
+
+    imageORIG.convertTo(imageF, CV_32FC3, 1/255.0); 
+
+    //split channels
+    cv::split(imageF,bgr);//split source
+
+    //write channels
+    for(int i=0; i<netRT->input_dim.c; i++) {
+        int idx = i*imageF.rows*imageF.cols;
+        int ch = netRT->input_dim.c-1 -i;
+        memcpy((void*)&input[idx], (void*)bgr[ch].data, imageF.rows*imageF.cols*sizeof(dnnType));
+    }
+
+
+    //DO INFERENCE
+    dnnType *rt_out[3]; 
+    tk::dnn::dataDim_t dim = netRT->input_dim;
+    checkCuda(cudaMemcpyAsync(input_d, input, dim.tot()*sizeof(dnnType), cudaMemcpyHostToDevice, netRT->stream));
+
+    printCenteredTitle(" TENSORRT inference ", '=', 30); {
+        dim.print();
+        TIMER_START
+        netRT->infer(dim, input_d);
+        TIMER_STOP
+        dim.print();
+    
+        stats.push_back(t_ns);
+    }
+    
+    TIMER_START
+    // compute dets
+    ndets = 0;
+    for(int i=0; i<3; i++) {
+        rt_out[i] = (dnnType*)netRT->buffersRT[i+1];
+        yolo[i]->dstData = rt_out[i];
+        yolo[i]->computeDetections(dets, ndets, netRT->input_dim.w, netRT->input_dim.h, thresh);
+    }
+    tk::dnn::Yolo::mergeDetections(dets, ndets, classes);
+    TIMER_STOP
+
+    // 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;
+        for(int c=0; c<classes; c++) {
+            if(dets[j].prob[c] >= thresh) {
+                obj_class = c;
+                prob = dets[j].prob[c];
+            }
+        }
+
+        if(obj_class >= 0) {
+            //std::cout<<obj_class<<" ("<<prob<<"): "<<x0<<" "<<y0<<" "<<x1<<" "<<y1<<"\n";
+            //cv::rectangle(image, cv::Point(x0, y0), cv::Point(x1, y1), colors[obj_class], 2);
+            
+            // convert to image coords
+            x0 = xRatio*x0;
+            x1 = xRatio*x1;
+            y0 = yRatio*y0;
+            y1 = yRatio*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;
+            detected.push_back(res);
+        }
+    }
+
+}
+
+}}

src/kernels/fill.cu

@@ -0,0 +1,21 @@
+#include "kernels.h"
+
+__global__
+void fill_kernel(dnnType *data, int size, dnnType val) {
+
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+
+    if(i<size) {    
+        data[i] = val;
+    }
+ }
+
+
+void fill(dnnType* data, int size, dnnType val, cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    fill_kernel<<<blocks, threads, 0, stream>>>(data, size, val);
+}
+
+

src/kernels/shortcut.cu

@@ -0,0 +1,47 @@
+#include "kernels.h"
+#include "assert.h"
+
+__global__ void shortcut_kernel(int size, int minw, int minh, int minc, int stride, int sample, int batch, 
+                                int w1, int h1, int c1, dnnType *add, 
+                                int w2, int h2, int c2, float s1, float s2, dnnType *out)
+{
+    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if (id >= size) return;
+    int i = id % minw;
+    id /= minw;
+    int j = id % minh;
+    id /= minh;
+    int k = id % minc;
+    id /= minc;
+    int b = id % batch;
+
+    int out_index = i*sample + w2*(j*sample + h2*(k + c2*b));
+    int add_index = i*stride + w1*(j*stride + h1*(k + c1*b));
+    out[out_index] = s1*out[out_index] + s2*add[add_index];
+    //out[out_index] += add[add_index];
+}
+
+void shortcutForward(dnnType* srcData, dnnType* dstData, int n1, int c1, int h1, int w1, int s1,
+                                                         int n2, int c2, int h2, int w2, int s2, 
+                     cudaStream_t stream)
+{
+    assert(n1 == n2);
+    int batch = n1;
+
+    int minw = (w1 < w2) ? w1 : w2;
+    int minh = (h1 < h2) ? h1 : h2;
+    int minc = (c1 < c2) ? c1 : c2;
+
+    int stride = w1/w2;
+    int sample = w2/w1;
+    assert(stride == h1/h2);
+    assert(sample == h2/h1);
+    if(stride < 1) stride = 1;
+    if(sample < 1) sample = 1;
+
+    int size = batch * minw * minh * minc;
+    int blocks = (size+255)/256;
+    int threads = 256;
+    shortcut_kernel<<<blocks, threads, 0, stream>>>(size, minw, minh, minc, stride, sample, batch, 
+        w1, h1, c1, srcData, w2, h2, c2, s1, s2, dstData);
+}

src/kernels/upsample.cu

@@ -0,0 +1,35 @@
+#include "kernels.h"
+
+__global__ void upsample_kernel(size_t N, dnnType *x, int w, int h, int c, int batch, int stride, int forward, float scale, dnnType *out)
+{
+    size_t i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if(i >= N) return;
+    int out_index = i;
+    int out_w = i%(w*stride);
+    i = i/(w*stride);
+    int out_h = i%(h*stride);
+    i = i/(h*stride);
+    int out_c = i%c;
+    i = i/c;
+    int b = i%batch;
+
+    int in_w = out_w / stride;
+    int in_h = out_h / stride;
+    int in_c = out_c;
+
+    int in_index = b*w*h*c + in_c*w*h + in_h*w + in_w;
+
+
+    if(forward) out[out_index] += scale * x[in_index];
+    else atomicAdd(x+in_index, scale * out[out_index]);
+}
+
+void upsampleForward(dnnType* srcData, dnnType* dstData, 
+                     int n, int c, int h, int w, int s, int forward, float scale,                                   
+                     cudaStream_t stream) {
+
+    int size = w*h*c*n*s*s;
+    int blocks = (size+255)/256;
+    int threads = 256;
+    upsample_kernel<<<blocks, threads, 0, stream>>>(size, srcData, w, h, c, n, s, forward, scale, dstData);
+}

src/utils.cpp

@@ -21,7 +21,7 @@ bool fileExist(const char *fname) {
 }
 
 
-void readBinaryFile(const char* fname, int size, dnnType** data_h, dnnType** data_d, int seek)
+void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** data_d, int seek)
 {
     std::ifstream dataFile (fname, std::ios::in | std::ios::binary);
     std::stringstream error_s;
@@ -70,7 +70,7 @@ void printDeviceVector(int size, dnnType* vec_d, bool device)
 int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) {
 
     dnnType *data_h, *correct_h;
-    const float eps = 0.001f;
+    const float eps = 0.02f;
 
     if(device) {
         data_h = new dnnType[size];

tests/mnist/test_mnist.cpp

@@ -11,18 +11,18 @@ const char *output_bin   = "../tests/mnist/output.bin";
 int main() {
 
     // Network layout
-    tkDNN::dataDim_t dim(1, 1, 28, 28, 1);
-    tkDNN::Network net(dim);
-    tkDNN::Conv2d     l0(&net, 20, 5, 5, 1, 1, 0, 0, c0_bin);
-    tkDNN::Pooling    l1(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
-    tkDNN::Conv2d     l2(&net, 50, 5, 5, 1, 1, 0, 0, c1_bin);
-    tkDNN::Pooling    l3(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
-    tkDNN::Dense      l4(&net, 500, d2_bin);
-    tkDNN::Activation l5(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Dense      l6(&net, 10, d3_bin);
-    tkDNN::Softmax    l7(&net);
+    tk::dnn::dataDim_t dim(1, 1, 28, 28, 1);
+    tk::dnn::Network net(dim);
+    tk::dnn::Conv2d     l0(&net, 20, 5, 5, 1, 1, 0, 0, c0_bin);
+    tk::dnn::Pooling    l1(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+    tk::dnn::Conv2d     l2(&net, 50, 5, 5, 1, 1, 0, 0, c1_bin);
+    tk::dnn::Pooling    l3(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+    tk::dnn::Dense      l4(&net, 500, d2_bin);
+    tk::dnn::Activation l5(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Dense      l6(&net, 10, d3_bin);
+    tk::dnn::Softmax    l7(&net);
 
-    tkDNN::NetworkRT netRT(&net, "mnist.rt");
+    tk::dnn::NetworkRT netRT(&net, "mnist.rt");
 
     // Load input
     dnnType *data;
@@ -43,7 +43,7 @@ int main() {
     //std::cout<<"\n======= CUDNN RESULT =======\n";
     //printDeviceVector(10, out_data);
  
-    tkDNN::dataDim_t dim2(1, 1, 28, 28, 1);
+    tk::dnn::dataDim_t dim2(1, 1, 28, 28, 1);
 
     std::cout<<"TENSORRT inference:\n"; {
         dim2.print();

tests/mnist/test_mnistRT.cpp

@@ -27,16 +27,16 @@ int main() {
 
     std::cout<<"\n==== CUDNN ====\n";
     // Network layout
-    tkDNN::dataDim_t dim(1, 1, 28, 28, 1);
-	tkDNN::Network net(dim);
-    tkDNN::Conv2d     l0(&net, 20, 5, 5, 1, 1, 0, 0, c0_bin);
-    tkDNN::Pooling    l1(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
-    tkDNN::Conv2d     l2(&net, 50, 5, 5, 1, 1, 0, 0, c1_bin);
-    tkDNN::Pooling    l3(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
-    tkDNN::Dense      l4(&net, 500, d2_bin);
-    tkDNN::Activation l5(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Dense      l6(&net, 10, d3_bin);
-    tkDNN::Softmax    l7(&net);
+    tk::dnn::dataDim_t dim(1, 1, 28, 28, 1);
+	tk::dnn::Network net(dim);
+    tk::dnn::Conv2d     l0(&net, 20, 5, 5, 1, 1, 0, 0, c0_bin);
+    tk::dnn::Pooling    l1(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+    tk::dnn::Conv2d     l2(&net, 50, 5, 5, 1, 1, 0, 0, c1_bin);
+    tk::dnn::Pooling    l3(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+    tk::dnn::Dense      l4(&net, 500, d2_bin);
+    tk::dnn::Activation l5(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Dense      l6(&net, 10, d3_bin);
+    tk::dnn::Softmax    l7(&net);
  
     // Load input
     dnnType *data;
@@ -71,7 +71,7 @@ int main() {
 	auto input = network->addInput("data", dt, DimsCHW{ 1, 28, 28});
 	assert(input != nullptr);
 
-    tkDNN::Conv2d *c0 = &l0; 
+    tk::dnn::Conv2d *c0 = &l0; 
     Weights w { dt, c0->data_h, c0->inputs*c0->outputs*c0->kernelH*c0->kernelW};
     Weights b { dt, c0->bias_h, c0->outputs};
 	// Add a convolution layer with 20 outputs and a 5x5 filter.
@@ -84,7 +84,7 @@ int main() {
 	assert(pool1 != nullptr);
 	pool1->setStride(DimsHW{2, 2});
 
-    tkDNN::Conv2d *c1 = &l2; 
+    tk::dnn::Conv2d *c1 = &l2; 
     Weights w1 { dt, c1->data_h, c1->inputs*c1->outputs*c1->kernelH*c1->kernelW};
     Weights b1 { dt, c1->bias_h, c1->outputs};
 	// Add a second convolution layer with 50 outputs and a 5x5 filter.
@@ -97,7 +97,7 @@ int main() {
 	assert(pool2 != nullptr);
 	pool2->setStride(DimsHW{2, 2});
 
-    tkDNN::Dense *d2 = &l4; 
+    tk::dnn::Dense *d2 = &l4; 
     Weights w2 { dt, d2->data_h, d2->inputs*d2->outputs};
     Weights b2 { dt, d2->bias_h, d2->outputs};
 	// Add a fully connected layer with 500 outputs.
@@ -108,7 +108,7 @@ int main() {
 	auto relu1 = network->addActivation(*ip1->getOutput(0), ActivationType::kRELU);
 	assert(relu1 != nullptr);
 
-    tkDNN::Dense *d3 = &l6; 
+    tk::dnn::Dense *d3 = &l6; 
     Weights w3 { dt, d3->data_h, d3->inputs*d3->outputs};
     Weights b3 { dt, d3->bias_h, d3->outputs};
 	// Add a second fully connected layer with 20 outputs.

tests/simple/test_simple.cpp

@@ -10,15 +10,15 @@ const char *output_bin  = "../tests/simple/output.bin";
 int main() {
 
     // Network layout
-    tkDNN::dataDim_t dim(1, 1, 10, 10, 1);
-    tkDNN::Network net(dim);
-    tkDNN::Conv2d     l0(&net, 2, 4, 4, 2, 2, 0, 0, c0_bin);
-    tkDNN::Activation l1(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     l2(&net, 4, 2, 2, 1, 1, 0, 0, c1_bin);
-    tkDNN::Activation l3(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Flatten    l4(&net);
-    tkDNN::Dense      l5(&net, 4, d2_bin);
-    tkDNN::Activation l6(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::dataDim_t dim(1, 1, 10, 10, 1);
+    tk::dnn::Network net(dim);
+    tk::dnn::Conv2d     l0(&net, 2, 4, 4, 2, 2, 0, 0, c0_bin);
+    tk::dnn::Activation l1(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     l2(&net, 4, 2, 2, 1, 1, 0, 0, c1_bin);
+    tk::dnn::Activation l3(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Flatten    l4(&net);
+    tk::dnn::Dense      l5(&net, 4, d2_bin);
+    tk::dnn::Activation l6(&net, CUDNN_ACTIVATION_RELU);
 
     // Load input
     dnnType *data;

tests/test_rtinference/rtinference.cpp

@@ -11,7 +11,7 @@ int main(int argc, char *argv[]) {
     srand (0); 
 
     //convert network to tensorRT
-    tkDNN::NetworkRT netRT(NULL, argv[1]);
+    tk::dnn::NetworkRT netRT(NULL, argv[1]);
 
     dnnType *input = new float[netRT.input_dim.tot()];
     dnnType *output = new float[netRT.input_dim.tot()];

tests/yolo/yolo.cpp

@@ -31,75 +31,75 @@ const char *output_bin = "../tests/yolo/layers/output.bin";
 int main() {
 
     // Network layout
-    tkDNN::dataDim_t dim(1, 3, 608, 608, 1);
-    tkDNN::Network net(dim);
+    tk::dnn::dataDim_t dim(1, 3, 608, 608, 1);
+    tk::dnn::Network net(dim);
 
-    tkDNN::Conv2d     c0 (&net, 32, 3, 3, 1, 1, 1, 1,   c0_bin, true);
-    tkDNN::Activation a0 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p1 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c0 (&net, 32, 3, 3, 1, 1, 1, 1,   c0_bin, true);
+    tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p1 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c2 (&net, 64, 3, 3, 1, 1, 1, 1,   c2_bin, true);
-    tkDNN::Activation a2 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p3 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c2 (&net, 64, 3, 3, 1, 1, 1, 1,   c2_bin, true);
+    tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p3 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c4 (&net, 128, 3, 3, 1, 1, 1, 1,  c4_bin, true);
-    tkDNN::Activation a4 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c5 (&net, 64, 1, 1, 1, 1, 0, 0,   c5_bin, true);
-    tkDNN::Activation a5 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
-    tkDNN::Activation a6 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p7 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c4 (&net, 128, 3, 3, 1, 1, 1, 1,  c4_bin, true);
+    tk::dnn::Activation a4 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c5 (&net, 64, 1, 1, 1, 1, 0, 0,   c5_bin, true);
+    tk::dnn::Activation a5 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
+    tk::dnn::Activation a6 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p7 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c8 (&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
-    tkDNN::Activation a8 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c9 (&net, 128, 1, 1, 1, 1, 0, 0,  c9_bin, true);
-    tkDNN::Activation a9 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c10(&net, 256, 3, 3, 1, 1, 1, 1,  c10_bin, true);
-    tkDNN::Activation a10(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p11(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c8 (&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
+    tk::dnn::Activation a8 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c9 (&net, 128, 1, 1, 1, 1, 0, 0,  c9_bin, true);
+    tk::dnn::Activation a9 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c10(&net, 256, 3, 3, 1, 1, 1, 1,  c10_bin, true);
+    tk::dnn::Activation a10(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p11(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1,  c12_bin, true);
-    tkDNN::Activation a12(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c13(&net, 256, 1, 1, 1, 1, 0, 0,  c13_bin, true);
-    tkDNN::Activation a13(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c14(&net, 512, 3, 3, 1, 1, 1, 1,  c14_bin, true);
-    tkDNN::Activation a14(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c15(&net, 256, 1, 1, 1, 1, 0, 0,  c15_bin, true);
-    tkDNN::Activation a15(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c16(&net, 512, 3, 3, 1, 1, 1, 1,  c16_bin, true); 
-    tkDNN::Activation a16(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p17(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1,  c12_bin, true);
+    tk::dnn::Activation a12(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c13(&net, 256, 1, 1, 1, 1, 0, 0,  c13_bin, true);
+    tk::dnn::Activation a13(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c14(&net, 512, 3, 3, 1, 1, 1, 1,  c14_bin, true);
+    tk::dnn::Activation a14(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c15(&net, 256, 1, 1, 1, 1, 0, 0,  c15_bin, true);
+    tk::dnn::Activation a15(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c16(&net, 512, 3, 3, 1, 1, 1, 1,  c16_bin, true); 
+    tk::dnn::Activation a16(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p17(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
-    tkDNN::Activation a18(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c19(&net, 512, 1, 1, 1, 1, 0, 0,  c19_bin, true);
-    tkDNN::Activation a19(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
-    tkDNN::Activation a20(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c21(&net, 512, 1, 1, 1, 1, 0, 0,  c21_bin, true);
-    tkDNN::Activation a21(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
-    tkDNN::Activation a22(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
-    tkDNN::Activation a23(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
-    tkDNN::Activation a24(&net, tkDNN::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
+    tk::dnn::Activation a18(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c19(&net, 512, 1, 1, 1, 1, 0, 0,  c19_bin, true);
+    tk::dnn::Activation a19(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
+    tk::dnn::Activation a20(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c21(&net, 512, 1, 1, 1, 1, 0, 0,  c21_bin, true);
+    tk::dnn::Activation a21(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
+    tk::dnn::Activation a22(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
+    tk::dnn::Activation a23(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
+    tk::dnn::Activation a24(&net, tk::dnn::ACTIVATION_LEAKY);
 
-    tkDNN::Layer *m25_layers[1] = { &a16 };
-    tkDNN::Route      m25(&net, m25_layers, 1);
-    tkDNN::Conv2d     c26(&net, 64, 1, 1, 1, 1, 0, 0,   c26_bin, true);
-    tkDNN::Activation a26(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Reorg      r27(&net, 2);
+    tk::dnn::Layer *m25_layers[1] = { &a16 };
+    tk::dnn::Route      m25(&net, m25_layers, 1);
+    tk::dnn::Conv2d     c26(&net, 64, 1, 1, 1, 1, 0, 0,   c26_bin, true);
+    tk::dnn::Activation a26(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Reorg      r27(&net, 2);
 
-    tkDNN::Layer *m28_layers[2] = { &r27, &a24 };
-    tkDNN::Route      m28(&net, m28_layers, 2);
+    tk::dnn::Layer *m28_layers[2] = { &r27, &a24 };
+    tk::dnn::Route      m28(&net, m28_layers, 2);
 
-    tkDNN::Conv2d     c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
-    tkDNN::Activation a29(&net, tkDNN::ACTIVATION_LEAKY);   
-    tkDNN::Conv2d     c30(&net, 425, 1, 1, 1, 1, 0, 0,  c30_bin, false);
-    tkDNN::Region     g31(&net, 80, 4, 5);
+    tk::dnn::Conv2d     c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
+    tk::dnn::Activation a29(&net, tk::dnn::ACTIVATION_LEAKY);   
+    tk::dnn::Conv2d     c30(&net, 425, 1, 1, 1, 1, 0, 0,  c30_bin, false);
+    tk::dnn::Region     g31(&net, 80, 4, 5);
     
-    tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.6f, g31_bin);
+    tk::dnn::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.6f, g31_bin);
 
     // Load input
     dnnType *data;
@@ -110,11 +110,11 @@ int main() {
     net.print();
 
     //convert network to tensorRT
-    tkDNN::NetworkRT netRT(&net, "yolo.rt");
+    tk::dnn::NetworkRT netRT(&net, "yolo.rt");
 
     dnnType *out_data, *out_data2; // cudnn output, tensorRT output
 
-    tkDNN::dataDim_t dim1 = dim; //input dim
+    tk::dnn::dataDim_t dim1 = dim; //input dim
     printCenteredTitle(" CUDNN inference ", '=', 30); {
         dim1.print();
         TIMER_START
@@ -123,7 +123,7 @@ int main() {
         dim1.print();   
     }
     
-    tkDNN::dataDim_t dim2 = dim;
+    tk::dnn::dataDim_t dim2 = dim;
     printCenteredTitle(" TENSORRT inference ", '=', 30); {
         dim2.print();
         TIMER_START

tests/yolo3_berkeley/yolo3_berkeley.cfg

@@ -0,0 +1,785 @@
+[net]
+# Testing
+batch=1
+subdivisions=1
+# Training
+#batch=32
+#subdivisions=8
+width=544
+height=320
+channels=3
+momentum=0.9
+decay=0.0005
+angle=0
+saturation = 1.5
+exposure = 1.5
+hue=.1
+
+learning_rate=0.001
+burn_in=1000
+max_batches = 50200
+policy=steps
+steps=40000,45000
+scales=.1,.1
+
+
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=2
+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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+######################
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=45
+activation=linear
+
+[yolo]
+mask = 6,7,8
+anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648
+classes=10
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=0
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 61
+
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=45
+activation=linear
+
+[yolo]
+mask = 3,4,5
+anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648
+classes=10
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=0
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 36
+
+
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=45
+activation=linear
+
+[yolo]
+mask = 0,1,2
+anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648
+classes=10
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=0
+

tests/yolo3_berkeley/yolo3_berkeley.cpp

@@ -0,0 +1,92 @@
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+
+int main() {
+
+    // Network layout
+    tk::dnn::dataDim_t dim(1, 3, 320, 544, 1);
+    tk::dnn::Network net(dim);
+
+    // create yolo3 model
+    std::string bin_path  = "../tests/yolo3_berkeley";
+    int classes = 10;
+    tk::dnn::Yolo *yolo [3];
+    #include "models/Yolo3.h"
+
+    // fill classes names
+    for(int i=0; i<3; i++) {
+        yolo[i]->classesNames = {"person", "car", "truck", "bus", "motor", "bike", "rider", "traffic light", "traffic sign", "train"};
+    }
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+    
+    //print network model
+    net.print();
+
+    //convert network to tensorRT
+    tk::dnn::NetworkRT netRT(&net, "yolo3_berkeley.rt");
+
+    // the network have 3 outputs
+    tk::dnn::dataDim_t out_dim[3];
+    for(int i=0; i<3; i++) out_dim[i] = yolo[i]->output_dim; 
+    dnnType *cudnn_out[3], *rt_out[3]; 
+
+    tk::dnn::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+    for(int i=0; i<3; i++) cudnn_out[i] = yolo[i]->dstData;
+    
+    printCenteredTitle(" compute detections ", '=', 30);
+    TIMER_START
+    int ndets = 0;
+    tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
+    for(int i=0; i<3; i++) yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5);
+    tk::dnn::Yolo::mergeDetections(dets, ndets, classes);
+
+    for(int j=0; j<ndets; j++) {
+        tk::dnn::Yolo::box b = dets[j].bbox;
+        int x0   = (b.x-b.w/2.);
+        int x1   = (b.x+b.w/2.);
+        int y0   = (b.y-b.h/2.);
+        int y1   = (b.y+b.h/2.);
+
+        int cl = 0;
+        for(int c = 0; c < classes; ++c){
+            float prob = dets[j].prob[c];
+            if(prob > 0)
+                cl = c;
+        }
+        std::cout<<cl<<": "<<x0<<" "<<y0<<" "<<x1<<" "<<y1<<"\n";
+    }
+    TIMER_STOP
+    
+    tk::dnn::dataDim_t dim2 = dim;
+    printCenteredTitle(" TENSORRT inference ", '=', 30); {
+        dim2.print();
+        TIMER_START
+        netRT.infer(dim2, data);
+        TIMER_STOP
+        dim2.print();
+    }
+    for(int i=0; i<3; i++) rt_out[i] = (dnnType*)netRT.buffersRT[i+1];
+
+    for(int i=0; i<3; i++) {
+        printCenteredTitle((std::string(" YOLO ") + std::to_string(i) + " CHECK RESULTS ").c_str(), '=', 30);
+        dnnType *out, *out_h;
+        int odim = out_dim[i].tot();
+        readBinaryFile(output_bins[i], odim, &out_h, &out);
+        std::cout<<"CUDNN vs correct"; checkResult(odim, cudnn_out[i], out);
+        std::cout<<"TRT   vs correct"; checkResult(odim, rt_out[i], out);
+        std::cout<<"CUDNN vs TRT    "; checkResult(odim, cudnn_out[i], rt_out[i]);
+    }
+    return 0;
+}

tests/yolo3_coco4/yolo3_coco4.cpp

@@ -0,0 +1,87 @@
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+
+int main() {
+
+    // Network layout
+    tk::dnn::dataDim_t dim(1, 3, 416, 416, 1);
+    tk::dnn::Network net(dim);
+
+    // create yolo3 model
+    std::string bin_path  = "../tests/yolo3_coco4";
+    int classes = 4;
+    tk::dnn::Yolo *yolo [3];
+    #include "models/Yolo3.h"
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+    
+    //print network model
+    net.print();
+
+    //convert network to tensorRT
+    tk::dnn::NetworkRT netRT(&net, "yolo3_coco4.rt");
+
+    // the network have 3 outputs
+    tk::dnn::dataDim_t out_dim[3];
+    for(int i=0; i<3; i++) out_dim[i] = yolo[i]->output_dim; 
+    dnnType *cudnn_out[3], *rt_out[3]; 
+
+    tk::dnn::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+    for(int i=0; i<3; i++) cudnn_out[i] = yolo[i]->dstData;
+    
+    printCenteredTitle(" compute detections ", '=', 30);
+    TIMER_START
+    int ndets = 0;
+    tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
+    for(int i=0; i<3; i++) yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5);
+    tk::dnn::Yolo::mergeDetections(dets, ndets, classes);
+
+    for(int j=0; j<ndets; j++) {
+        tk::dnn::Yolo::box b = dets[j].bbox;
+        int x0   = (b.x-b.w/2.);
+        int x1   = (b.x+b.w/2.);
+        int y0   = (b.y-b.h/2.);
+        int y1   = (b.y+b.h/2.);
+
+        int cl = 0;
+        for(int c = 0; c < classes; ++c){
+            float prob = dets[j].prob[c];
+            if(prob > 0)
+                cl = c;
+        }
+        std::cout<<cl<<": "<<x0<<" "<<y0<<" "<<x1<<" "<<y1<<"\n";
+    }
+    TIMER_STOP
+    
+    tk::dnn::dataDim_t dim2 = dim;
+    printCenteredTitle(" TENSORRT inference ", '=', 30); {
+        dim2.print();
+        TIMER_START
+        netRT.infer(dim2, data);
+        TIMER_STOP
+        dim2.print();
+    }
+    for(int i=0; i<3; i++) rt_out[i] = (dnnType*)netRT.buffersRT[i+1];
+
+    for(int i=0; i<3; i++) {
+        printCenteredTitle((std::string(" YOLO ") + std::to_string(i) + " CHECK RESULTS ").c_str(), '=', 30);
+        dnnType *out, *out_h;
+        int odim = out_dim[i].tot();
+        readBinaryFile(output_bins[i], odim, &out_h, &out);
+        std::cout<<"CUDNN vs correct"; checkResult(odim, cudnn_out[i], out);
+        std::cout<<"TRT   vs correct"; checkResult(odim, rt_out[i], out);
+        std::cout<<"CUDNN vs TRT    "; checkResult(odim, cudnn_out[i], rt_out[i]);
+    }
+    return 0;
+}

tests/yolo3_coco4/yolov3-coco4.cfg

@@ -0,0 +1,785 @@
+[net]
+# Testing
+batch=1
+subdivisions=1
+# Training
+#batch=32
+#subdivisions=8
+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.001
+burn_in=1000
+max_batches = 50200
+policy=steps
+steps=40000,45000
+scales=.1,.1
+
+
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=2
+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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+######################
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=27
+activation=linear
+
+[yolo]
+mask = 6,7,8
+anchors = 10,13,  16,30,  33,23,  30,61,  62,45,  59,119,  116,90,  156,198,  373,326
+classes=4
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=1
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 61
+
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=27
+activation=linear
+
+[yolo]
+mask = 3,4,5
+anchors = 10,13,  16,30,  33,23,  30,61,  62,45,  59,119,  116,90,  156,198,  373,326
+classes=4
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=1
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 36
+
+
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=27
+activation=linear
+
+[yolo]
+mask = 0,1,2
+anchors = 10,13,  16,30,  33,23,  30,61,  62,45,  59,119,  116,90,  156,198,  373,326
+classes=4
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=1
+

tests/yolo3_flir/yolo3_flir.cfg

@@ -0,0 +1,785 @@
+[net]
+# Testing
+#batch=1
+#subdivisions=1
+# Training
+batch=32
+subdivisions=8
+width=544
+height=320
+channels=1
+momentum=0.9
+decay=0.0005
+angle=0
+saturation = 1.5
+exposure = 1.5
+hue=.1
+
+learning_rate=0.001
+burn_in=1000
+max_batches = 20000
+policy=steps
+steps=8000,9000
+scales=.1,.1
+
+
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=2
+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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+[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
+
+[shortcut]
+from=-3
+activation=linear
+
+# Downsample
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=2
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[shortcut]
+from=-3
+activation=linear
+
+######################
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=24
+activation=linear
+
+[yolo]
+mask = 6,7,8
+anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648
+classes=3
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=0
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 61
+
+
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=512
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=24
+activation=linear
+
+[yolo]
+mask = 3,4,5
+anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648
+classes=3
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=0
+
+[route]
+layers = -4
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[upsample]
+stride=2
+
+[route]
+layers = -1, 36
+
+
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=256
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=24
+activation=linear
+
+[yolo]
+mask = 0,1,2
+anchors = 8.2087,8.5515, 18.4134,20.3391, 40.2194,29.2990, 31.6137,69.2240, 69.8497,48.3838, 108.8817,76.6316, 96.5753,145.5743, 165.9182,117.4493, 215.7497,198.4648
+classes=3
+num=9
+jitter=.3
+ignore_thresh = .5
+truth_thresh = 1
+random=0
+

tests/yolo3_flir/yolo3_flir.cpp

@@ -0,0 +1,93 @@
+#include<iostream>
+#include<vector>
+#include "tkdnn.h"
+
+
+int main() {
+
+    // Network layout
+    tk::dnn::dataDim_t dim(1, 1, 320, 544, 1);
+    tk::dnn::Network net(dim);
+
+    // create yolo3 model
+    std::string bin_path  = "../tests/yolo3_flir";
+    int classes = 3;
+    tk::dnn::Yolo *yolo [3];
+    #include "models/Yolo3.h"
+
+    // fill classes names
+    for(int i=0; i<3; i++) {
+        yolo[i]->classesNames = {"person", "bike", "car"};
+    }
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+    
+    //print network model
+    net.print();
+
+    //convert network to tensorRT
+    tk::dnn::NetworkRT netRT(&net, "yolo3_flir.rt");
+
+    // the network have 3 outputs
+    tk::dnn::dataDim_t out_dim[3];
+    for(int i=0; i<3; i++) out_dim[i] = yolo[i]->output_dim; 
+    dnnType *cudnn_out[3], *rt_out[3]; 
+
+    tk::dnn::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+    for(int i=0; i<3; i++) cudnn_out[i] = yolo[i]->dstData;
+    
+    printCenteredTitle(" compute detections ", '=', 30);
+    TIMER_START
+    int ndets = 0;
+    tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes);
+    for(int i=0; i<3; i++) yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5);
+    tk::dnn::Yolo::mergeDetections(dets, ndets, classes);
+
+    for(int j=0; j<ndets; j++) {
+        tk::dnn::Yolo::box b = dets[j].bbox;
+        int x0   = (b.x-b.w/2.);
+        int x1   = (b.x+b.w/2.);
+        int y0   = (b.y-b.h/2.);
+        int y1   = (b.y+b.h/2.);
+
+        int cl = 0;
+        for(int c = 0; c < classes; ++c){
+            float prob = dets[j].prob[c];
+            if(prob > 0)
+                cl = c;
+        }
+        std::cout<<cl<<": "<<x0<<" "<<y0<<" "<<x1<<" "<<y1<<"\n";
+    }
+    TIMER_STOP
+    
+    tk::dnn::dataDim_t dim2 = dim;
+    printCenteredTitle(" TENSORRT inference ", '=', 30); {
+        dim2.print();
+        TIMER_START
+        netRT.infer(dim2, data);
+        TIMER_STOP
+        dim2.print();
+    }
+    for(int i=0; i<3; i++) rt_out[i] = (dnnType*)netRT.buffersRT[i+1];
+
+    for(int i=0; i<3; i++) {
+        printCenteredTitle((std::string(" YOLO ") + std::to_string(i) + " CHECK RESULTS ").c_str(), '=', 30);
+        dnnType *out, *out_h;
+        int odim = out_dim[i].tot();
+        readBinaryFile(output_bins[i], odim, &out_h, &out);
+        std::cout<<"CUDNN vs correct"; checkResult(odim, cudnn_out[i], out);
+        std::cout<<"TRT   vs correct"; checkResult(odim, rt_out[i], out);
+        std::cout<<"CUDNN vs TRT    "; checkResult(odim, cudnn_out[i], rt_out[i]);
+    }
+    return 0;
+}

tests/yolo_224/yolo_224.cpp

@@ -31,75 +31,75 @@ const char *output_bin = "../tests/yolo_224/layers/output.bin";
 int main() {
 
     // Network layout
-    tkDNN::dataDim_t dim(1, 3, 224, 224, 1);
-    tkDNN::Network net(dim);
+    tk::dnn::dataDim_t dim(1, 3, 224, 224, 1);
+    tk::dnn::Network net(dim);
 
-    tkDNN::Conv2d     c0 (&net, 32, 3, 3, 1, 1, 1, 1,   c0_bin, true);
-    tkDNN::Activation a0 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p1 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c0 (&net, 32, 3, 3, 1, 1, 1, 1,   c0_bin, true);
+    tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p1 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c2 (&net, 64, 3, 3, 1, 1, 1, 1,   c2_bin, true);
-    tkDNN::Activation a2 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p3 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c2 (&net, 64, 3, 3, 1, 1, 1, 1,   c2_bin, true);
+    tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p3 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c4 (&net, 128, 3, 3, 1, 1, 1, 1,  c4_bin, true);
-    tkDNN::Activation a4 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c5 (&net, 64, 1, 1, 1, 1, 0, 0,   c5_bin, true);
-    tkDNN::Activation a5 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
-    tkDNN::Activation a6 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p7 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c4 (&net, 128, 3, 3, 1, 1, 1, 1,  c4_bin, true);
+    tk::dnn::Activation a4 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c5 (&net, 64, 1, 1, 1, 1, 0, 0,   c5_bin, true);
+    tk::dnn::Activation a5 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
+    tk::dnn::Activation a6 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p7 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c8 (&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
-    tkDNN::Activation a8 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c9 (&net, 128, 1, 1, 1, 1, 0, 0,  c9_bin, true);
-    tkDNN::Activation a9 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c10(&net, 256, 3, 3, 1, 1, 1, 1,  c10_bin, true);
-    tkDNN::Activation a10(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p11(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c8 (&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
+    tk::dnn::Activation a8 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c9 (&net, 128, 1, 1, 1, 1, 0, 0,  c9_bin, true);
+    tk::dnn::Activation a9 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c10(&net, 256, 3, 3, 1, 1, 1, 1,  c10_bin, true);
+    tk::dnn::Activation a10(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p11(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1,  c12_bin, true);
-    tkDNN::Activation a12(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c13(&net, 256, 1, 1, 1, 1, 0, 0,  c13_bin, true);
-    tkDNN::Activation a13(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c14(&net, 512, 3, 3, 1, 1, 1, 1,  c14_bin, true);
-    tkDNN::Activation a14(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c15(&net, 256, 1, 1, 1, 1, 0, 0,  c15_bin, true);
-    tkDNN::Activation a15(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c16(&net, 512, 3, 3, 1, 1, 1, 1,  c16_bin, true); 
-    tkDNN::Activation a16(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p17(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1,  c12_bin, true);
+    tk::dnn::Activation a12(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c13(&net, 256, 1, 1, 1, 1, 0, 0,  c13_bin, true);
+    tk::dnn::Activation a13(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c14(&net, 512, 3, 3, 1, 1, 1, 1,  c14_bin, true);
+    tk::dnn::Activation a14(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c15(&net, 256, 1, 1, 1, 1, 0, 0,  c15_bin, true);
+    tk::dnn::Activation a15(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c16(&net, 512, 3, 3, 1, 1, 1, 1,  c16_bin, true); 
+    tk::dnn::Activation a16(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p17(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
-    tkDNN::Activation a18(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c19(&net, 512, 1, 1, 1, 1, 0, 0,  c19_bin, true);
-    tkDNN::Activation a19(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
-    tkDNN::Activation a20(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c21(&net, 512, 1, 1, 1, 1, 0, 0,  c21_bin, true);
-    tkDNN::Activation a21(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
-    tkDNN::Activation a22(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
-    tkDNN::Activation a23(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
-    tkDNN::Activation a24(&net, tkDNN::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
+    tk::dnn::Activation a18(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c19(&net, 512, 1, 1, 1, 1, 0, 0,  c19_bin, true);
+    tk::dnn::Activation a19(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
+    tk::dnn::Activation a20(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c21(&net, 512, 1, 1, 1, 1, 0, 0,  c21_bin, true);
+    tk::dnn::Activation a21(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
+    tk::dnn::Activation a22(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
+    tk::dnn::Activation a23(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
+    tk::dnn::Activation a24(&net, tk::dnn::ACTIVATION_LEAKY);
 
-    tkDNN::Layer *m25_layers[1] = { &a16 };
-    tkDNN::Route      m25(&net, m25_layers, 1);
-    tkDNN::Conv2d     c26(&net, 64, 1, 1, 1, 1, 0, 0,   c26_bin, true);
-    tkDNN::Activation a26(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Reorg      r27(&net, 2);
+    tk::dnn::Layer *m25_layers[1] = { &a16 };
+    tk::dnn::Route      m25(&net, m25_layers, 1);
+    tk::dnn::Conv2d     c26(&net, 64, 1, 1, 1, 1, 0, 0,   c26_bin, true);
+    tk::dnn::Activation a26(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Reorg      r27(&net, 2);
 
-    tkDNN::Layer *m28_layers[2] = { &r27, &a24 };
-    tkDNN::Route      m28(&net, m28_layers, 2);
+    tk::dnn::Layer *m28_layers[2] = { &r27, &a24 };
+    tk::dnn::Route      m28(&net, m28_layers, 2);
 
-    tkDNN::Conv2d     c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
-    tkDNN::Activation a29(&net, tkDNN::ACTIVATION_LEAKY);   
-    tkDNN::Conv2d     c30(&net, 425, 1, 1, 1, 1, 0, 0,  c30_bin, false);
-    tkDNN::Region     g31(&net, 80, 4, 5);
+    tk::dnn::Conv2d     c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
+    tk::dnn::Activation a29(&net, tk::dnn::ACTIVATION_LEAKY);   
+    tk::dnn::Conv2d     c30(&net, 425, 1, 1, 1, 1, 0, 0,  c30_bin, false);
+    tk::dnn::Region     g31(&net, 80, 4, 5);
     
-    tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.6f, g31_bin);
+    tk::dnn::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.6f, g31_bin);
 
     // Load input
     dnnType *data;
@@ -110,11 +110,11 @@ int main() {
     net.print();
 
     //convert network to tensorRT
-    tkDNN::NetworkRT netRT(&net, "yolo_224.rt");
+    tk::dnn::NetworkRT netRT(&net, "yolo_224.rt");
 
     dnnType *out_data, *out_data2; // cudnn output, tensorRT output
 
-    tkDNN::dataDim_t dim1 = dim; //input dim
+    tk::dnn::dataDim_t dim1 = dim; //input dim
     printCenteredTitle(" CUDNN inference ", '=', 30); {
         dim1.print();
         TIMER_START
@@ -123,7 +123,7 @@ int main() {
         dim1.print();   
     }
     
-    tkDNN::dataDim_t dim2 = dim;
+    tk::dnn::dataDim_t dim2 = dim;
     printCenteredTitle(" TENSORRT inference ", '=', 30); {
         dim2.print();
         TIMER_START

tests/yolo_berkeley/yolo_berkeley.cpp

@@ -0,0 +1,150 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin  = "../tests/yolo_berkeley/layers/input.bin";
+const char *c0_bin     = "../tests/yolo_berkeley/layers/c0.bin";
+const char *c2_bin     = "../tests/yolo_berkeley/layers/c2.bin";
+const char *c4_bin     = "../tests/yolo_berkeley/layers/c4.bin";
+const char *c5_bin     = "../tests/yolo_berkeley/layers/c5.bin";
+const char *c6_bin     = "../tests/yolo_berkeley/layers/c6.bin";
+const char *c8_bin     = "../tests/yolo_berkeley/layers/c8.bin";
+const char *c9_bin     = "../tests/yolo_berkeley/layers/c9.bin";
+const char *c10_bin    = "../tests/yolo_berkeley/layers/c10.bin";
+const char *c12_bin    = "../tests/yolo_berkeley/layers/c12.bin";
+const char *c13_bin    = "../tests/yolo_berkeley/layers/c13.bin";
+const char *c14_bin    = "../tests/yolo_berkeley/layers/c14.bin";
+const char *c15_bin    = "../tests/yolo_berkeley/layers/c15.bin";
+const char *c16_bin    = "../tests/yolo_berkeley/layers/c16.bin";
+const char *c18_bin    = "../tests/yolo_berkeley/layers/c18.bin";
+const char *c19_bin    = "../tests/yolo_berkeley/layers/c19.bin";
+const char *c20_bin    = "../tests/yolo_berkeley/layers/c20.bin";
+const char *c21_bin    = "../tests/yolo_berkeley/layers/c21.bin";
+const char *c22_bin    = "../tests/yolo_berkeley/layers/c22.bin";
+const char *c23_bin    = "../tests/yolo_berkeley/layers/c23.bin";
+const char *c24_bin    = "../tests/yolo_berkeley/layers/c24.bin";
+const char *c26_bin    = "../tests/yolo_berkeley/layers/c26.bin";
+const char *c29_bin    = "../tests/yolo_berkeley/layers/c29.bin";
+const char *c30_bin    = "../tests/yolo_berkeley/layers/c30.bin";
+const char *g31_bin    = "../tests/yolo_berkeley/layers/g31.bin";
+const char *output_bin = "../tests/yolo_berkeley/layers/output.bin";
+
+int main() {
+
+    // Network layout
+    tk::dnn::dataDim_t dim(1, 3, 416, 736, 1);
+    tk::dnn::Network net(dim);
+
+    tk::dnn::Conv2d     c0 (&net, 32, 3, 3, 1, 1, 1, 1,   c0_bin, true);
+    tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p1 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c2 (&net, 64, 3, 3, 1, 1, 1, 1,   c2_bin, true);
+    tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p3 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c4 (&net, 128, 3, 3, 1, 1, 1, 1,  c4_bin, true);
+    tk::dnn::Activation a4 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c5 (&net, 64, 1, 1, 1, 1, 0, 0,   c5_bin, true);
+    tk::dnn::Activation a5 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
+    tk::dnn::Activation a6 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p7 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c8 (&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
+    tk::dnn::Activation a8 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c9 (&net, 128, 1, 1, 1, 1, 0, 0,  c9_bin, true);
+    tk::dnn::Activation a9 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c10(&net, 256, 3, 3, 1, 1, 1, 1,  c10_bin, true);
+    tk::dnn::Activation a10(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p11(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1,  c12_bin, true);
+    tk::dnn::Activation a12(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c13(&net, 256, 1, 1, 1, 1, 0, 0,  c13_bin, true);
+    tk::dnn::Activation a13(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c14(&net, 512, 3, 3, 1, 1, 1, 1,  c14_bin, true);
+    tk::dnn::Activation a14(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c15(&net, 256, 1, 1, 1, 1, 0, 0,  c15_bin, true);
+    tk::dnn::Activation a15(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c16(&net, 512, 3, 3, 1, 1, 1, 1,  c16_bin, true); 
+    tk::dnn::Activation a16(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p17(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
+    tk::dnn::Activation a18(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c19(&net, 512, 1, 1, 1, 1, 0, 0,  c19_bin, true);
+    tk::dnn::Activation a19(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
+    tk::dnn::Activation a20(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c21(&net, 512, 1, 1, 1, 1, 0, 0,  c21_bin, true);
+    tk::dnn::Activation a21(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
+    tk::dnn::Activation a22(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
+    tk::dnn::Activation a23(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
+    tk::dnn::Activation a24(&net, tk::dnn::ACTIVATION_LEAKY);
+
+    tk::dnn::Layer *m25_layers[1] = { &a16 };
+    tk::dnn::Route      m25(&net, m25_layers, 1);
+    tk::dnn::Conv2d     c26(&net, 64, 1, 1, 1, 1, 0, 0,   c26_bin, true);
+    tk::dnn::Activation a26(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Reorg      r27(&net, 2);
+
+    tk::dnn::Layer *m28_layers[2] = { &r27, &a24 };
+    tk::dnn::Route      m28(&net, m28_layers, 2);
+
+    tk::dnn::Conv2d     c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
+    tk::dnn::Activation a29(&net, tk::dnn::ACTIVATION_LEAKY);   
+    tk::dnn::Conv2d     c30(&net, 75, 1, 1, 1, 1, 0, 0,  c30_bin, false);
+    tk::dnn::Region     g31(&net, 10, 4, 5);
+    
+    tk::dnn::RegionInterpret rI(dim, g31.output_dim, 10, 4, 5, 0.3f, g31_bin);
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+    
+    //print network model
+    net.print();
+
+    //convert network to tensorRT
+    tk::dnn::NetworkRT netRT(&net, "yolo_berkeley.rt");
+
+    dnnType *out_data, *out_data2; // cudnn output, tensorRT output
+
+    tk::dnn::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        out_data = net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+    
+    tk::dnn::dataDim_t dim2 = dim;
+    printCenteredTitle(" TENSORRT inference ", '=', 30); {
+        dim2.print();
+        TIMER_START
+        out_data2 = netRT.infer(dim2, data);
+        TIMER_STOP
+        dim2.print();
+    }
+
+    printCenteredTitle(" CHECK RESULTS ", '=', 30);
+    dnnType *out, *out_h;
+    int out_dim = net.getOutputDim().tot();
+    readBinaryFile(output_bin, out_dim, &out_h, &out);
+    std::cout<<"CUDNN vs correct"; checkResult(out_dim, out_data, out);
+    std::cout<<"TRT   vs correct"; checkResult(out_dim, out_data2, out);
+    std::cout<<"CUDNN vs TRT    "; checkResult(out_dim, out_data, out_data2);
+    
+    std::cout<<"\n\nDetected objects: \n";
+    dnnType *output_h = new dnnType[rI.output_dim.tot()];
+    checkCuda(cudaMemcpy(output_h, out_data2, 
+              rI.output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost));
+    rI.interpretData(output_h);
+    rI.showImageResult(input_h);
+    return 0;
+}

tests/yolo_berkeley/yolov2-voc-10-resize-test.cfg

@@ -0,0 +1,259 @@
+[net]
+# Testing
+batch=1
+subdivisions=1
+# Training
+#batch=64
+#subdivisions=8
+height=416
+width=736
+channels=3
+momentum=0.9
+decay=0.0005
+angle=0
+saturation = 1.5
+exposure = 1.5
+hue=.1
+
+learning_rate=0.001
+burn_in=1000
+max_batches = 80200
+policy=steps
+steps=40000,60000
+scales=.1,.1
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[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]
+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
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+
+#######
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[route]
+layers=-9
+
+[convolutional]
+batch_normalize=1
+size=1
+stride=1
+pad=1
+filters=64
+activation=leaky
+
+[reorg]
+stride=2
+
+[route]
+layers=-1,-4
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=75
+activation=linear
+
+
+[region]
+anchors =  0.4043,0.4167, 1.2109,1.1018, 2.7258,2.1215, 4.9477,3.9132, 7.9508,6.6806
+bias_match=1
+classes=10
+coords=4
+num=5
+softmax=1
+jitter=.3
+rescore=1
+
+object_scale=5
+noobject_scale=1
+class_scale=1
+coord_scale=1
+
+absolute=1
+thresh = .6
+random=0
+flip=1

tests/yolo_relu/yolo_relu.cpp

@@ -31,75 +31,75 @@ const char *output_bin = "../tests/yolo_relu/layers/output.bin";
 int main() {
 
     // Network layout
-    tkDNN::dataDim_t dim(1, 3, 608, 608, 1);
-    tkDNN::Network net(dim);
+    tk::dnn::dataDim_t dim(1, 3, 608, 608, 1);
+    tk::dnn::Network net(dim);
 
-    tkDNN::Conv2d     c0 (&net, 32, 3, 3, 1, 1, 1, 1,   c0_bin, true);
-    tkDNN::Activation a0 (&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Pooling    p1 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c0 (&net, 32, 3, 3, 1, 1, 1, 1,   c0_bin, true);
+    tk::dnn::Activation a0 (&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Pooling    p1 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c2 (&net, 64, 3, 3, 1, 1, 1, 1,   c2_bin, true);
-    tkDNN::Activation a2 (&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Pooling    p3 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c2 (&net, 64, 3, 3, 1, 1, 1, 1,   c2_bin, true);
+    tk::dnn::Activation a2 (&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Pooling    p3 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c4 (&net, 128, 3, 3, 1, 1, 1, 1,  c4_bin, true);
-    tkDNN::Activation a4 (&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c5 (&net, 64, 1, 1, 1, 1, 0, 0,   c5_bin, true);
-    tkDNN::Activation a5 (&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
-    tkDNN::Activation a6 (&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Pooling    p7 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c4 (&net, 128, 3, 3, 1, 1, 1, 1,  c4_bin, true);
+    tk::dnn::Activation a4 (&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c5 (&net, 64, 1, 1, 1, 1, 0, 0,   c5_bin, true);
+    tk::dnn::Activation a5 (&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
+    tk::dnn::Activation a6 (&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Pooling    p7 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c8 (&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
-    tkDNN::Activation a8 (&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c9 (&net, 128, 1, 1, 1, 1, 0, 0,  c9_bin, true);
-    tkDNN::Activation a9 (&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c10(&net, 256, 3, 3, 1, 1, 1, 1,  c10_bin, true);
-    tkDNN::Activation a10(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Pooling    p11(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c8 (&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
+    tk::dnn::Activation a8 (&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c9 (&net, 128, 1, 1, 1, 1, 0, 0,  c9_bin, true);
+    tk::dnn::Activation a9 (&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c10(&net, 256, 3, 3, 1, 1, 1, 1,  c10_bin, true);
+    tk::dnn::Activation a10(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Pooling    p11(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1,  c12_bin, true);
-    tkDNN::Activation a12(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c13(&net, 256, 1, 1, 1, 1, 0, 0,  c13_bin, true);
-    tkDNN::Activation a13(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c14(&net, 512, 3, 3, 1, 1, 1, 1,  c14_bin, true);
-    tkDNN::Activation a14(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c15(&net, 256, 1, 1, 1, 1, 0, 0,  c15_bin, true);
-    tkDNN::Activation a15(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c16(&net, 512, 3, 3, 1, 1, 1, 1,  c16_bin, true); 
-    tkDNN::Activation a16(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Pooling    p17(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1,  c12_bin, true);
+    tk::dnn::Activation a12(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c13(&net, 256, 1, 1, 1, 1, 0, 0,  c13_bin, true);
+    tk::dnn::Activation a13(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c14(&net, 512, 3, 3, 1, 1, 1, 1,  c14_bin, true);
+    tk::dnn::Activation a14(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c15(&net, 256, 1, 1, 1, 1, 0, 0,  c15_bin, true);
+    tk::dnn::Activation a15(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c16(&net, 512, 3, 3, 1, 1, 1, 1,  c16_bin, true); 
+    tk::dnn::Activation a16(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Pooling    p17(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
-    tkDNN::Activation a18(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c19(&net, 512, 1, 1, 1, 1, 0, 0,  c19_bin, true);
-    tkDNN::Activation a19(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
-    tkDNN::Activation a20(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c21(&net, 512, 1, 1, 1, 1, 0, 0,  c21_bin, true);
-    tkDNN::Activation a21(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
-    tkDNN::Activation a22(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
-    tkDNN::Activation a23(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Conv2d     c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
-    tkDNN::Activation a24(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
+    tk::dnn::Activation a18(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c19(&net, 512, 1, 1, 1, 1, 0, 0,  c19_bin, true);
+    tk::dnn::Activation a19(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
+    tk::dnn::Activation a20(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c21(&net, 512, 1, 1, 1, 1, 0, 0,  c21_bin, true);
+    tk::dnn::Activation a21(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
+    tk::dnn::Activation a22(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
+    tk::dnn::Activation a23(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Conv2d     c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
+    tk::dnn::Activation a24(&net, CUDNN_ACTIVATION_RELU);
 
-    tkDNN::Layer *m25_layers[1] = { &a16 };
-    tkDNN::Route      m25(&net, m25_layers, 1);
-    tkDNN::Conv2d     c26(&net, 64, 1, 1, 1, 1, 0, 0,   c26_bin, true);
-    tkDNN::Activation a26(&net, CUDNN_ACTIVATION_RELU);
-    tkDNN::Reorg      r27(&net, 2);
+    tk::dnn::Layer *m25_layers[1] = { &a16 };
+    tk::dnn::Route      m25(&net, m25_layers, 1);
+    tk::dnn::Conv2d     c26(&net, 64, 1, 1, 1, 1, 0, 0,   c26_bin, true);
+    tk::dnn::Activation a26(&net, CUDNN_ACTIVATION_RELU);
+    tk::dnn::Reorg      r27(&net, 2);
 
-    tkDNN::Layer *m28_layers[2] = { &r27, &a24 };
-    tkDNN::Route      m28(&net, m28_layers, 2);
+    tk::dnn::Layer *m28_layers[2] = { &r27, &a24 };
+    tk::dnn::Route      m28(&net, m28_layers, 2);
 
-    tkDNN::Conv2d     c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
-    tkDNN::Activation a29(&net, CUDNN_ACTIVATION_RELU);   
-    tkDNN::Conv2d     c30(&net, 425, 1, 1, 1, 1, 0, 0,  c30_bin, false);
-    tkDNN::Region     g31(&net, 80, 4, 5);
+    tk::dnn::Conv2d     c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
+    tk::dnn::Activation a29(&net, CUDNN_ACTIVATION_RELU);   
+    tk::dnn::Conv2d     c30(&net, 425, 1, 1, 1, 1, 0, 0,  c30_bin, false);
+    tk::dnn::Region     g31(&net, 80, 4, 5);
     
-    tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.3f, g31_bin);
+    tk::dnn::RegionInterpret rI(dim, g31.output_dim, 80, 4, 5, 0.3f, g31_bin);
 
     // Load input
     dnnType *data;
@@ -110,11 +110,11 @@ int main() {
     net.print();
 
     //convert network to tensorRT
-    tkDNN::NetworkRT netRT(&net, "yolo_relu.rt");
+    tk::dnn::NetworkRT netRT(&net, "yolo_relu.rt");
 
     dnnType *out_data, *out_data2; // cudnn output, tensorRT output
 
-    tkDNN::dataDim_t dim1 = dim; //input dim
+    tk::dnn::dataDim_t dim1 = dim; //input dim
     printCenteredTitle(" CUDNN inference ", '=', 30); {
         dim1.print();
         TIMER_START
@@ -123,7 +123,7 @@ int main() {
         dim1.print();   
     }
     
-    tkDNN::dataDim_t dim2 = dim;
+    tk::dnn::dataDim_t dim2 = dim;
     printCenteredTitle(" TENSORRT inference ", '=', 30); {
         dim2.print();
         TIMER_START

tests/yolo_tiny/yolo_tiny.cpp

@@ -18,38 +18,38 @@ const char *output_bin = "../tests/yolo_tiny/layers/output.bin";
 int main() {
 
     // Network layout
-    tkDNN::dataDim_t dim(1, 3, 416, 416, 1);
-    tkDNN::Network net(dim);
+    tk::dnn::dataDim_t dim(1, 3, 416, 416, 1);
+    tk::dnn::Network net(dim);
 
-    tkDNN::Conv2d     c0 (&net, 16, 3, 3, 1, 1, 1, 1,   c0_bin, true);
-    tkDNN::Activation a0 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p1 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c0 (&net, 16, 3, 3, 1, 1, 1, 1,   c0_bin, true);
+    tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p1 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c2 (&net, 32, 3, 3, 1, 1, 1, 1,   c2_bin, true);
-    tkDNN::Activation a2 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p3 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c2 (&net, 32, 3, 3, 1, 1, 1, 1,   c2_bin, true);
+    tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p3 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c4 (&net, 64, 3, 3, 1, 1, 1, 1,  c4_bin, true);
-    tkDNN::Activation a4 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p5 (&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c4 (&net, 64, 3, 3, 1, 1, 1, 1,  c4_bin, true);
+    tk::dnn::Activation a4 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p5 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
-    tkDNN::Activation a6 (&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p7(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
+    tk::dnn::Activation a6 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p7(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c8(&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
-    tkDNN::Activation a8(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Pooling    p9(&net, 2, 2, 2, 2, tkDNN::POOLING_MAX);
+    tk::dnn::Conv2d     c8(&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
+    tk::dnn::Activation a8(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p9(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
 
-    tkDNN::Conv2d     c10(&net, 512, 3, 3, 1, 1, 1, 1, c10_bin, true);
-    tkDNN::Activation a10(&net, tkDNN::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c10(&net, 512, 3, 3, 1, 1, 1, 1, c10_bin, true);
+    tk::dnn::Activation a10(&net, tk::dnn::ACTIVATION_LEAKY);
 
-    tkDNN::Conv2d     c11(&net, 1024, 3, 3, 1, 1, 1, 1, c11_bin, true);
-    tkDNN::Activation a11(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1, c12_bin, true);
-    tkDNN::Activation a12(&net, tkDNN::ACTIVATION_LEAKY);
-    tkDNN::Conv2d     c13(&net, 425, 1, 1, 1, 1, 0, 0, c13_bin, false);
-    tkDNN::Region     g14(&net, 80, 4, 5);
+    tk::dnn::Conv2d     c11(&net, 1024, 3, 3, 1, 1, 1, 1, c11_bin, true);
+    tk::dnn::Activation a11(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1, c12_bin, true);
+    tk::dnn::Activation a12(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c13(&net, 425, 1, 1, 1, 1, 0, 0, c13_bin, false);
+    tk::dnn::Region     g14(&net, 80, 4, 5);
 
     // Load input
     dnnType *data;
@@ -60,11 +60,11 @@ int main() {
     net.print();
 
     //convert network to tensorRT
-    tkDNN::NetworkRT netRT(&net, "yolo_tiny.rt");
+    tk::dnn::NetworkRT netRT(&net, "yolo_tiny.rt");
 
     dnnType *out_data, *out_data2; // cudnn output, tensorRT output
 
-    tkDNN::dataDim_t dim1 = dim; //input dim
+    tk::dnn::dataDim_t dim1 = dim; //input dim
     printCenteredTitle(" CUDNN inference ", '=', 30); {
         dim1.print();
         TIMER_START
@@ -73,7 +73,7 @@ int main() {
         dim1.print();   
     }
  
-    tkDNN::dataDim_t dim2 = dim;
+    tk::dnn::dataDim_t dim2 = dim;
     printCenteredTitle(" TENSORRT inference ", '=', 30); {
         dim2.print();
         TIMER_START

tests/yolo_voc/yolo_voc.cfg

@@ -0,0 +1,258 @@
+[net]
+# Testing
+batch=1
+subdivisions=1
+# Training
+# batch=64
+# subdivisions=8
+height=416
+width=416
+channels=3
+momentum=0.9
+decay=0.0005
+angle=0
+saturation = 1.5
+exposure = 1.5
+hue=.1
+
+learning_rate=0.001
+burn_in=1000
+max_batches = 80200
+policy=steps
+steps=40000,60000
+scales=.1,.1
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[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]
+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
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=leaky
+
+
+#######
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[route]
+layers=-9
+
+[convolutional]
+batch_normalize=1
+size=1
+stride=1
+pad=1
+filters=64
+activation=leaky
+
+[reorg]
+stride=2
+
+[route]
+layers=-1,-4
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=125
+activation=linear
+
+
+[region]
+anchors =  1.3221, 1.73145, 3.19275, 4.00944, 5.05587, 8.09892, 9.47112, 4.84053, 11.2364, 10.0071
+bias_match=1
+classes=20
+coords=4
+num=5
+softmax=1
+jitter=.3
+rescore=1
+
+object_scale=5
+noobject_scale=1
+class_scale=1
+coord_scale=1
+
+absolute=1
+thresh = .6
+random=1

tests/yolo_voc/yolo_voc.cpp

@@ -0,0 +1,150 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin  = "../tests/yolo_voc/layers/input.bin";
+const char *c0_bin     = "../tests/yolo_voc/layers/c0.bin";
+const char *c2_bin     = "../tests/yolo_voc/layers/c2.bin";
+const char *c4_bin     = "../tests/yolo_voc/layers/c4.bin";
+const char *c5_bin     = "../tests/yolo_voc/layers/c5.bin";
+const char *c6_bin     = "../tests/yolo_voc/layers/c6.bin";
+const char *c8_bin     = "../tests/yolo_voc/layers/c8.bin";
+const char *c9_bin     = "../tests/yolo_voc/layers/c9.bin";
+const char *c10_bin    = "../tests/yolo_voc/layers/c10.bin";
+const char *c12_bin    = "../tests/yolo_voc/layers/c12.bin";
+const char *c13_bin    = "../tests/yolo_voc/layers/c13.bin";
+const char *c14_bin    = "../tests/yolo_voc/layers/c14.bin";
+const char *c15_bin    = "../tests/yolo_voc/layers/c15.bin";
+const char *c16_bin    = "../tests/yolo_voc/layers/c16.bin";
+const char *c18_bin    = "../tests/yolo_voc/layers/c18.bin";
+const char *c19_bin    = "../tests/yolo_voc/layers/c19.bin";
+const char *c20_bin    = "../tests/yolo_voc/layers/c20.bin";
+const char *c21_bin    = "../tests/yolo_voc/layers/c21.bin";
+const char *c22_bin    = "../tests/yolo_voc/layers/c22.bin";
+const char *c23_bin    = "../tests/yolo_voc/layers/c23.bin";
+const char *c24_bin    = "../tests/yolo_voc/layers/c24.bin";
+const char *c26_bin    = "../tests/yolo_voc/layers/c26.bin";
+const char *c29_bin    = "../tests/yolo_voc/layers/c29.bin";
+const char *c30_bin    = "../tests/yolo_voc/layers/c30.bin";
+const char *g31_bin    = "../tests/yolo_voc/layers/g31.bin";
+const char *output_bin = "../tests/yolo_voc/layers/output.bin";
+
+int main() {
+
+    // Network layout
+    tk::dnn::dataDim_t dim(1, 3, 416, 416, 1);
+    tk::dnn::Network net(dim);
+
+    tk::dnn::Conv2d     c0 (&net, 32, 3, 3, 1, 1, 1, 1,   c0_bin, true);
+    tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p1 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c2 (&net, 64, 3, 3, 1, 1, 1, 1,   c2_bin, true);
+    tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p3 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c4 (&net, 128, 3, 3, 1, 1, 1, 1,  c4_bin, true);
+    tk::dnn::Activation a4 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c5 (&net, 64, 1, 1, 1, 1, 0, 0,   c5_bin, true);
+    tk::dnn::Activation a5 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c6 (&net, 128, 3, 3, 1, 1, 1, 1,  c6_bin, true);
+    tk::dnn::Activation a6 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p7 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c8 (&net, 256, 3, 3, 1, 1, 1, 1,  c8_bin, true);
+    tk::dnn::Activation a8 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c9 (&net, 128, 1, 1, 1, 1, 0, 0,  c9_bin, true);
+    tk::dnn::Activation a9 (&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c10(&net, 256, 3, 3, 1, 1, 1, 1,  c10_bin, true);
+    tk::dnn::Activation a10(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p11(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c12(&net, 512, 3, 3, 1, 1, 1, 1,  c12_bin, true);
+    tk::dnn::Activation a12(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c13(&net, 256, 1, 1, 1, 1, 0, 0,  c13_bin, true);
+    tk::dnn::Activation a13(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c14(&net, 512, 3, 3, 1, 1, 1, 1,  c14_bin, true);
+    tk::dnn::Activation a14(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c15(&net, 256, 1, 1, 1, 1, 0, 0,  c15_bin, true);
+    tk::dnn::Activation a15(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c16(&net, 512, 3, 3, 1, 1, 1, 1,  c16_bin, true); 
+    tk::dnn::Activation a16(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Pooling    p17(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX);
+
+    tk::dnn::Conv2d     c18(&net, 1024, 3, 3, 1, 1, 1, 1, c18_bin, true);
+    tk::dnn::Activation a18(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c19(&net, 512, 1, 1, 1, 1, 0, 0,  c19_bin, true);
+    tk::dnn::Activation a19(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c20(&net, 1024, 3, 3, 1, 1, 1, 1, c20_bin, true);
+    tk::dnn::Activation a20(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c21(&net, 512, 1, 1, 1, 1, 0, 0,  c21_bin, true);
+    tk::dnn::Activation a21(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c22(&net, 1024, 3, 3, 1, 1, 1, 1, c22_bin, true);
+    tk::dnn::Activation a22(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c23(&net, 1024, 3, 3, 1, 1, 1, 1, c23_bin, true);
+    tk::dnn::Activation a23(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Conv2d     c24(&net, 1024, 3, 3, 1, 1, 1, 1, c24_bin, true);
+    tk::dnn::Activation a24(&net, tk::dnn::ACTIVATION_LEAKY);
+
+    tk::dnn::Layer *m25_layers[1] = { &a16 };
+    tk::dnn::Route      m25(&net, m25_layers, 1);
+    tk::dnn::Conv2d     c26(&net, 64, 1, 1, 1, 1, 0, 0,   c26_bin, true);
+    tk::dnn::Activation a26(&net, tk::dnn::ACTIVATION_LEAKY);
+    tk::dnn::Reorg      r27(&net, 2);
+
+    tk::dnn::Layer *m28_layers[2] = { &r27, &a24 };
+    tk::dnn::Route      m28(&net, m28_layers, 2);
+
+    tk::dnn::Conv2d     c29(&net, 1024, 3, 3, 1, 1, 1, 1, c29_bin, true);
+    tk::dnn::Activation a29(&net, tk::dnn::ACTIVATION_LEAKY);   
+    tk::dnn::Conv2d     c30(&net, 125, 1, 1, 1, 1, 0, 0,  c30_bin, false);
+    tk::dnn::Region     g31(&net, 20, 4, 5);
+    
+    tk::dnn::RegionInterpret rI(dim, g31.output_dim, 20, 4, 5, 0.6f, g31_bin);
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+    
+    //print network model
+    net.print();
+
+    //convert network to tensorRT
+    tk::dnn::NetworkRT netRT(&net, "yolo_voc.rt");
+
+    dnnType *out_data, *out_data2; // cudnn output, tensorRT output
+
+    tk::dnn::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        out_data = net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+    
+    tk::dnn::dataDim_t dim2 = dim;
+    printCenteredTitle(" TENSORRT inference ", '=', 30); {
+        dim2.print();
+        TIMER_START
+        out_data2 = netRT.infer(dim2, data);
+        TIMER_STOP
+        dim2.print();
+    }
+
+    printCenteredTitle(" CHECK RESULTS ", '=', 30);
+    dnnType *out, *out_h;
+    int out_dim = net.getOutputDim().tot();
+    readBinaryFile(output_bin, out_dim, &out_h, &out);
+    std::cout<<"CUDNN vs correct"; checkResult(out_dim, out_data, out);
+    std::cout<<"TRT   vs correct"; checkResult(out_dim, out_data2, out);
+    std::cout<<"CUDNN vs TRT    "; checkResult(out_dim, out_data, out_data2);
+    
+    std::cout<<"\n\nDetected objects: \n";
+    dnnType *output_h = new dnnType[rI.output_dim.tot()];
+    checkCuda(cudaMemcpy(output_h, out_data2, 
+              rI.output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost));
+    rI.interpretData(output_h);
+    rI.showImageResult(input_h);
+    return 0;
+}