fp16 implementation, TODO deallocate in LayerWgs

.gitignore

@@ -2,4 +2,9 @@
 build/
 .vscode/
 *.bin
-*.pyc
+*.pyc
+*.prototxt
+*.caffemodel
+*.h5
+*.tar.gz
+*.weights

CMakeLists.txt

@@ -1,20 +1,91 @@
 cmake_minimum_required(VERSION 2.8)
-
 project (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()
+
+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)
+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)
 
-include_directories(${CMAKE_CURRENT_SOURCE_DIR}/include ${CUDA_INCLUDE_DIRS})
-add_library(tkDNN SHARED src/Layer.cpp src/LayerWgs.cpp
-    src/Dense.cpp src/Activation.cpp src/Conv2d.cpp src/Flatten.cpp src/MulAdd.cpp src/Pooling.cpp src/Softmax.cpp
-    src/Network.cpp src/utils.cpp)
-target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn)
+file(GLOB tkdnn_SRC "src/*.cpp")
+set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer ${OpenCV_LIBS})
 
-add_executable(test_simple tests/test/test.cpp)
+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})
+add_library(tkDNN SHARED ${tkdnn_SRC})
+target_link_libraries(tkDNN ${tkdnn_LIBS})
+
+#static 
+#add_library(tkDNN_static STATIC ${tkdnn_SRC})
+#target_link_libraries(tkDNN_static ${tkdnn_LIBS})
+
+add_executable(test_simple tests/simple/test_simple.cpp)
 target_link_libraries(test_simple tkDNN)
 
-add_executable(test_mnist tests/mnist/test.cpp)
+add_executable(test_mnist tests/mnist/test_mnist.cpp)
 target_link_libraries(test_mnist tkDNN)
+
+add_executable(test_mnistRT tests/mnist/test_mnistRT.cpp)
+target_link_libraries(test_mnistRT tkDNN)
+
+## YOLO NETS
+add_executable(test_yolo tests/yolo/yolo.cpp)
+target_link_libraries(test_yolo tkDNN)
+
+add_executable(test_yolo_tiny tests/yolo_tiny/yolo_tiny.cpp)
+target_link_libraries(test_yolo_tiny tkDNN)
+
+add_executable(test_yolo_relu tests/yolo_relu/yolo_relu.cpp)
+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_rtinference tests/test_rtinference/rtinference.cpp)
+target_link_libraries(test_rtinference tkDNN)
+
+add_executable(detection demo/detection/detection.cpp)
+target_link_libraries(detection tkDNN)
+
+#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(TARGETS tkDNN kernels DESTINATION lib)

README.md

@@ -1,15 +1,11 @@
 # tkDNN
 tkDNN is a Deep Neural Network library built with cuDNN primitives specifically thought to work on NVIDIA TK1 board.<br>
 The main scope is to do high performance inference on already trained models.
-Currently supports the following layers:
 
-* Dense, fully interconnected
-* Activation (RELU, ELU, SIGMOID, TANH)
-* Convolutional 2D
-* Convolutional 3D
-* Max and Average Pooling
-* Flatten 
-* Data preprocessing
+this branch is actually work on every NVIDIA GPU that support the dependencies:
+* CUDA 8
+* CUDNN 6
+* TENSORRT 2
 
 ## Workflow
 The recommended workflow follow these step:
@@ -26,60 +22,14 @@ cd build
 cmake ..
 make
 ```
+during the cmake configuration it will be dowloaded the weights needed for running
+the tests
 
 ## Test
-There is a ready to use example on *test* directory, to try it you must generate the weights with Keras
-```
-cd tests
-python test_model.py
-```
-And then execute the inference on build directory
-```
-cd build
-./tkDNNtest 
-```
-this should output the same prediction as Keras.
+Assumiung you have correctly builded the library these are the test ready to exec:
+* test_simple: a simple convolutional and dense network (CUDNN only)
+* test_mnist: the famous mnist netwok (CUDNN and TENSORRT)
+* 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)
 
-## Simple example
-Here is a example of the entire workflow on a simple model.
-Using the following Keras model save it to a file
-```python
-model = Sequential()
-model.add(Reshape((20, 1), input_shape=(20)))
-model.add(Dense(256))
-model.compile()
-
-# save model
-model.save("path/to/model.h5")
-```
-
-After the model is created the weights can be exported for tkDNN inference
-```
-python weights_exporter model.h5 dense --output=weights/path
-```
-the exporter take as arguments, in order:
-* input model
-* layer type ["dense", "conv2d", conv3d"]
-* { layer type ["dense", "conv2d", conv3d"] for each layer to export }
-* optional argument --output define path where export weights   
-
-Then we can create a c++ program to do inference on tk1
-```c++
-#include<tkdnn.h>   //library include
-
-//Network object 
-tkDNN::Network net;
-//input dimension
-tkDNN::dataDim_t dim(1, 20, 1, 1, 1);
-//Dense layer
-tkDNN::Dense d0(&net, dim, 256, "weights/path", "bias/path");
-
-//here load the input data to CUDA
-//value_type is an alias of "float"
-value_type *data_d = [...] 
-
-//do inference
-value_type *output_d = d0.infer(dim, data_d);
-//dim will be updated with the output dimension
-``` 
-The result is finally stored on output_d in device memory.

demo/detection/detection.cpp

@@ -0,0 +1,249 @@
+#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/Layer.h

@@ -7,52 +7,58 @@
 
 namespace tkDNN {
 
-/**
-    Data rapresentation beetween layers
-    n = batch size
-    c = channels
-    h = heigth (lines)
-    w = width  (rows)
-    l = lenght (3rd dimension)
-*/
-struct dataDim_t {
-
-    int n, c, h, w, l;
-
-    dataDim_t() : n(1), c(1), h(1), w(1), l(1) {};
-
-    dataDim_t(int _n, int _c, int _h, int _w, int _l = 1) :
-        n(_n), c(_c), h(_h), w(_w), l(_l) {};
-
-    void print() {
-        std::cout<<"Data dim: "<<n<<" "<<c<<" "<<h<<" "<<w<<" "<<l<<"\n";
-    }
-
-    int tot() {
-        return n*c*h*w*l;
-    }
+enum layerType_t {
+    LAYER_DENSE,
+    LAYER_CONV2D,
+    LAYER_ACTIVATION,
+    LAYER_FLATTEN,
+    LAYER_MULADD,
+    LAYER_POOLING,
+    LAYER_SOFTMAX,
+    LAYER_ROUTE,
+    LAYER_REORG,
+    LAYER_REGION
 };
 
-
 /** 
     Simple layer Father class
 */
 class Layer {
 
 public:
-    Layer(Network *net, dataDim_t input_dim);
+    Layer(Network *net);
     virtual ~Layer();
+    virtual layerType_t getLayerType() = 0;
 
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData) {
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData) {
         std::cout<<"No infer action for this layer\n";
         return NULL;
     }
 
     dataDim_t input_dim, output_dim;
+    dnnType *dstData;  //where results will be putted
+
+    std::string getLayerName() {
+        layerType_t type = getLayerType();
+        switch(type) {
+            case LAYER_DENSE:       return "Dense";
+            case LAYER_CONV2D:      return "Conv2d";
+            case LAYER_ACTIVATION:  return "Activation";
+            case LAYER_FLATTEN:     return "Flatten";
+            case LAYER_MULADD:      return "MulAdd";
+            case LAYER_POOLING:     return "Pooling";
+            case LAYER_SOFTMAX:     return "Softmax";
+            case LAYER_ROUTE:       return "Route";            
+            case LAYER_REORG:       return "Reorg";
+            case LAYER_REGION:      return "Region";
+            default:                return "unknown";
+        }
+    }
 
 protected:
     Network *net;
     cudnnTensorDescriptor_t srcTensorDesc, dstTensorDesc;
+
 };
 
 
@@ -62,17 +68,31 @@ protected:
 class LayerWgs : public Layer {
 
 public:
-    LayerWgs(Network *net, dataDim_t input_dim, 
-             int inputs, int outputs, int kh, int kw, int kt, 
-             const char* fname_weights, const char* fname_bias); 
+    LayerWgs(Network *net, int inputs, int outputs, int kh, int kw, int kt, 
+             const char* fname_weights, bool batchnorm = false); 
     virtual ~LayerWgs();
 
-protected:
     int inputs, outputs;
-    std::string weights_path, bias_path;
+    std::string weights_path;
 
-    value_type *data_h, *data_d;
-    value_type *bias_h, *bias_d;
+    dnnType *data_h, *data_d;
+    dnnType *bias_h, *bias_d;
+
+    //batchnorm
+    bool batchnorm;
+    dnnType *power_h;
+    dnnType *scales_h,   *scales_d;
+    dnnType *mean_h,     *mean_d;
+    dnnType *variance_h, *variance_d;
+
+    //fp16
+    __half *data16_h, *bias16_h;
+    __half *data16_d, *bias16_d;
+
+    __half *power16_h,    *power16_d;
+    __half *scales16_h,   *scales16_d;
+    __half *mean16_h,     *mean16_d;
+    __half *variance16_h, *variance16_d;
 };
 
 
@@ -82,32 +102,38 @@ protected:
 class Dense : public LayerWgs {
 
 public:
-    Dense(Network *net, dataDim_t in_dim, int out_ch, 
-          const char* fname_weights, const char* fname_bias); 
+    Dense(Network *net, int out_ch, const char* fname_weights); 
     virtual ~Dense();
+    virtual layerType_t getLayerType() { return LAYER_DENSE; };
 
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
-
-protected:
-    value_type *dstData;  //where results will be putted
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
 };
 
 
+/**
+    Avaible activation functions
+*/
+typedef enum {
+    ACTIVATION_ELU     = 100,
+    ACTIVATION_LEAKY   = 101
+} tkdnnActivationMode_t;
+
 /**
     Activation layer (it doesnt need weigths)
 */
 class Activation : public Layer {
 
 public:
-    Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode); 
-    virtual ~Activation();
+    int act_mode;
 
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
+    Activation(Network *net, int act_mode); 
+    virtual ~Activation();
+    virtual layerType_t getLayerType() { return LAYER_ACTIVATION; };
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
 
 protected:
-    cudnnActivationMode_t act_mode;
     cudnnActivationDescriptor_t activDesc;
-    value_type *dstData;  //where results will be putted
 };
 
 
@@ -117,17 +143,17 @@ protected:
 class Conv2d : public LayerWgs {
 
 public:
-    Conv2d(Network *net, dataDim_t in_dim, int out_ch,
-            int kernelH, int kernelW, int strideH, int strideW,
-            const char* fname_weights, const char* fname_bias); 
+    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); 
     virtual ~Conv2d();
+    virtual layerType_t getLayerType() { return LAYER_CONV2D; };
 
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+
+    int kernelH, kernelW, strideH, strideW, paddingH, paddingW;
 
 protected:
-    value_type *dstData;  //where results will be putted
-    int kernelH, kernelW, strideH, strideW;
-
     cudnnFilterDescriptor_t filterDesc;
     cudnnConvolutionDescriptor_t convDesc;
     cudnnConvolutionFwdAlgo_t algo;
@@ -145,13 +171,11 @@ protected:
 class Flatten : public Layer {
 
 public:
-    Flatten(Network *net, dataDim_t input_dim); 
+    Flatten(Network *net); 
     virtual ~Flatten();
+    virtual layerType_t getLayerType() { return LAYER_FLATTEN; };
 
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
-
-protected:
-    value_type *dstData;  //where results will be putted
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
 };
 
 
@@ -162,14 +186,15 @@ protected:
 class MulAdd : public Layer {
 
 public:
-    MulAdd(Network *net, dataDim_t input_dim, value_type mul, value_type add); 
+    MulAdd(Network *net, dnnType mul, dnnType add); 
     virtual ~MulAdd();
+    virtual layerType_t getLayerType() { return LAYER_MULADD; };
 
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
 
 protected:
-    value_type mul, add;
-    value_type *dstData, *add_vector;  //where results will be putted
+    dnnType mul, add;
+    dnnType *add_vector;
 };
 
 
@@ -190,20 +215,22 @@ typedef enum {
 class Pooling : public Layer {
 
 public:
-    Pooling(Network *net, dataDim_t input_dim, int winH, int winW, 
+    int winH, winW;
+    int strideH, strideW;
+    int paddingH, paddingW;
+
+    Pooling(Network *net, int winH, int winW, 
             int strideH, int strideW, tkdnnPoolingMode_t pool_mode); 
     virtual ~Pooling();
+    virtual layerType_t getLayerType() { return LAYER_POOLING; };
 
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
 
 protected:
 
     cudnnPoolingDescriptor_t poolingDesc;
-
-    int winH, winW;
-    int strideH, strideW;
     tkdnnPoolingMode_t pool_mode;
-    value_type *dstData, *tmpInputData, *tmpOutputData;  //where results will be putted
+    dnnType *tmpInputData, *tmpOutputData;
     bool poolOn3d;
 };
 
@@ -213,13 +240,104 @@ protected:
 class Softmax : public Layer {
 
 public:
-    Softmax(Network *net, dataDim_t input_dim); 
+    Softmax(Network *net); 
     virtual ~Softmax();
+    virtual layerType_t getLayerType() { return LAYER_SOFTMAX; };
 
-    virtual value_type* infer(dataDim_t &dim, value_type* srcData);
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+};
 
-protected:
-    value_type *dstData;  //where results will be putted
+/**
+    Route layer
+    Merge a list of layers
+*/
+class Route : public Layer {
+
+public:
+    Route(Network *net, Layer **layers, int layers_n); 
+    virtual ~Route();
+    virtual layerType_t getLayerType() { return LAYER_ROUTE; };
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+
+public:
+    Layer **layers;  //ids of layers to be merged
+    int layers_n; //number of layers
+};
+
+
+/**
+    Reorg layer
+    Mantain same dimension but change C*H*W distribution
+*/
+class Reorg : public Layer {
+
+public:
+    Reorg(Network *net, int stride);
+    virtual ~Reorg();
+    virtual layerType_t getLayerType() { return LAYER_REORG; };
+
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+
+    int stride;
+};
+
+
+struct box {
+    int cl;
+    float x, y, w, h;
+    float prob;
+};
+struct sortable_bbox {
+    int index;
+    int cl;
+    float **probs;
+};
+
+/**
+    Region layer
+    Mantain same dimension but change C*H*W distribution
+*/
+class Region : public Layer {
+
+public:
+    Region(Network *net, int classes, int coords, int num);
+    virtual ~Region();
+    virtual layerType_t getLayerType() { return LAYER_REGION; };
+
+    int classes, coords, num;
+    
+    virtual dnnType* infer(dataDim_t &dim, dnnType* srcData);
+};
+
+class RegionInterpret {
+
+public:
+    RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, 
+                    int classes, int coords, int num, float thresh, const char* fname_weights);
+    ~RegionInterpret();
+
+    dataDim_t input_dim, output_dim;
+    dnnType *bias_h, *bias_d; //anchors
+    int classes, coords, num;
+    float thresh;
+
+    
+    box *boxes;
+    float **probs;
+    sortable_bbox *s;
+    box res_boxes[256];
+    int res_boxes_n;
+
+    box get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h, int stride);
+    void get_region_boxes(  float *input, int w, int h, int netw, int neth, float thresh, 
+                            float **probs, box *boxes, int only_objectness, 
+                            int *map, float tree_thresh, int relative); 
+    void correct_region_boxes(box *boxes, int n, int w, int h, int netw, int neth, int relative);
+    void interpretData(dnnType *data_h, int imageW = 0, int imageH = 0);
+    void showImageResult(dnnType *input_h);
+
+    static float box_iou(box a, box b);
 };
 
 }

include/Network.h

@@ -5,32 +5,62 @@
 
 namespace tkDNN {
 
-struct dataDim_t;
+/**
+    Data rapresentation beetween layers
+    n = batch size
+    c = channels
+    h = heigth (lines)
+    w = width  (rows)
+    l = lenght (3rd dimension)
+*/
+struct dataDim_t {
+
+    int n, c, h, w, l;
+
+    dataDim_t() : n(1), c(1), h(1), w(1), l(1) {};
+
+    dataDim_t(int _n, int _c, int _h, int _w, int _l = 1) :
+        n(_n), c(_c), h(_h), w(_w), l(_l) {};
+
+    void print() {
+        std::cout<<"Data dim: "<<n<<" "<<c<<" "<<h<<" "<<w<<" "<<l<<"\n";
+    }
+
+    int tot() {
+        return n*c*h*w*l;
+    }
+};
+
 class Layer;
 const int MAX_LAYERS = 256;
 
 class Network {
 
 public:
-    Network();
+    Network(dataDim_t input_dim);
     virtual ~Network();
 
     /**
         Do inferece for every added layer
     */
-    value_type* infer(dataDim_t &dim, value_type* data);
+    dnnType* infer(dataDim_t &dim, dnnType* data);
 
     bool addLayer(Layer *l);
+    void print();
 
     cudnnDataType_t dataType;
     cudnnTensorFormat_t tensorFormat;
     cudnnHandle_t cudnnHandle;
     cublasHandle_t cublasHandle;
 
-private:
     Layer* layers[MAX_LAYERS];  //contains layers of the net
     int num_layers; //current number of layers
+
+    dataDim_t input_dim;
+    dataDim_t getOutputDim();
+
+    bool fp16;
 };
 
 }
-#endif //NETWORK_H
+#endif //NETWORK_H

include/NetworkRT.h

@@ -0,0 +1,66 @@
+#ifndef NETWORKRT_H
+#define NETWORKRT_H
+
+#include "utils.h"
+#include "Network.h"
+#include "Layer.h"
+#include "NvInfer.h"
+
+namespace tkDNN {
+
+class NetworkRT {
+
+public:
+    nvinfer1::DataType dtRT;
+    nvinfer1::IBuilder *builderRT;
+    nvinfer1::IRuntime *runtimeRT;
+    nvinfer1::INetworkDefinition *networkRT; 
+    
+    nvinfer1::ICudaEngine *engineRT;
+    nvinfer1::IExecutionContext *contextRT;
+    void* buffersRT[2];
+    int buf_input_idx, buf_output_idx;
+
+    dataDim_t input_dim, output_dim;
+    dnnType *output;
+    cudaStream_t stream;
+
+    NetworkRT(Network *net, const char *name);
+    virtual ~NetworkRT();
+
+    /**
+        Do inferece
+    */
+    dnnType* infer(dataDim_t &dim, dnnType* data);
+    void enqueue();    
+
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l);
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l);
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Activation *l);
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Dense *l);
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Pooling *l);
+    nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Softmax *l);
+    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);
+
+    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/kernels.h

@@ -1,3 +1,17 @@
+#ifndef KERNELS_H
+#define KERNELS_H
+
 #include "utils.h"
 
-void activationELUForward(value_type* srcData, value_type* dstData, int size);
+void activationELUForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
+void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
+void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, 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 float2half(float* srcData, __half* dstData, int size, const cudaStream_t stream = cudaStream_t(0));
+#endif //KERNELS_H

include/tkdnn.h

@@ -3,14 +3,6 @@
 */
 #include "Network.h"
 #include "Layer.h"
+#include "NetworkRT.h"
 
-namespace tkDNN {
-
-    /**
-        Return the tkDNN version
-    */
-    int getVersion() {
-
-        return 100;
-    }
-}
+#define TKDNN_VERSION 200

include/utils.h

@@ -12,17 +12,35 @@
 #include <cublas_v2.h>
 #include <cudnn.h>
 
-#define value_type float
+#define dnnType float
 
+// Colored output
+#define COL_END "\033[0m"
 
+#define COL_RED "\033[31m"
+#define COL_GREEN "\033[32m"
+#define COL_ORANGE "\033[33m"
+#define COL_BLUE "\033[34m"
+#define COL_PURPLE "\033[35m"
+#define COL_CYAN "\033[36m"
+
+#define COL_REDB "\033[1;31m"
+#define COL_GREENB "\033[1;32m"
+#define COL_ORANGEB "\033[1;33m"
+#define COL_BLUEB "\033[1;34m"
+#define COL_PURPLEB "\033[1;35m"
+#define COL_CYANB "\033[1;36m"
+
+// Simple Timer 
 #define TIMER_START timespec start, end;                               \
                     clock_gettime(CLOCK_MONOTONIC, &start);            
 
-#define TIMER_STOP  clock_gettime(CLOCK_MONOTONIC, &end);              \
+#define TIMER_STOP_C(col)  clock_gettime(CLOCK_MONOTONIC, &end);       \
     double t_ns = ((double)(end.tv_sec - start.tv_sec) * 1.0e9 +       \
                   (double)(end.tv_nsec - start.tv_nsec))/1.0e6;        \
-    std::cout<<"Time:"<<std::setw(16)<<t_ns<<" ms\n"; 
+    std::cout<<col<<"Time:"<<std::setw(16)<<t_ns<<" ms\n"<<COL_END; 
 
+#define TIMER_STOP TIMER_STOP_C(COL_CYANB)
 
 /********************************************************
  * Prints the error message, and exits
@@ -62,12 +80,23 @@
     }                                                                  \
 }
 
-void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d);
-void printDeviceVector(int size, value_type* vec_d);
-void resize(int size, value_type **data);
+#define checkNULL(ptr) {                                               \
+    std::stringstream _error;                                          \
+    if (ptr == nullptr) {                                              \
+      _error << "Null pointer";                                        \
+      FatalError(_error.str());                                        \
+    }                                                                  \
+}
 
-void matrixTranspose(cublasHandle_t handle, value_type* srcData, value_type* dstData, int rows, int cols);
+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);
+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);
 
-void matrixMulAdd(  cublasHandle_t handle, value_type* srcData, value_type* dstData, 
-                    value_type* add_vector, int dim, value_type mul);
-#endif //UTILS_H
+void matrixTranspose(cublasHandle_t handle, dnnType* srcData, dnnType* dstData, int rows, int cols);
+
+void matrixMulAdd(  cublasHandle_t handle, dnnType* srcData, dnnType* dstData, 
+                    dnnType* add_vector, int dim, dnnType mul);
+#endif //UTILS_H

src/Activation.cpp

@@ -5,52 +5,61 @@
 
 namespace tkDNN {
 
-Activation::Activation(Network *net, dataDim_t input_dim, cudnnActivationMode_t act_mode) : 
-    Layer(net, input_dim) {
+Activation::Activation(Network *net, int act_mode) : 
+    Layer(net) {
 
     this->act_mode = act_mode;
-    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
+    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
 
-    checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc,
-                                        net->tensorFormat,
-                                        net->dataType,
-                                        input_dim.n*input_dim.l, 
-                                        input_dim.c,
-                                        input_dim.h, input_dim.w) );
-    checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
+    if(int(act_mode) < 100) {
+
+        checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc,
                                             net->tensorFormat,
                                             net->dataType,
                                             input_dim.n*input_dim.l, 
                                             input_dim.c,
                                             input_dim.h, input_dim.w) );
+        checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
+                                                net->tensorFormat,
+                                                net->dataType,
+                                                input_dim.n*input_dim.l, 
+                                                input_dim.c,
+                                                input_dim.h, input_dim.w) );
 
 
-    checkCUDNN( cudnnCreateActivationDescriptor(&activDesc) );
-    checkCUDNN( cudnnSetActivationDescriptor(activDesc,
-                                             act_mode,
-                                             CUDNN_PROPAGATE_NAN,
-                                             0.0) );
+        checkCUDNN( cudnnCreateActivationDescriptor(&activDesc) );
+        checkCUDNN( cudnnSetActivationDescriptor(activDesc,
+                                                (cudnnActivationMode_t) act_mode,
+                                                CUDNN_PROPAGATE_NAN,
+                                                0.0) );
+    }
 }
 
 Activation::~Activation() {
 
     checkCuda( cudaFree(dstData) );
 
-    checkCUDNN( cudnnDestroyActivationDescriptor(activDesc) );
+    if(int(act_mode) < 100)
+        checkCUDNN( cudnnDestroyActivationDescriptor(activDesc) );
 }
 
-value_type* Activation::infer(dataDim_t &dim, value_type* srcData) {
+dnnType* Activation::infer(dataDim_t &dim, dnnType* srcData) {
 
-    value_type alpha = value_type(1);
-    value_type beta  = value_type(0);
-    checkCUDNN( cudnnActivationForward(net->cudnnHandle,
-                                        activDesc,
-                                        &alpha,
-                                        srcTensorDesc,
-                                        srcData,
-                                        &beta,
-                                        dstTensorDesc,
-                                        dstData) );    
+    if(act_mode == ACTIVATION_LEAKY) {
+        activationLEAKYForward(srcData, dstData, dim.tot());
+    
+    } else {
+        dnnType alpha = dnnType(1);
+        dnnType beta  = dnnType(0);
+        checkCUDNN( cudnnActivationForward(net->cudnnHandle,
+                                            activDesc,
+                                            &alpha,
+                                            srcTensorDesc,
+                                            srcData,
+                                            &beta,
+                                            dstTensorDesc,
+                                            dstData) );   
+    } 
     return dstData;
 }
 

src/Conv2d.cpp

@@ -4,17 +4,19 @@
 
 namespace tkDNN {
 
-Conv2d::Conv2d( Network *net, dataDim_t in_dim, int out_ch, 
-                int kernelH, int kernelW, int strideH, int strideW,
-                const char* fname_weights, const char* fname_bias) : 
+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) : 
     
-    LayerWgs(net, in_dim, in_dim.c, out_ch, kernelH, kernelW, 1, 
-             fname_weights, fname_bias) {
+    LayerWgs(net, net->getOutputDim().c, out_ch, kernelH, kernelW, 1, 
+             fname_weights, batchnorm) {
 
     this->kernelH = kernelH;
     this->kernelW = kernelW;
     this->strideH = strideH;
     this->strideW = strideW;
+    this->paddingH = paddingH;
+    this->paddingW = paddingW;
 
     checkCUDNN( cudnnCreateFilterDescriptor(&filterDesc) );
     checkCUDNN( cudnnCreateConvolutionDescriptor(&convDesc) );
@@ -33,10 +35,10 @@ Conv2d::Conv2d( Network *net, dataDim_t in_dim, int out_ch,
                 kernelH, kernelW) );
 
     checkCUDNN( cudnnSetConvolution2dDescriptor(convDesc,
-                0,0, // padding
+                paddingH, paddingW, // padding
                 strideH, strideW, // stride
                 1,1, // upscale
-                CUDNN_CROSS_CORRELATION) );
+                CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT) );
 
     // find dimension of convolution output
     checkCUDNN( cudnnGetConvolution2dForwardOutputDim(
@@ -74,7 +76,7 @@ Conv2d::Conv2d( Network *net, dataDim_t in_dim, int out_ch,
     output_dim.l = 1;
 
     //allocate data for infer result
-    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) );
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
 }
 
 Conv2d::~Conv2d() {
@@ -89,24 +91,34 @@ Conv2d::~Conv2d() {
     checkCuda( cudaFree(dstData) );
 }
 
-value_type* Conv2d::infer(dataDim_t &dim, value_type* srcData) {
+dnnType* Conv2d::infer(dataDim_t &dim, dnnType* srcData) {
 
 
     // convolution
-    value_type alpha = value_type(1);
-    value_type beta  = value_type(0);
+    dnnType alpha = dnnType(1);
+    dnnType beta  = dnnType(0);
     checkCUDNN( cudnnConvolutionForward(net->cudnnHandle,
                 &alpha, srcTensorDesc, srcData, filterDesc,
                 data_d, convDesc, algo, workSpace, ws_sizeInBytes,
                 &beta, dstTensorDesc, dstData) );
 
-    // bias
-    alpha = value_type(1);
-    beta  = value_type(1);
-    checkCUDNN( cudnnAddTensor(net->cudnnHandle,
-                &alpha, biasTensorDesc, bias_d,
-                &beta, dstTensorDesc, dstData) );
-
+    if(!batchnorm) {
+        // bias
+        alpha = dnnType(1);
+        beta  = dnnType(1);
+        checkCUDNN( cudnnAddTensor(net->cudnnHandle,
+                    &alpha, biasTensorDesc, bias_d,
+                    &beta, dstTensorDesc, dstData) );
+    } else {
+        float one = 1;
+        float zero = 0;
+        cudnnBatchNormalizationForwardInference(net->cudnnHandle,
+            CUDNN_BATCHNORM_SPATIAL, &one, &zero, 
+            dstTensorDesc, dstData, dstTensorDesc, 
+            dstData, biasTensorDesc, //same tensor descriptor as bias 
+            scales_d, bias_d, mean_d, variance_d, 
+            CUDNN_BN_MIN_EPSILON);
+    }
     //update data dimensions    
     dim = output_dim;
 

src/Dense.cpp

@@ -4,9 +4,8 @@
 
 namespace tkDNN {
 
-Dense::Dense(Network *net, dataDim_t in_dim, 
-    int out_ch, const char* fname_weights, const char* fname_bias) : 
-    LayerWgs(net, in_dim, in_dim.tot(), out_ch, 1, 1, 1, fname_weights, fname_bias) {
+Dense::Dense(Network *net, int out_ch, const char* fname_weights) : 
+    LayerWgs(net, net->getOutputDim().tot(), out_ch, 1, 1, 1, fname_weights) {
 
     output_dim.n = 1;
     output_dim.c = out_ch;
@@ -15,7 +14,7 @@ Dense::Dense(Network *net, dataDim_t in_dim,
     output_dim.l = 1;
 
     //allocate data for infer result
-    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) );
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
 }
 
 Dense::~Dense() {
@@ -23,7 +22,7 @@ Dense::~Dense() {
     checkCuda( cudaFree(dstData) );
 }
 
-value_type* Dense::infer(dataDim_t &dim, value_type* srcData) {
+dnnType* Dense::infer(dataDim_t &dim, dnnType* srcData) {
 
     if (dim.n != 1)
         FatalError("Not Implemented"); 
@@ -34,9 +33,9 @@ value_type* Dense::infer(dataDim_t &dim, value_type* srcData) {
     if (dim_x != input_dim.tot())
         FatalError("Input mismatch");
 
-    value_type alpha = value_type(1), beta = value_type(1);
+    dnnType alpha = dnnType(1), beta = dnnType(1);
     // place bias into dstData
-    checkCuda( cudaMemcpy(dstData, bias_d, dim_y*sizeof(value_type), cudaMemcpyDeviceToDevice) );
+    checkCuda( cudaMemcpy(dstData, bias_d, dim_y*sizeof(dnnType), cudaMemcpyDeviceToDevice) );
     
     //do matrix moltiplication
     checkERROR( cublasSgemv(net->cublasHandle, CUBLAS_OP_T,

src/Flatten.cpp

@@ -5,10 +5,9 @@
 
 namespace tkDNN {
 
-Flatten::Flatten(Network *net, dataDim_t input_dim) : 
-    Layer(net, input_dim) {
+Flatten::Flatten(Network *net) : Layer(net) {
 
-    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
+    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
 
     output_dim.n = 1;
     output_dim.c = input_dim.tot();
@@ -23,7 +22,7 @@ Flatten::~Flatten() {
     checkCuda( cudaFree(dstData) );
 }
 
-value_type* Flatten::infer(dataDim_t &dim, value_type* srcData) {
+dnnType* Flatten::infer(dataDim_t &dim, dnnType* srcData) {
 
     //transpose per channel
     matrixTranspose(net->cublasHandle, srcData, dstData, dim.c, dim.h*dim.w*dim.l);

src/Layer.cpp

@@ -4,11 +4,11 @@
 
 namespace tkDNN {
 
-Layer::Layer(Network *net, dataDim_t in_dim) {
+Layer::Layer(Network *net) {
 
     this->net = net;
-    this->input_dim = in_dim;
-    this->output_dim = in_dim;
+    this->input_dim = net->getOutputDim();
+    this->output_dim = input_dim;
     
     checkCUDNN( cudnnCreateTensorDescriptor(&srcTensorDesc) );
     checkCUDNN( cudnnCreateTensorDescriptor(&dstTensorDesc) );

src/LayerWgs.cpp

@@ -1,21 +1,100 @@
 #include <iostream>
+#include <string.h>
 
 #include "Layer.h"
+#include "kernels.h"
 
 namespace tkDNN {
 
-LayerWgs::LayerWgs(Network *net, dataDim_t in_dim, 
-                int inputs, int outputs, int kh, int kw, int kl, 
-                const char* fname_weights, const char* fname_bias) : Layer(net, in_dim) {
+LayerWgs::LayerWgs(Network *net, int inputs, int outputs, 
+                   int kh, int kw, int kl, 
+                   const char* fname_weights, bool batchnorm) : Layer(net) {
 
     this->inputs  = inputs;
     this->outputs = outputs;    
     this->weights_path  = std::string(fname_weights);
-    this->bias_path     = std::string(fname_bias);
-
+   
     std::cout<<"Reading weights: I="<<inputs<<" O="<<outputs<<" KERNEL="<<kh<<"x"<<kw<<"x"<<kl<<"\n";
-    readBinaryFile(weights_path.c_str(), inputs*outputs*kh*kw*kl, &data_h, &data_d);
-    readBinaryFile(bias_path.c_str(), outputs, &bias_h, &bias_d);
+    int seek = 0;
+    readBinaryFile(weights_path.c_str(), inputs*outputs*kh*kw*kl, &data_h, &data_d, seek);
+    seek += inputs*outputs*kh*kw*kl;
+    readBinaryFile(weights_path.c_str(), outputs, &bias_h, &bias_d, seek);
+    
+    this->batchnorm = batchnorm;
+    if(batchnorm) {
+        seek += outputs;
+        readBinaryFile(weights_path.c_str(), outputs, &scales_h, &scales_d, seek);
+        seek += outputs;
+        readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek);
+        seek += outputs;
+        readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek);
+
+        float eps = CUDNN_BN_MIN_EPSILON;
+
+        power_h = new dnnType[outputs];
+        for(int i=0; i<outputs; i++) power_h[i] = 1.0f;
+
+        for(int i=0; i<outputs; i++)
+            mean_h[i] = mean_h[i] / -sqrt(eps + variance_h[i]); 
+
+        for(int i=0; i<outputs; i++)
+            variance_h[i] = 1.0f / sqrt(eps + variance_h[i]);
+    }
+
+
+    if(!net->fp16)
+        return;
+
+    //convert to fp16
+    int w_size = inputs*outputs*kh*kw*kl;
+    data16_h = new __half[w_size];
+    cudaMalloc(&data16_d, w_size*sizeof(__half));
+    float2half(data_d, data16_d, w_size);
+    cudaMemcpy(data16_h, data16_d, w_size*sizeof(__half), cudaMemcpyDeviceToHost);
+
+    int b_size = outputs;
+    bias16_h = new __half[b_size];
+    cudaMalloc(&bias16_d, w_size*sizeof(__half));
+    float2half(bias_d, bias16_d, b_size);
+    cudaMemcpy(bias16_h, bias16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
+
+    if(batchnorm) {
+
+        power16_h    = new __half[b_size];  
+        mean16_h     = new __half[b_size];
+        variance16_h = new __half[b_size];
+        scales16_h   = new __half[b_size];
+
+        cudaMalloc(&power16_d, b_size*sizeof(__half));
+        cudaMalloc(&mean16_d, b_size*sizeof(__half));
+        cudaMalloc(&variance16_d, b_size*sizeof(__half));
+        cudaMalloc(&scales16_d, b_size*sizeof(__half));
+
+        //temporary buffers
+        float *tmp_d;
+        cudaMalloc(&tmp_d, b_size*sizeof(float));
+
+        //init power array of ones
+        cudaMemcpy(tmp_d, power_h, b_size*sizeof(float), cudaMemcpyHostToDevice);
+        float2half(tmp_d, power16_d, b_size);
+        cudaMemcpy(power16_h, power16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
+
+        //mean array
+
+        cudaMemcpy(tmp_d, mean_h, b_size*sizeof(float), cudaMemcpyHostToDevice);
+        float2half(tmp_d, mean16_d, b_size);
+        cudaMemcpy(mean16_h, mean16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
+
+        //convert variance
+ 
+        cudaMemcpy(tmp_d, variance_h, b_size*sizeof(float), cudaMemcpyHostToDevice);
+        float2half(tmp_d, variance16_d, b_size);
+        cudaMemcpy(variance16_h, variance16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
+
+        //conver scales
+        float2half(scales_d, scales16_d, b_size);
+        cudaMemcpy(scales16_h, scales16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost);
+    }
 }
 
 LayerWgs::~LayerWgs() {
@@ -24,6 +103,15 @@ LayerWgs::~LayerWgs() {
     delete [] bias_h;
     checkCuda( cudaFree(data_d) );
     checkCuda( cudaFree(bias_d) );
+
+    if(batchnorm) {
+        delete [] scales_h;
+        delete [] mean_h;
+        delete [] variance_h;
+        checkCuda( cudaFree(scales_d) );
+        checkCuda( cudaFree(mean_d) );
+        checkCuda( cudaFree(variance_d) );
+    }
 }
 
-}
+}

src/MulAdd.cpp

@@ -5,8 +5,7 @@
 
 namespace tkDNN {
 
-MulAdd::MulAdd(Network *net, dataDim_t input_dim, value_type mul, value_type add) : 
-    Layer(net, input_dim) {
+MulAdd::MulAdd(Network *net, dnnType mul, dnnType add) : Layer(net) {
 
     this->mul = mul;
     this->add = add;
@@ -14,16 +13,16 @@ MulAdd::MulAdd(Network *net, dataDim_t input_dim, value_type mul, value_type add
     int size = input_dim.tot();
 
     // create a vector with all value setted to add 
-    value_type *add_vector_h = new value_type[size];
+    dnnType *add_vector_h = new dnnType[size];
     for(int i=0; i<size; i++)
         add_vector_h[i] = add;
 
-    checkCuda( cudaMalloc(&add_vector, size*sizeof(value_type)));
-    checkCuda( cudaMemcpy(add_vector, add_vector_h, size*sizeof(value_type), cudaMemcpyHostToDevice));
+    checkCuda( cudaMalloc(&add_vector, size*sizeof(dnnType)));
+    checkCuda( cudaMemcpy(add_vector, add_vector_h, size*sizeof(dnnType), cudaMemcpyHostToDevice));
     delete [] add_vector_h;
 
 
-    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
+    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
 }
 
 MulAdd::~MulAdd() {
@@ -32,7 +31,7 @@ MulAdd::~MulAdd() {
     checkCuda( cudaFree(dstData) );
 }
 
-value_type* MulAdd::infer(dataDim_t &dim, value_type* srcData) {
+dnnType* MulAdd::infer(dataDim_t &dim, dnnType* srcData) {
 
     matrixMulAdd(net->cublasHandle, srcData, dstData, add_vector, input_dim.tot(), mul);
     

src/Network.cpp

@@ -1,4 +1,5 @@
 #include <iostream>
+#include <string.h>
 
 #include "tkdnn.h"
 #include "Network.h"
@@ -6,12 +7,14 @@
 
 namespace tkDNN {
 
-Network::Network() {
+Network::Network(dataDim_t input_dim) {
+    this->input_dim = input_dim;
 
-    float tk_ver = float(tkDNN::getVersion())/1000;
+    float tk_ver = float(TKDNN_VERSION)/1000;
     float cu_ver = float(cudnnGetVersion())/1000;
 
-    std::cout<<"New NETWORK (tkDNN v"<<tk_ver<<", CUDNN v"<<cu_ver<<")\n";
+    std::cout<<"New NETWORK (tkDNN v"<<tk_ver
+             <<", CUDNN v"<<cu_ver<<")\n";
     dataType = CUDNN_DATA_FLOAT;
     tensorFormat = CUDNN_TENSOR_NCHW;
 
@@ -19,6 +22,14 @@ Network::Network() {
     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() {
@@ -27,12 +38,13 @@ Network::~Network() {
     checkERROR( cublasDestroy(cublasHandle) );
 }
 
-value_type* Network::infer(dataDim_t &dim, value_type* data) {
+dnnType* Network::infer(dataDim_t &dim, dnnType* data) {
 
     //do infer for every layer
-    for(int i=0; i<num_layers; i++)
+    for(int i=0; i<num_layers; i++) {
         data = layers[i]->infer(dim, data);
-
+    }
+    checkCuda(cudaDeviceSynchronize());
     return data;
 }
 
@@ -44,4 +56,47 @@ bool Network::addLayer(Layer *l) {
     return true;
 }
 
-}
+dataDim_t Network::getOutputDim() {
+
+        if(num_layers == 0)
+            return input_dim;
+        else
+            return layers[num_layers-1]->output_dim;
+}
+
+void Network::print() {
+
+    printCenteredTitle(" NETWORK MODEL ", '=', 60);
+    std::cout.width(3); std::cout<<std::left<<"N.";
+    std::cout<<" ";
+    std::cout.width(17); std::cout<<std::left<<"Layer type";
+    std::cout.width(22); std::cout<<std::left<<"input (H*W,CH)";
+    std::cout.width(16); std::cout<<std::left<<"output (H*W,CH)";
+    std::cout<<"\n";
+
+    for(int i=0; i<num_layers; i++) {
+        dataDim_t in = layers[i]->input_dim;
+        dataDim_t out = layers[i]->output_dim;
+
+        std::cout.width(3); std::cout<<std::right<<i;
+        std::cout<<" ";
+        std::cout.width(16); std::cout<<std::left<<layers[i]->getLayerName();
+        std::cout.width(4);  std::cout<<std::right<<in.h;
+        std::cout<<" x ";
+        std::cout.width(4);  std::cout<<std::right<<in.w;
+        std::cout<<", ";
+        std::cout.width(4);  std::cout<<std::right<<in.c;
+        std::cout<<"  -> ";
+        std::cout.width(4);  std::cout<<std::right<<out.h;
+        std::cout<<" x ";
+        std::cout.width(4);  std::cout<<std::right<<out.w;
+        std::cout<<", ";
+        std::cout.width(4);  std::cout<<std::right<<out.c;
+        std::cout<<"\n";
+    }
+    printCenteredTitle("", '=', 60);
+    std::cout<<"\n";
+}
+
+
+}

src/NetworkRT.cpp

@@ -0,0 +1,402 @@
+#include <iostream>
+#include <map>
+#include <errno.h>
+#include <string.h> // memcpy
+#include <stdlib.h>
+
+#include "kernels.h"
+
+#include "utils.h"
+#include "NvInfer.h"
+#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 {
+    void log(Severity severity, const char* msg) override {
+#ifdef DEBUG
+        std::cout <<"TENSORRT LOG: "<< msg << std::endl;
+#endif
+    }
+} loggerRT;
+
+namespace tkDNN {
+
+std::map<Layer*, nvinfer1::ITensor*>tensors; 
+
+NetworkRT::NetworkRT(Network *net, const char *name) {
+
+    float rt_ver = float(NV_TENSORRT_MAJOR) + 
+                   float(NV_TENSORRT_MINOR)/10 + 
+                   float(NV_TENSORRT_PATCH)/100;
+    std::cout<<"New NetworkRT (TensorRT v"<<rt_ver<<")\n";
+  
+    builderRT = createInferBuilder(loggerRT);
+    std::cout<<"Float16 support: "<<builderRT->platformHasFastFp16()<<"\n";
+	std::cout<<"Int8 support: "<<builderRT->platformHasFastInt8()<<"\n";
+    networkRT = builderRT->createNetwork();
+
+    if(!fileExist(name)) {
+
+        //input and dataType
+        dataDim_t dim = net->layers[0]->input_dim;
+        dtRT = DataType::kFLOAT;
+
+        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()) {
+            dtRT = DataType::kHALF;
+            builderRT->setHalf2Mode(true);
+        }
+
+        //add input layer
+        ITensor *input = networkRT->addInput("data", DataType::kFLOAT, 
+                        DimsCHW{ dim.c, dim.h, dim.w});
+        checkNULL(input);
+
+        //add other layers
+        for(int i=0; i<net->num_layers; i++) {
+            Layer *l = net->layers[i];
+            ILayer *Ilay = convert_layer(input, l);
+            Ilay->setName( (l->getLayerName() + std::to_string(i)).c_str() );
+            
+            input = Ilay->getOutput(0);
+            tensors[l] = input;
+        }
+        if(input == NULL)
+            FatalError("conversion failed");
+
+        //build tensorRT
+        input->setName("out");
+        networkRT->markOutput(*input);
+
+        std::cout<<"Building tensorRT cuda engine...\n";
+        engineRT = builderRT->buildCudaEngine(*networkRT);
+        // we don't need the network any more
+        //networkRT->destroy();
+        serialize(name);
+    } else {
+        deserialize(name);
+    }
+
+    std::cout<<"create execution context\n";
+	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");
+
+	// 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()
+	buf_input_idx = engineRT->getBindingIndex("data"); 
+    buf_output_idx = engineRT->getBindingIndex("out");
+    std::cout<<"input idex = "<<buf_input_idx<<" -> output index = "<<buf_output_idx<<"\n";
+
+
+    Dims iDim = engineRT->getBindingDimensions(buf_input_idx);
+    input_dim.n = 1;
+    input_dim.c = iDim.d[0];
+    input_dim.h = iDim.d[1];
+    input_dim.w = iDim.d[2];
+    input_dim.print();
+
+    Dims oDim = engineRT->getBindingDimensions(buf_output_idx);
+    output_dim.n = 1;
+    output_dim.c = oDim.d[0];
+    output_dim.h = oDim.d[1];
+    output_dim.w = oDim.d[2];
+	
+    // 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)));
+    checkCuda(cudaMalloc(&output, output_dim.tot()*sizeof(dnnType)));
+	checkCuda(cudaStreamCreate(&stream));
+}
+
+NetworkRT::~NetworkRT() {
+
+}
+
+dnnType* NetworkRT::infer(dataDim_t &dim, dnnType* data) {
+
+    checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, input_dim.tot()*sizeof(float), cudaMemcpyDeviceToDevice, stream));
+    contextRT->enqueue(1, buffersRT, stream, nullptr);
+    checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], output_dim.tot()*sizeof(float), cudaMemcpyDeviceToDevice, stream));
+    cudaStreamSynchronize(stream);
+
+    dim = output_dim;
+
+    return output;
+}
+
+void NetworkRT::enqueue() {
+    contextRT->enqueue(1, buffersRT, stream, nullptr);
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) {
+
+    layerType_t type = l->getLayerType();
+
+    if(type == LAYER_DENSE)
+        return convert_layer(input, (Dense*) l);
+    if(type == LAYER_CONV2D)
+        return convert_layer(input, (Conv2d*) l);
+    if(type == LAYER_POOLING)
+        return convert_layer(input, (Pooling*) l);
+    if(type == LAYER_ACTIVATION)
+        return convert_layer(input, (Activation*) l);
+    if(type == LAYER_SOFTMAX)
+        return convert_layer(input, (Softmax*) l);
+    if(type == LAYER_ROUTE)
+        return convert_layer(input, (Route*) l);
+    if(type == LAYER_REORG)
+        return convert_layer(input, (Reorg*) l);
+    if(type == LAYER_REGION)
+        return convert_layer(input, (Region*) l);
+
+    FatalError("Layer not implemented in tensorRT");
+    return NULL;
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Dense *l) {
+    //std::cout<<"convert Dense\n";
+    void *data_b, *bias_b;
+    if(dtRT == DataType::kHALF) {
+        data_b     = l->data16_h;    
+        bias_b     = l->bias16_h;
+    } else {
+        data_b     = l->data_h;    
+        bias_b     = l->bias_h;
+    }
+
+    Weights w { dtRT, data_b, l->inputs*l->outputs};
+    Weights b = { dtRT, bias_b, l->outputs};
+    IFullyConnectedLayer *lRT = networkRT->addFullyConnected(*input, l->outputs, w, b);
+
+    checkNULL(lRT);
+    return lRT;
+}
+
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
+    //std::cout<<"convert conv2D\n";
+
+    void *data_b, *bias_b, *power_b, *mean_b, *variance_b, *scales_b;
+    if(dtRT == DataType::kHALF) {
+        data_b     = l->data16_h;    
+        bias_b     = l->bias16_h;
+        power_b    = l->power16_h;
+        mean_b     = l->mean16_h;
+        variance_b = l->variance16_h;
+        scales_b   = l->scales16_h;
+    } else {
+        data_b     = l->data_h;    
+        bias_b     = l->bias_h;
+        power_b    = l->power_h;
+        mean_b     = l->mean_h;
+        variance_b = l->variance_h;
+        scales_b   = l->scales_h;
+    }
+
+
+    Weights w { dtRT, data_b, l->inputs*l->outputs*l->kernelH*l->kernelW};
+    Weights b;
+    if(!l->batchnorm)
+        b = { dtRT, bias_b, l->outputs};
+    else
+        b = { dtRT, nullptr, 0}; //on batchnorm bias are added later
+
+    // Add a convolution layer with 20 outputs and a 5x5 filter.
+    IConvolutionLayer *lRT = networkRT->addConvolution(*input, 
+               l->outputs, DimsHW{l->kernelH, l->kernelW}, w, b);
+    checkNULL(lRT);
+
+    lRT->setStride(DimsHW{l->strideH, l->strideW});
+    lRT->setPadding(DimsHW{l->paddingH, l->paddingW});
+
+    if(l->batchnorm) {
+        Weights power{dtRT, power_b, l->outputs};
+        Weights shift{dtRT, mean_b, l->outputs};
+        Weights scale{dtRT, variance_b, l->outputs};
+        IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, 
+                    shift, scale, power);
+        checkNULL(lRT2);
+
+        Weights shift2{dtRT, bias_b, l->outputs};
+        Weights scale2{dtRT, scales_b, l->outputs};
+        IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, 
+                    shift2, scale2, power);
+        checkNULL(lRT3);
+
+        return lRT3;
+    }
+
+    return lRT;
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) {
+    //std::cout<<"convert Pooling\n";
+
+    IPoolingLayer *lRT = networkRT->addPooling(*input, 
+        PoolingType::kMAX, DimsHW{l->winH, l->winW});
+    checkNULL(lRT);
+    lRT->setStride(DimsHW{l->strideH, l->strideW});
+
+    return lRT;
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Activation *l) {
+    //std::cout<<"convert Activation\n";
+
+    if(l->act_mode == ACTIVATION_LEAKY) {
+        //std::cout<<"New plugin LEAKY\n";
+        IPlugin *plugin = new ActivationLeakyRT();
+        IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+        checkNULL(lRT);
+        return lRT;
+
+    } else if(l->act_mode == CUDNN_ACTIVATION_RELU) {
+        IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU);
+        checkNULL(lRT);
+        return lRT;
+    
+    } else {
+        FatalError("this Activation mode is not yet implemented");
+        return NULL;
+    }
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Softmax *l) {
+    //std::cout<<"convert softmax\n";
+
+    ISoftMaxLayer *lRT = networkRT->addSoftMax(*input);
+    checkNULL(lRT);
+
+    return lRT;
+}
+
+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++)
+        tens[i] = tensors[l->layers[i]];
+    IConcatenationLayer *lRT = networkRT->addConcatenation(tens, l->layers_n);
+    checkNULL(lRT);
+
+    return lRT;
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Reorg *l) {
+    //std::cout<<"convert Reorg\n";
+
+    //std::cout<<"New plugin REORG\n";
+    IPlugin *plugin = new ReorgRT(l->stride);
+    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    checkNULL(lRT);
+    return lRT;
+}
+
+ILayer* NetworkRT::convert_layer(ITensor *input, Region *l) {
+    //std::cout<<"convert Region\n";
+
+    //std::cout<<"New plugin REGION\n";
+    IPlugin *plugin = new RegionRT(l->classes, l->coords, l->num);
+    IPluginLayer *lRT = networkRT->addPlugin(&input, 1, *plugin);
+    checkNULL(lRT);
+    return lRT;
+}
+
+bool NetworkRT::serialize(const char *filename) {
+
+    std::ofstream p(filename);
+    if (!p) {
+        FatalError("could not open plan output file");
+        return false;
+    }
+
+    IHostMemory *ptr = engineRT->serialize();
+    if(ptr == nullptr)
+        FatalError("Cant serialize network");
+
+    p.write(reinterpret_cast<const char*>(ptr->data()), ptr->size());
+    ptr->destroy();
+    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};
+    size_t size{0};
+    std::ifstream file(filename, std::ios::binary);
+    if (file.good()) {
+        file.seekg(0, file.end);
+        size = file.tellg();
+        file.seekg(0, file.beg);
+        gieModelStream = new char[size];
+        file.read(gieModelStream, size);
+        file.close();
+    }
+
+    PluginFactory plfact;
+
+    runtimeRT = createInferRuntime(loggerRT);
+    engineRT = runtimeRT->deserializeCudaEngine(gieModelStream, size, (IPluginFactory *) &plfact);
+    //if (gieModelStream) delete [] gieModelStream;
+
+    return true;
+}
+
+}

src/Pooling.cpp

@@ -5,20 +5,18 @@
 
 namespace tkDNN {
 
-Pooling::Pooling(   Network *net, dataDim_t input_dim, 
-                    int winH, int winW, int strideH, int strideW, tkdnnPoolingMode_t pool_mode) : 
-    Layer(net, input_dim) {
-
-
-    if(winH != strideH || winW != strideW)
-        FatalError("stride pooling not yet implemented");
+Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW,
+                  tkdnnPoolingMode_t pool_mode) : 
+    Layer(net) {
 
     this->winH = winH;
     this->winW = winW;
     this->strideH = strideH;
     this->strideW = strideW;
     this->pool_mode = pool_mode;
-    
+    this->paddingH = 0;
+    this->paddingW = 0;
+
     checkCUDNN( cudnnCreatePoolingDescriptor(&poolingDesc) );
 
     int n = input_dim.n;
@@ -46,27 +44,28 @@ Pooling::Pooling(   Network *net, dataDim_t input_dim,
                 net->tensorFormat, net->dataType, n, c, h, w) );
 
     //get out dim
-    h = h / winH; w = w / winW;
-    
+    checkCUDNN( cudnnGetPooling2dForwardOutputDim(poolingDesc, srcTensorDesc, &n, &c, &h, &w)); 
+    //h = (h + winH*this->paddingH)/strideH;
+    //w = (w + winW*this->paddingW)/strideW;
+
     checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
                 net->tensorFormat, net->dataType, n, c, h, w) );
        
-
     output_dim.n = n;
     output_dim.c = c;
     output_dim.h = h;
     output_dim.w = w;
     output_dim.l = l;
 
-    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(value_type)) );
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
 
     //pool on 3d data need transposition at the enter and on the exit
     //allocate for initial and final transposition
     if(poolOn3d) {
         output_dim.n = 1;
 
-        checkCuda( cudaMalloc(&tmpInputData, input_dim.tot()*sizeof(value_type)) );
-        checkCuda( cudaMalloc(&tmpOutputData, output_dim.tot()*sizeof(value_type)) );
+        checkCuda( cudaMalloc(&tmpInputData, input_dim.tot()*sizeof(dnnType)) );
+        checkCuda( cudaMalloc(&tmpOutputData, output_dim.tot()*sizeof(dnnType)) );
     }
 
 }
@@ -82,10 +81,10 @@ Pooling::~Pooling() {
     checkCuda( cudaFree(dstData) );
 }
 
-value_type* Pooling::infer(dataDim_t &dim, value_type* srcData) {
+dnnType* Pooling::infer(dataDim_t &dim, dnnType* srcData) {
 
-    value_type *poolSrc = srcData;
-    value_type *poolDst = dstData;
+    dnnType *poolSrc = srcData;
+    dnnType *poolDst = dstData;
 
     if(poolOn3d) {
         matrixTranspose(net->cublasHandle, srcData, tmpInputData, dim.h*dim.w*dim.c, dim.l);
@@ -93,8 +92,8 @@ value_type* Pooling::infer(dataDim_t &dim, value_type* srcData) {
         poolDst = tmpOutputData;
     }
 
-    value_type alpha = value_type(1);
-    value_type beta = value_type(0);
+    dnnType alpha = dnnType(1);
+    dnnType beta = dnnType(0);
     checkCUDNN( cudnnPoolingForward(net->cudnnHandle, poolingDesc,
                 &alpha, srcTensorDesc, poolSrc,
                 &beta, dstTensorDesc, poolDst) );

src/Region.cpp

@@ -0,0 +1,341 @@
+#include <iostream>
+
+#ifdef OPENCV
+    #include <opencv2/core/core.hpp>
+    #include <opencv2/highgui/highgui.hpp>
+#endif
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tkDNN {
+
+Region::Region(Network *net, int classes, int coords, int num) : 
+    Layer(net) {
+
+    this->classes = classes;
+    this->coords = coords;
+    this->num = num;
+
+    // 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, input_dim.tot()*sizeof(dnnType)) );
+}
+
+Region::~Region() {
+    checkCuda( cudaFree(dstData) );
+}
+
+int entry_index(int batch, int location, int entry, 
+            int coords, 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*(coords+classes+1) +
+           entry*input_dim.w*input_dim.h + loc;
+}
+
+dnnType* Region::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, coords, classes, input_dim, output_dim);
+            activationLOGISTICForward(srcData + index, dstData + index, 2*dim.w*dim.h);
+            
+            index = entry_index(b, n*dim.w*dim.h, coords, coords, classes, input_dim, output_dim);
+            activationLOGISTICForward(srcData + index, dstData + index, dim.w*dim.h);
+        }
+    }
+
+    //softmax start
+    int index = entry_index(0, 0, coords + 1, coords, classes, input_dim, output_dim);
+    softmaxForward(srcData + index, classes, output_dim.n*num, output_dim.tot()/num, 
+                   output_dim.w*output_dim.h, 1, output_dim.w*output_dim.h, 1, dstData + index);
+
+    dim = output_dim;
+    return dstData;
+}
+
+
+/* Intepret class */
+RegionInterpret::RegionInterpret(dataDim_t input_dim, dataDim_t output_dim, 
+    int classes, int coords, int num, float thresh, const char* fname_weights) {
+
+    this->input_dim = input_dim;
+    this->output_dim = output_dim;
+
+    this->classes = classes;
+    this->coords = coords;
+    this->num = num;
+    this->thresh = thresh;
+    this->res_boxes_n = 0;
+
+    int tot = output_dim.w*output_dim.h*num;
+    boxes = (box*)    malloc(tot*sizeof(box));
+    probs = (float**) malloc(tot*sizeof(float *));
+    for(int j = 0; j < tot; ++j) probs[j] = (float*) malloc((classes + 1)*sizeof(float *));
+    s = (sortable_bbox*) malloc(tot*sizeof(sortable_bbox));
+
+    //load anchors
+    readBinaryFile(fname_weights, 2*num, &bias_h, &bias_d);
+}
+
+RegionInterpret::~RegionInterpret() {
+
+    delete [] boxes;
+    for(int j = 0; j < output_dim.w*output_dim.h*num; ++j) 
+        delete [] probs[j];
+    delete [] probs;
+    delete [] s;
+
+    delete [] bias_h;
+    checkCuda( cudaFree(bias_d) );
+}
+
+box RegionInterpret::get_region_box(float *x, float *biases, int n, int index, int i, int j, int w, int h, int stride)
+{
+    box b;
+    b.x = (i + x[index + 0*stride]) / w;
+    b.y = (j + x[index + 1*stride]) / h;
+    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;
+}
+
+void RegionInterpret::get_region_boxes(  float *input, int w, int h, int netw, int neth, float thresh, 
+                                float **probs, box *boxes, int only_objectness, 
+                                int *map, float tree_thresh, int relative) {
+    int lh = output_dim.h;
+    int lw = output_dim.w;
+    float *predictions = input;
+    for (int i = 0; i < lw*lh; ++i){
+
+        int row = i / lw;
+        int col = i % lw;
+        for(int n = 0; n < num; ++n){
+
+            int index = n*lw*lh + i;
+            for(int j = 0; j < classes; ++j){
+                probs[index][j] = 0;
+            }
+            int obj_index  = entry_index(0, n*lw*lh + i, 
+                                         coords, coords, classes, output_dim, output_dim);
+            int box_index  = entry_index(0, n*lw*lh + i, 0, 
+                                         coords, classes, output_dim, output_dim);
+            float scale = predictions[obj_index];
+            boxes[index] = get_region_box(predictions, bias_h, n, box_index, col, row, lw, lh, lw*lh);
+
+            float max = 0;
+            for(int j = 0; j < classes; ++j){
+                int class_index = entry_index(0, n*lw*lh + i, coords + 1 + j,
+                                              coords, classes, output_dim, output_dim);
+                float prob = scale*predictions[class_index];
+                probs[index][j] = (prob > thresh) ? prob : 0;
+                if(prob > max) max = prob;
+            }
+            probs[index][classes] = max;
+        }
+    }
+    correct_region_boxes(boxes, lw*lh*num, w, h, netw, neth, relative);
+}
+
+
+void RegionInterpret::correct_region_boxes(box *boxes, 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){
+        box b = boxes[i];
+        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;
+        }
+        boxes[i] = b;
+    }
+}
+
+
+
+
+//############################ BOX PROBABILITY UTILS ############################
+int nms_comparator(const void *pa, const void *pb) {
+    sortable_bbox a = *(sortable_bbox *)pa;
+    sortable_bbox b = *(sortable_bbox *)pb;
+    float diff = a.probs[a.index][b.cl] - b.probs[b.index][b.cl];
+    if(diff < 0) return 1;
+    else if(diff > 0) return -1;
+    return 0;
+}
+float overlap(float x1, float w1, float x2, float w2) {
+    /*
+    //SLOW METHOD
+    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;
+    */
+
+    //SPALLA METHOD
+    float l;
+    w1 < w2? l=w1 : l=w2;
+    float d = fabs(x1 - x2);
+    float k = fabs(w1 - w2)/2;
+    if      (d <= k)    return l;
+    else if (d <= k +l) return l - (d-k);
+    else                return 0;
+}
+float box_intersection(box a, box b) {
+    float w = overlap(a.x, a.w, b.x, b.w);
+    if(w <= 0) return 0;
+    float h = overlap(a.y, a.h, b.y, b.h);
+    if(h <= 0) return 0;
+    float area = w*h;
+    return area;
+}
+float box_union(box a, box b) {
+    float i = box_intersection(a, b);
+    float u = a.w*a.h + b.w*b.h - i;
+    return u;
+}
+int max_index(float *a, int n) {
+    if(n <= 0) return -1;
+    int i, max_i = 0;
+    float max = a[0];
+    for(i = 1; i < n; ++i){
+        if(a[i] > max){
+            max = a[i];
+            max_i = i;
+        }
+    }
+    return max_i;
+}
+//###############################################################################
+float RegionInterpret::box_iou(box a, box b) {
+    if(fabs(a.x - b.x) > (a.w+b.w)/2 || fabs(a.y - b.y) > (a.h+b.h)/2)
+        return 0;
+    return box_intersection(a, b)/box_union(a, b);
+}
+
+
+
+void RegionInterpret::interpretData(dnnType *data_h, int imageW, int imageH) {
+
+    int imW, imH;
+    if(imageW <= 0 || imageH <= 0) {
+        imW = input_dim.w;
+        imH = input_dim.h;
+    } else {
+        imW = imageW;
+        imH = imageH;
+    }
+
+    int tot = output_dim.w*output_dim.h*num;
+
+    get_region_boxes(data_h, imW, imH, output_dim.w, output_dim.h, thresh, probs, boxes, 0, 0, 0.5, 1);
+
+    //delete repeats
+    for(int i = 0; i < tot; ++i){
+        s[i].index = i;       
+        s[i].cl = classes;
+        s[i].probs = probs;
+    }
+    qsort(s, tot, sizeof(sortable_bbox), nms_comparator);
+
+    for(int i = 0; i < tot; ++i){
+        if(probs[s[i].index][classes] == 0) continue;
+        box a = boxes[s[i].index];
+        for(int j = i+1; j < tot; ++j){
+            box b = boxes[s[j].index];
+            if (box_iou(a, b) > 0.3f){
+                for(int k = 0; k < classes+1; ++k){
+                    probs[s[j].index][k] = 0;
+                }
+            }
+        }
+    }
+
+    res_boxes_n = 0;
+    //print results
+    for(int i = 0; i < tot; ++i){
+        int cl = max_index(probs[i], classes);
+        float prob = probs[i][cl];
+
+        if(prob > thresh) {
+            box b = boxes[i];
+            int x = (b.x)*imW;
+            int w = (b.w)*imW - b.x;
+            int y = (b.y)*imH;
+            int h = (b.h)*imH - b.y;
+
+            //if(x < 0) x = 0;
+            //if(y < 0) y = 0;
+            //if(w > imW) w = imW;
+            //if(h > imH) h = imH;
+
+            //printf("%d: %.0f%% box(x1, y1, x2, y2): %d %d %d %d\n", cl, prob*100, x, y, w, h);
+            b.x = x;
+            b.y = y;
+            b.h = h;
+            b.w = w;
+            b.cl = cl;
+            b.prob = prob;
+            res_boxes[res_boxes_n] = b;
+            res_boxes_n++;
+        }
+    }
+}
+
+void RegionInterpret::showImageResult(dnnType *input_h) {
+
+#ifdef OPENCV
+    dataDim_t dim = input_dim;
+    // read an image
+    cv::Mat r(dim.h, dim.w, CV_32F, input_h);
+    cv::Mat g(dim.h, dim.w, CV_32F, input_h + dim.h*dim.w);
+    cv::Mat b(dim.h, dim.w, CV_32F, input_h + dim.h*dim.w*2);
+    std::vector<cv::Mat> array_to_merge;
+    array_to_merge.push_back(b);
+    array_to_merge.push_back(g);
+    array_to_merge.push_back(r);
+    cv::Mat color;
+    cv::merge(array_to_merge, color);
+
+    for(int i=0; i<res_boxes_n; i++) {
+        box bx = res_boxes[i];
+        cv::rectangle(color, 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);
+    }
+    cv::namedWindow("result");
+    // show the image on window
+    cv::imshow("result", color);
+    // wait key for 5000 ms
+    cv::waitKey(0);
+#else
+    std::cout<<"Visualization not supported, please recompile with OpenCV\n";
+#endif
+}
+
+}

src/Reorg.cpp

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

src/Route.cpp

@@ -0,0 +1,56 @@
+#include <iostream>
+
+#include "Layer.h"
+#include "kernels.h"
+
+namespace tkDNN {
+
+Route::Route(Network *net, Layer **layers, int layers_n) : Layer(net) {
+
+    this->layers = layers;
+    this->layers_n = layers_n;
+
+    //get dims
+    output_dim.l = 1;
+    output_dim.c = 0;
+    for(int i=0; i<layers_n; i++) {
+
+        if(i==0) {
+            output_dim.w = layers[i]->output_dim.w;
+            output_dim.h = layers[i]->output_dim.h;
+        } else {
+            if( layers[i]->output_dim.w != output_dim.w ||
+                layers[i]->output_dim.h != output_dim.h   )
+                FatalError("Route Output dim missmatch");
+        }
+        output_dim.c += layers[i]->output_dim.c;
+    }
+
+    input_dim = output_dim;
+
+    checkCuda( cudaMalloc(&dstData, output_dim.tot()*sizeof(dnnType)) );
+}
+
+Route::~Route() {
+
+    checkCuda( cudaFree(dstData) );
+}
+
+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();
+        checkCuda( cudaMemcpy(dstData + offset, input, in_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
+        offset += in_dim;
+    }
+
+    //update data dimensions    
+    dim = output_dim;
+
+    return dstData;
+}
+
+}

src/Softmax.cpp

@@ -5,10 +5,9 @@
 
 namespace tkDNN {
 
-Softmax::Softmax(Network *net, dataDim_t input_dim) : 
-    Layer(net, input_dim) {
+Softmax::Softmax(Network *net) : Layer(net) {
 
-    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(value_type)) );
+    checkCuda( cudaMalloc(&dstData, input_dim.tot()*sizeof(dnnType)) );
 
     checkCUDNN( cudnnSetTensor4dDescriptor(srcTensorDesc,
                                         net->tensorFormat,
@@ -29,10 +28,10 @@ Softmax::~Softmax() {
     checkCuda( cudaFree(dstData) );
 }
 
-value_type* Softmax::infer(dataDim_t &dim, value_type* srcData) {
+dnnType* Softmax::infer(dataDim_t &dim, dnnType* srcData) {
 
-    value_type alpha = value_type(1);
-    value_type beta  = value_type(0);
+    dnnType alpha = dnnType(1);
+    dnnType beta  = dnnType(0);
     checkCUDNN( cudnnSoftmaxForward(net->cudnnHandle,
                                     CUDNN_SOFTMAX_ACCURATE ,
                                     CUDNN_SOFTMAX_MODE_CHANNEL,

src/kernels/activation_elu.cu

@@ -7,12 +7,12 @@
         x > 0 :    y = x
 */
 __global__
-void activation_elu(value_type *input, value_type *output, int size) {
+void activation_elu(dnnType *input, dnnType *output, int size) {
 
     int i = blockDim.x*blockIdx.x + threadIdx.x;
 
     if(i<size) {    
-        value_type k0, k1;
+        dnnType k0, k1;
         
         if (input[i]>0)
             k0 = 1.0f;
@@ -28,11 +28,10 @@ void activation_elu(value_type *input, value_type *output, int size) {
 /**
     ELU activation function
 */
-void activationELUForward(value_type* srcData, value_type* dstData, int size)
+void activationELUForward(dnnType* srcData, dnnType* dstData, int size, const cudaStream_t stream)
 {
     int blocks = (size+255)/256;
     int threads = 256;
     
-    activation_elu<<<blocks, threads>>>(srcData, dstData, size);
-    checkCuda( cudaDeviceSynchronize() );
-}
+    activation_elu<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+}

src/kernels/activation_leaky.cu

@@ -0,0 +1,28 @@
+#include "kernels.h"
+
+__global__
+void activation_leaky(dnnType *input, dnnType *output, int size) {
+
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+
+    if(i<size) {    
+        if (input[i]>0)
+            output[i] = input[i];
+        else
+            output[i] = 0.1f*input[i];
+    }
+ }
+
+
+/**
+    ELU activation function
+*/
+void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    activation_leaky<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+}
+
+

src/kernels/activation_logistic.cu

@@ -0,0 +1,25 @@
+#include "kernels.h"
+
+__global__
+void activation_logistic(dnnType *input, dnnType *output, int size) {
+
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+
+    if(i<size) {    
+        output[i] =  1.0f/(1.0f + exp(-input[i]));;
+    }
+ }
+
+
+/**
+    LOGISTIC activation function
+*/
+void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    activation_logistic<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+}
+
+

src/kernels/convert.cu

@@ -0,0 +1,21 @@
+#include "kernels.h"
+
+__global__
+void float2half_device(float *input, __half *output, int size) {
+
+    int i = blockDim.x*blockIdx.x + threadIdx.x;
+
+    if(i<size) {    
+        output[i] = __float2half(input[i]);
+    }
+}
+
+
+void float2half(float* srcData, __half *dstData, int size, const cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    float2half_device<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+    cudaDeviceSynchronize();
+}

src/kernels/reorg.cu

@@ -0,0 +1,49 @@
+#include "kernels.h"
+
+__global__ void reorg_kernel(int N, float *x, int w, int h, int c, int batch, int stride, int forward, float *out)
+{
+    int i = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if(i >= N) return;
+    int in_index = i;
+    int in_w = i%w;
+    i = i/w;
+    int in_h = i%h;
+    i = i/h;
+    int in_c = i%c;
+    i = i/c;
+    int b = i%batch;
+
+    int out_c = c/(stride*stride);
+
+    int c2 = in_c % out_c;
+    int offset = in_c / out_c;
+    int w2 = in_w*stride + offset % stride;
+    int h2 = in_h*stride + offset / stride;
+    //printf("%d\n", offset);
+    int out_index = w2 + w*stride*(h2 + h*stride*(c2 + out_c*b));
+
+   // printf("%d %d %d\n", w2, h2, c2);
+    //printf("%d %d\n", in_index, out_index);
+    //if(out_index >= N || out_index < 0) printf("bad bad bad \n");
+
+    if(forward) out[out_index] = x[in_index];
+    else out[in_index] = x[out_index];
+    //if(forward) out[1] = x[1];
+    //else out[0] = x[0];
+}
+
+/**
+    reorg function function
+*/
+void reorgForward(dnnType* srcData, dnnType* dstData, 
+                  int n, int c, int h, int w, int stride, cudaStream_t stream) {
+
+    int size = n*c*h*w;
+
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    reorg_kernel<<<blocks, threads, 0, stream>>>(size, srcData, w, h, c, n, stride, false, dstData);
+}
+
+

src/kernels/softmax.cu

@@ -0,0 +1,42 @@
+#include "kernels.h"
+
+__device__ void softmax_device(float *input, int n, float temp, int stride, float *output)
+{
+    int i;
+    float sum = 0;
+    float largest = -INFINITY;
+    for(i = 0; i < n; ++i){
+        int val = input[i*stride];
+        largest = (val>largest) ? val : largest;
+    }
+    for(i = 0; i < n; ++i){
+        float e = exp(input[i*stride]/temp - largest/temp);
+        sum += e;
+        output[i*stride] = e;
+    }
+    for(i = 0; i < n; ++i){
+        output[i*stride] /= sum;
+    }
+}
+
+__global__ void softmax_kernel(float *input, int n, int batch, int batch_offset, int groups, int group_offset, int stride, float temp, float *output)
+{
+    int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
+    if (id >= batch*groups) return;
+    int b = id / groups;
+    int g = id % groups;
+    softmax_device(input + b*batch_offset + g*group_offset, n, temp, stride, output + b*batch_offset + g*group_offset);
+}
+
+/**
+    softmax function
+*/
+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)
+{
+    int size = groups*batch;
+    int blocks = (size+255)/256;
+    int threads = 256;
+
+    softmax_kernel<<<blocks, threads, 0, stream>>>(input, n, batch, batch_offset, groups, group_offset, stride, temp, output);
+}

src/pluginsRT/ActivationLeakyRT.cpp

@@ -0,0 +1,60 @@
+#include<cassert>
+#include "kernels.h"
+
+class ActivationLeakyRT : public IPlugin {
+
+public:
+	ActivationLeakyRT() {
+
+
+	}
+
+	~ActivationLeakyRT(){
+
+	}
+
+	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 {
+		size = 1;
+		for(int i=0; i<outputDims[0].nbDims; i++)
+			size *= outputDims[0].d[i];
+	}
+
+	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 {
+
+		activationLEAKYForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
+											reinterpret_cast<dnnType*>(outputs[0]), size, stream);
+		return 0;
+	}
+
+
+	virtual size_t getSerializationSize() override {
+		return 1*sizeof(int);
+	}
+
+	virtual void serialize(void* buffer) override {
+		char *buf = reinterpret_cast<char*>(buffer);
+		tkDNN::writeBUF(buf, size);
+	}
+
+	int size;
+};

src/pluginsRT/Int8Calibrator.cpp

@@ -0,0 +1,168 @@
+#include <vector>
+#include <assert.h>
+#include <algorithm>
+#include <iterator>
+
+#include "NvInfer.h"
+
+class BatchStream
+{
+public:
+	BatchStream(tkDNN::dataDim_t dim, int batchSize, int maxBatches)
+	{
+        mBatchSize = batchSize;
+        mMaxBatches = maxBatches;
+		mDims = nvinfer1::DimsNCHW{ dim.n, dim.c, dim.h, dim.w };
+		mImageSize = mDims.c()*mDims.h()*mDims.w();
+		mBatch.resize(mBatchSize*mImageSize, 0);
+		mLabels.resize(mBatchSize, 0);
+		mFileBatch.resize(mDims.n()*mImageSize, 0);
+		mFileLabels.resize(mDims.n(), 0);
+		reset(0);
+	}
+
+	void reset(int firstBatch)
+	{
+		mBatchCount = 0;
+		mFileCount = 0;
+		mFileBatchPos = mDims.n();
+		skip(firstBatch);
+	}
+
+	bool next()
+	{
+        std::cout<<"Next batch: "<<mBatchCount<<" of "<<mMaxBatches<<"\n";
+		if (mBatchCount == mMaxBatches)
+			return false;
+
+		for (int csize = 1, batchPos = 0; batchPos < mBatchSize; batchPos += csize, mFileBatchPos += csize)
+		{
+			assert(mFileBatchPos > 0 && mFileBatchPos <= mDims.n());
+			if (mFileBatchPos == mDims.n() && !update())
+				return false;
+
+			// copy the smaller of: elements left to fulfill the request, or elements left in the file buffer.
+			csize = std::min(mBatchSize - batchPos, mDims.n() - mFileBatchPos);
+			std::copy_n(getFileBatch() + mFileBatchPos * mImageSize, csize * mImageSize, getBatch() + batchPos * mImageSize);
+			std::copy_n(getFileLabels() + mFileBatchPos, csize, getLabels() + batchPos);
+		}
+		mBatchCount++;
+		return true;
+	}
+
+	void skip(int skipCount)
+	{
+		if (mBatchSize >= mDims.n() && mBatchSize%mDims.n() == 0 && mFileBatchPos == mDims.n())
+		{
+			mFileCount += skipCount * mBatchSize / mDims.n();
+            std::cout<<mFileCount<<"\n";
+			return;
+		}
+
+		int x = mBatchCount;
+		for (int i = 0; i < skipCount; i++)
+			next();
+		mBatchCount = x;
+	}
+
+	float *getBatch() { return &mBatch[0]; }
+	float *getLabels() { return &mLabels[0]; }
+	int getBatchesRead() const { return mBatchCount; }
+	int getBatchSize() const { return mBatchSize; }
+	nvinfer1::DimsNCHW getDims() const { return mDims; }
+private:
+	float* getFileBatch() { return &mFileBatch[0]; }
+	float* getFileLabels() { return &mFileLabels[0]; }
+
+	bool update()
+	{
+		std::string inputFileName = std::string("calibBatches/batch") + std::to_string(mFileCount++);
+		FILE * file = fopen(inputFileName.c_str(), "rb");
+		if (!file) {
+            FatalError("cant open batch calib file: " + inputFileName);
+        	return false;
+        }
+
+		size_t readInputCount = fread(getFileBatch(), sizeof(float), mDims.n()*mImageSize, file);
+		size_t readLabelCount = fread(getFileLabels(), sizeof(float), mDims.n(), file);;
+		assert(readInputCount == size_t(mDims.n()*mImageSize) && readLabelCount == size_t(mDims.n()));
+
+		fclose(file);
+		mFileBatchPos = 0;
+		return true;
+	}
+
+	int mBatchSize{ 0 };
+	int mMaxBatches{ 0 };
+	int mBatchCount{ 0 };
+
+	int mFileCount{ 0 }, mFileBatchPos{ 0 };
+	int mImageSize{ 0 };
+
+	nvinfer1::DimsNCHW mDims;
+	std::vector<float> mBatch;
+	std::vector<float> mLabels;
+	std::vector<float> mFileBatch;
+	std::vector<float> mFileLabels;
+};
+
+
+
+class Int8EntropyCalibrator : public IInt8EntropyCalibrator
+{
+public:
+	Int8EntropyCalibrator(BatchStream& stream, int firstBatch, bool readCache = true)
+		: mStream(stream), mReadCache(readCache)
+	{
+		DimsNCHW dims = mStream.getDims();
+		mInputCount = mStream.getBatchSize() * dims.c() * dims.h() * dims.w();
+		checkCuda(cudaMalloc(&mDeviceInput, mInputCount * sizeof(float)));
+		mStream.reset(firstBatch);
+	}
+
+	virtual ~Int8EntropyCalibrator()
+	{
+		checkCuda(cudaFree(mDeviceInput));
+	}
+
+	int getBatchSize() const override { return mStream.getBatchSize(); }
+
+	bool getBatch(void* bindings[], const char* names[], int nbBindings) override
+	{
+        std::cout<<"CALIB request batch\n";
+		if (!mStream.next())
+			return false;
+
+		checkCuda(cudaMemcpy(mDeviceInput, mStream.getBatch(), mInputCount * sizeof(float), cudaMemcpyHostToDevice));
+		bindings[0] = mDeviceInput;
+		return true;
+	}
+
+	const void* readCalibrationCache(size_t& length) override
+	{
+		mCalibrationCache.clear();
+		std::ifstream input("table.calib", std::ios::binary);
+		input >> std::noskipws;
+		
+        FatalError("rewrite different");
+        //if (mReadCache && input.good())
+		//	std::copy(std::istream_iterator<char>(input), std::istream_iterator<char>(), std::back_inserter(mCalibrationCache));
+
+		length = mCalibrationCache.size();
+		return length ? &mCalibrationCache[0] : nullptr;
+	}
+
+	void writeCalibrationCache(const void* cache, size_t length) override
+	{
+		std::ofstream output("table.calib", std::ios::binary);
+		output.write(reinterpret_cast<const char*>(cache), length);
+	}
+
+private:
+	BatchStream mStream;
+	bool mReadCache{ true };
+
+	size_t mInputCount;
+	void* mDeviceInput{ nullptr };
+	std::vector<char> mCalibrationCache;
+};

src/pluginsRT/RegionRT.cpp

@@ -0,0 +1,94 @@
+#include<cassert>
+#include "kernels.h"
+
+class RegionRT : public IPlugin {
+
+public:
+	RegionRT(int classes, int coords, int num) {
+
+		this->classes = classes;
+		this->coords = coords;
+		this->num = num;
+	}
+
+	~RegionRT(){
+
+	}
+
+	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, coords, batchSize);
+				activationLOGISTICForward(srcData + index, dstData + index, w*h, stream);
+			}
+		}
+
+		//softmax start
+		int index = entry_index(0, 0, coords + 1, batchSize);
+		softmaxForward(	srcData + index, classes, batchSize*num, 
+						(batchSize*c*h*w)/num, 
+						w*h, 1, w*h, 1, dstData + index, stream);
+
+		return 0;
+	}
+
+
+	virtual size_t getSerializationSize() override {
+		return 6*sizeof(int) + 1*sizeof(float);
+	}
+
+	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);
+	}
+
+	int c, h, w;
+    int classes, coords, num;
+
+	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*(coords+classes+1) + entry*w*h + loc;
+	}
+
+};

src/pluginsRT/ReorgRT.cpp

@@ -0,0 +1,63 @@
+#include<cassert>
+#include "kernels.h"
+
+class ReorgRT : public IPlugin {
+
+public:
+	ReorgRT(int stride) {
+		this->stride = stride;
+	}
+
+	~ReorgRT(){
+
+	}
+
+	int getNbOutputs() const override {
+		return 1;
+	}
+
+	Dims getOutputDimensions(int index, const Dims* inputs, int nbInputDims) override {
+		return DimsCHW{inputs[0].d[0]*stride*stride, 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 {
+
+		reorgForward((dnnType*)reinterpret_cast<const dnnType*>(inputs[0]), 
+					  reinterpret_cast<dnnType*>(outputs[0]), 
+					  batchSize, c, h, w, stride, stream);
+		return 0;
+	}
+
+
+	virtual size_t getSerializationSize() override {
+		return 4*sizeof(int);
+	}
+
+	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);
+	}
+
+	int c, h, w, stride;
+};

src/utils.cpp

@@ -1,6 +1,27 @@
 #include "utils.h"
+#include <string.h>
 
-void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d)
+void printCenteredTitle(const char *title, char fill, int dim) {
+
+    int len = strlen(title);
+    int first = dim/2 + len/2;
+    
+    if(len >0)
+        std::cout<<"\n";
+    std::cout.width(first); std::cout.fill(fill); std::cout<<std::right<<title;
+    std::cout.width(dim - first); std::cout<<"\n"; 
+    std::cout.fill(' ');
+}
+
+bool fileExist(const char *fname) {
+    std::ifstream dataFile (fname, std::ios::in | std::ios::binary);
+    if(!dataFile)
+        return false;
+    return true;
+}
+
+
+void readBinaryFile(const char* 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;
@@ -9,8 +30,13 @@ void readBinaryFile(const char* fname, int size, value_type** data_h, value_type
         error_s << "Error opening file " << fname; 
         FatalError(error_s.str());
     }
-    int size_b = size*sizeof(value_type);
-    *data_h = new value_type[size];
+
+    if(seek != 0) {
+        dataFile.seekg(seek*sizeof(dnnType), dataFile.cur);
+    }
+
+    int size_b = size*sizeof(dnnType);
+    *data_h = new dnnType[size];
     if (!dataFile.read ((char*) *data_h, size_b)) 
     {
         error_s << "Error reading file " << fname; 
@@ -18,49 +44,97 @@ void readBinaryFile(const char* fname, int size, value_type** data_h, value_type
     }
     
     checkCuda( cudaMalloc(data_d, size_b) );
-    checkCuda( cudaMemcpy(*data_d, *data_h,
-                                size_b,
-                                cudaMemcpyHostToDevice) );
+    checkCuda( cudaMemcpy(*data_d, *data_h, size_b, cudaMemcpyHostToDevice) );
 }
 
-void printDeviceVector(int size, value_type* vec_d)
+void printDeviceVector(int size, dnnType* vec_d, bool device)
 {
-    value_type *vec;
-    vec = new value_type[size];
-    cudaDeviceSynchronize();
-    cudaMemcpy(vec, vec_d, size*sizeof(value_type), cudaMemcpyDeviceToHost);
-    for (int i = 0; i < size; i++)
-    {
+    dnnType *vec;
+    if(device) {
+        vec = new dnnType[size];
+        cudaDeviceSynchronize();
+        cudaMemcpy(vec, vec_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost);
+    } else {
+        vec = vec_d;
+    }
+    
+    for (int i = 0; i < size; i++) {
         std::cout << vec[i] << " ";
     }
     std::cout << std::endl;
-    delete [] vec;
+
+    if(device)
+        delete [] vec;
 }
 
-void resize(int size, value_type **data)
+int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) {
+
+    dnnType *data_h, *correct_h;
+    const float eps = 0.001f;
+
+    if(device) {
+        data_h = new dnnType[size];
+        correct_h = new dnnType[size];
+        cudaDeviceSynchronize();
+        cudaMemcpy(data_h, data_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost);
+        cudaMemcpy(correct_h, correct_d, size*sizeof(dnnType), cudaMemcpyDeviceToHost);
+    
+    } else {
+        data_h = data_d;
+        correct_h = correct_d;
+    }
+
+    int diffs = 0;
+    for(int i=0; i<size; i++) {
+        if(data_h[i] != data_h[i] || correct_h[i] != correct_h[i] || //nan control
+           fabs(data_h[i] - correct_h[i]) > eps) {
+            diffs += 1;
+            if(diffs == 1)
+                std::cout<<"\n";
+            if(diffs < 10)
+                std::cout<<" | [ "<<i<<" ]: "<<data_h[i]<<" "<<correct_h[i]<<"\n";
+        }
+    }
+
+    if(device) {
+        delete [] data_h;
+        delete [] correct_h;
+    }
+
+    std::cout<<" | ";
+    if(diffs == 0)
+        std::cout<<COL_GREENB<<"OK";
+    else
+        std::cout<<COL_REDB<<"Wrongs: "<<diffs;
+
+    std::cout<<COL_END<<" ~"<<eps<<"\n";
+    return diffs;
+}
+
+void resize(int size, dnnType **data)
 {
     if (*data != NULL)
         checkCuda( cudaFree(*data) );
-    checkCuda( cudaMalloc(data, size*sizeof(value_type)) );
+    checkCuda( cudaMalloc(data, size*sizeof(dnnType)) );
 }
 
-void matrixTranspose(cublasHandle_t handle, value_type* srcData, value_type* dstData, int rows, int cols) {
+void matrixTranspose(cublasHandle_t handle, dnnType* srcData, dnnType* dstData, int rows, int cols) {
 
-    value_type *A = srcData, *clone = dstData;
+    dnnType *A = srcData, *clone = dstData;
     int m = rows, n= cols;
-    checkCuda( cudaMemcpy(clone, A, m*n*sizeof(value_type), cudaMemcpyDeviceToDevice));
+    checkCuda( cudaMemcpy(clone, A, m*n*sizeof(dnnType), cudaMemcpyDeviceToDevice));
    
     float const alpha(1.0);
     float const beta(0.0);
     checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, m, n, &alpha, A, n, &beta, A, m, clone, m ));
 }
 
-void matrixMulAdd(  cublasHandle_t handle, value_type* srcData, value_type* dstData, 
-                    value_type* add_vector, int dim, value_type mul) {
+void matrixMulAdd(  cublasHandle_t handle, dnnType* srcData, dnnType* dstData, 
+                    dnnType* add_vector, int dim, dnnType mul) {
 
-    checkCuda( cudaMemcpy(dstData, add_vector, dim*sizeof(value_type), cudaMemcpyDeviceToDevice));
+    checkCuda( cudaMemcpy(dstData, add_vector, dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
         
-    value_type alpha = mul;
+    dnnType alpha = mul;
     checkERROR( cublasSaxpy(handle, dim, &alpha, srcData, 1, dstData, 1));
 
-}
+}

tests/build_models.sh

@@ -1,4 +1,11 @@
 #!/bin/bash
+if [ "$1" == "download" ]; then
+    wget https://github.com/ceccocats/tkDNN/releases/download/testData/tkDNN_testwg.tar.gz --no-check-certificate
+    tar -xf tkDNN_testwg.tar.gz
+    rm tkDNN_testwg.tar.gz
+    exit
+fi
+
 echo "build test Model"
 cd test
 python test_model.py

tests/mnist/test.cpp

@@ -1,58 +0,0 @@
-#include<iostream>
-#include "tkdnn.h"
-
-const char *input_bin   = "../tests/mnist/input.bin";
-const char *c0_bin      = "../tests/mnist/layers/Convolution0.bin";
-const char *c0_bias_bin = "../tests/mnist/layers/Convolution0.bias.bin";
-const char *c1_bin      = "../tests/mnist/layers/Convolution1.bin";
-const char *c1_bias_bin = "../tests/mnist/layers/Convolution1.bias.bin";
-const char *d2_bin      = "../tests/mnist/layers/InnerProduct2.bin";
-const char *d2_bias_bin = "../tests/mnist/layers/InnerProduct2.bias.bin";
-const char *d3_bin      = "../tests/mnist/layers/InnerProduct3.bin";
-const char *d3_bias_bin = "../tests/mnist/layers/InnerProduct3.bias.bin";
-const char *output_bin   = "../tests/mnist/output.bin";
-
-int main() {
-
-    // Network layout
-    tkDNN::Network net;
-    tkDNN::dataDim_t dim(1, 1, 28, 28, 1);
-    tkDNN::Layer *l;
-    l = new tkDNN::Conv2d     (&net, dim, 20, 5, 5, 1, 1, c0_bin, c0_bias_bin);
-    l = new tkDNN::Pooling    (&net, l->output_dim, 2, 2, 2, 2, tkDNN::POOLING_MAX);
-    l = new tkDNN::Conv2d     (&net, l->output_dim, 50, 5, 5, 1, 1, c1_bin, c1_bias_bin);
-    l = new tkDNN::Pooling    (&net, l->output_dim, 2, 2, 2, 2, tkDNN::POOLING_MAX);
-    l = new tkDNN::Dense      (&net, l->output_dim, 500, d2_bin, d2_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU);
-    l = new tkDNN::Dense      (&net, l->output_dim, 10, d3_bin, d3_bias_bin);
-    l = new tkDNN::Softmax    (&net, l->output_dim);
- 
-    // Load input
-    value_type *data;
-    value_type *input_h;
-    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
-
-    printDeviceVector(dim.tot(), data);
-    dim.print(); //print initial dimension
-    
-    TIMER_START
-
-    // Inference
-    data = net.infer(dim, data);
-    
-    TIMER_STOP
-    dim.print();   
-
-    // Print result
-    std::cout<<"\n======= RESULT =======\n";
-    printDeviceVector(dim.tot(), data);
- 
-    // Print real test
-    std::cout<<"\n==== CHECK RESULT ====\n";
-    value_type *out;
-    value_type *out_h;
-    readBinaryFile(output_bin, dim.tot(), &out_h, &out);
-    printDeviceVector(dim.tot(), out);
- 
-    return 0;
-}

tests/mnist/test_mnist.cpp

@@ -0,0 +1,72 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin   = "../tests/mnist/input.bin";
+const char *c0_bin      = "../tests/mnist/layers/c0.bin";
+const char *c1_bin      = "../tests/mnist/layers/c1.bin";
+const char *d2_bin      = "../tests/mnist/layers/d2.bin";
+const char *d3_bin      = "../tests/mnist/layers/d3.bin";
+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);
+
+    tkDNN::NetworkRT netRT(&net, "mnist.rt");
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+
+    dnnType *out_data, *out_data2;
+
+    std::cout<<"CUDNN inference:\n"; {
+        dim.print(); //print initial dimension  
+        TIMER_START
+        out_data = net.infer(dim, data);    
+        TIMER_STOP
+        dim.print();   
+    }
+
+    // Print result
+    //std::cout<<"\n======= CUDNN RESULT =======\n";
+    //printDeviceVector(10, out_data);
+ 
+    tkDNN::dataDim_t dim2(1, 1, 28, 28, 1);
+
+    std::cout<<"TENSORRT inference:\n"; {
+        dim2.print();
+        TIMER_START
+        out_data2 = netRT.infer(dim2, data);
+        TIMER_STOP
+        dim2.print();
+    }
+
+    // Print result
+    //std::cout<<"\n======= TENRT RESULT =======\n";
+    //printDeviceVector(10, out_data);
+
+    std::cout<<"\n======= CHECK RESULT =======\n";
+    checkResult(dim.tot(), out_data, out_data2);
+
+ /*
+    // Print real test
+    std::cout<<"\n==== CHECK RESULT ====\n";
+    dnnType *out;
+    dnnType *out_h;
+    readBinaryFile(output_bin, dim.tot(), &out_h, &out);
+    printDeviceVector(dim.tot(), out);
+*/ 
+    return 0;
+}

tests/mnist/test_mnistRT.cpp

@@ -0,0 +1,178 @@
+#include<iostream>
+#include<cassert>
+#include "tkdnn.h"
+#include "NvInfer.h"
+
+const char *input_bin   = "../tests/mnist/input.bin";
+const char *c0_bin      = "../tests/mnist/layers/c0.bin";
+const char *c1_bin      = "../tests/mnist/layers/c1.bin";
+const char *d2_bin      = "../tests/mnist/layers/d2.bin";
+const char *d3_bin      = "../tests/mnist/layers/d3.bin";
+const char *output_bin   = "../tests/mnist/output.bin";
+
+using namespace nvinfer1;
+
+// Logger for info/warning/errors
+class Logger : public ILogger			
+{
+	void log(Severity severity, const char* msg) override
+	{
+		// suppress info-level messages
+		if (severity != Severity::kINFO)
+			std::cout << msg << std::endl;
+	}
+} gLogger;
+
+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);
+ 
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+
+    dim.print(); //print initial dimension
+    
+    // Inference
+    {
+        TIMER_START
+        data = net.infer(dim, data);
+        TIMER_STOP
+        dim.print();   
+    }
+
+    // Print real test
+    std::cout<<"\n==== CHECK CUDNN RESULT ====\n";
+    dnnType *out;
+    dnnType *out_h;
+    readBinaryFile(output_bin, dim.tot(), &out_h, &out);
+    std::cout<<"Diff: "<<checkResult(dim.tot(), out, data)<<"\n";
+ 
+
+    std::cout<<"\n==== TensorRT ====\n";
+	// create the builder
+	IBuilder* builder = nvinfer1::createInferBuilder(gLogger);
+	INetworkDefinition* network = builder->createNetwork();
+
+    DataType dt = DataType::kFLOAT;
+	//  Create input of shape { 1, 1, 28, 28 } with name referenced by "data"
+	auto input = network->addInput("data", dt, DimsCHW{ 1, 28, 28});
+	assert(input != nullptr);
+
+    tkDNN::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.
+	auto conv1 = network->addConvolution(*input, 20, DimsHW{5, 5}, w, b);
+	assert(conv1 != nullptr);
+	conv1->setStride(DimsHW{1, 1});
+
+	// Add a max pooling layer with stride of 2x2 and kernel size of 2x2.
+	auto pool1 = network->addPooling(*conv1->getOutput(0), PoolingType::kMAX, DimsHW{2, 2});
+	assert(pool1 != nullptr);
+	pool1->setStride(DimsHW{2, 2});
+
+    tkDNN::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.
+	auto conv2 = network->addConvolution(*pool1->getOutput(0), 50, DimsHW{5, 5}, w1, b1);
+	assert(conv2 != nullptr);
+	conv2->setStride(DimsHW{1, 1});
+
+	// Add a second max pooling layer with stride of 2x2 and kernel size of 2x3>
+	auto pool2 = network->addPooling(*conv2->getOutput(0), PoolingType::kMAX, DimsHW{2, 2});
+	assert(pool2 != nullptr);
+	pool2->setStride(DimsHW{2, 2});
+
+    tkDNN::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.
+	auto ip1 = network->addFullyConnected(*pool2->getOutput(0), 500, w2, b2);
+	assert(ip1 != nullptr);
+
+	// Add an activation layer using the ReLU algorithm.
+	auto relu1 = network->addActivation(*ip1->getOutput(0), ActivationType::kRELU);
+	assert(relu1 != nullptr);
+
+    tkDNN::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.
+	auto ip2 = network->addFullyConnected(*relu1->getOutput(0), 10, w3, b3);
+	assert(ip2 != nullptr);
+
+	// Add a softmax layer to determine the probability.
+	auto prob = network->addSoftMax(*ip2->getOutput(0));
+	assert(prob != nullptr);
+	prob->getOutput(0)->setName("out");
+
+	network->markOutput(*prob->getOutput(0));
+
+	// Build the engine
+	builder->setMaxBatchSize(1);
+	builder->setMaxWorkspaceSize(1 << 20);
+
+	auto engine = builder->buildCudaEngine(*network);
+	// we don't need the network any more
+	network->destroy();
+
+	IExecutionContext *context = engine->createExecutionContext();
+
+	// run inference
+	// 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.
+	assert(engine->getNbBindings() == 2);
+	void* buffers[2];
+
+	// 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()
+	int inputIndex = engine->getBindingIndex("data"); 
+    int outputIndex = engine->getBindingIndex("out");
+
+    float output[10];
+	// create GPU buffers and a stream
+	checkCuda(cudaMalloc(&buffers[inputIndex], 28*28*sizeof(float)));
+	checkCuda(cudaMalloc(&buffers[outputIndex], 10*sizeof(float)));
+
+	cudaStream_t stream;
+	checkCuda(cudaStreamCreate(&stream));
+
+	// DMA the input to the GPU,  execute the batch asynchronously, and DMA it back:
+    {
+        checkCuda(cudaMemcpyAsync(buffers[inputIndex], input_h, 1 * 28*28* sizeof(float), cudaMemcpyHostToDevice, stream));
+        cudaStreamSynchronize(stream);  //want to test only the inference time
+        TIMER_START
+        context->enqueue(1, buffers, stream, nullptr);
+        TIMER_STOP
+        checkCuda(cudaMemcpyAsync(output, buffers[outputIndex],10*sizeof(float), cudaMemcpyDeviceToHost, stream));
+        cudaStreamSynchronize(stream);
+    }
+
+    std::cout<<"\n==== CHECK CUDNN RESULT ====\n";
+    std::cout<<"Diff: "<<checkResult(dim.tot(), (float*)buffers[outputIndex], data)<<"\n";
+
+	// release the stream and the buffers
+	cudaStreamDestroy(stream);
+	checkCuda(cudaFree(buffers[inputIndex]));
+	checkCuda(cudaFree(buffers[outputIndex]));
+
+	// destroy the engine
+	context->destroy();
+	engine->destroy();
+
+    return 0;
+}

tests/simple/test_simple.cpp

@@ -0,0 +1,47 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin   = "../tests/simple/input.bin";
+const char *c0_bin      = "../tests/simple/layers/c0.bin";
+const char *c1_bin      = "../tests/simple/layers/c1.bin";
+const char *d2_bin      = "../tests/simple/layers/d2.bin";
+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);
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+
+    printDeviceVector(dim.tot(), data);
+    dim.print(); //print initial dimension
+    
+    TIMER_START
+    // Inference
+    data = net.infer(dim, data); dim.print();
+    TIMER_STOP
+
+    // Print result
+    std::cout<<"\n======= RESULT =======\n";
+    printDeviceVector(dim.tot(), data);
+
+    // Print real test
+    std::cout<<"\n==== CHECK RESULT ====\n";
+    dnnType *out;
+    dnnType *out_h;
+    readBinaryFile(output_bin, dim.tot(), &out_h, &out);
+    printDeviceVector(dim.tot(), out);
+    return 0;
+}

tests/test/test.cpp

@@ -1,54 +0,0 @@
-#include<iostream>
-#include "tkdnn.h"
-
-const char *input_bin   = "../tests/test/input.bin";
-const char *c0_bin      = "../tests/test/layers/conv0.bin";
-const char *c0_bias_bin = "../tests/test/layers/conv0.bias.bin";
-const char *c1_bin      = "../tests/test/layers/conv1.bin";
-const char *c1_bias_bin = "../tests/test/layers/conv1.bias.bin";
-const char *d2_bin      = "../tests/test/layers/dense2.bin";
-const char *d2_bias_bin = "../tests/test/layers/dense2.bias.bin";
-const char *output_bin   = "../tests/test/output.bin";
-
-int main() {
-
-    // Network layout
-    tkDNN::Network net;
-    tkDNN::dataDim_t dim(1, 1, 10, 10, 1);
-    tkDNN::Layer *l;
-    l = new tkDNN::Conv2d     (&net, dim, 2, 4, 4, 2, 2, c0_bin, c0_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU);
-    l = new tkDNN::Conv2d     (&net, l->output_dim, 4, 2, 2, 1, 1, c1_bin, c1_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU);
-    l = new tkDNN::Flatten      (&net, l->output_dim);
-    l = new tkDNN::Dense      (&net, l->output_dim, 4, d2_bin, d2_bias_bin);
-    l = new tkDNN::Activation (&net, l->output_dim, CUDNN_ACTIVATION_RELU);
-
-    // Load input
-    value_type *data;
-    value_type *input_h;
-    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
-
-    printDeviceVector(dim.tot(), data);
-    dim.print(); //print initial dimension
-    
-    TIMER_START
-
-    // Inference
-    data = net.infer(dim, data); dim.print();
-    
-
-    TIMER_STOP
-
-    // Print result
-    std::cout<<"\n======= RESULT =======\n";
-    printDeviceVector(dim.tot(), data);
-
-    // Print real test
-    std::cout<<"\n==== CHECK RESULT ====\n";
-    value_type *out;
-    value_type *out_h;
-    readBinaryFile(output_bin, dim.tot(), &out_h, &out);
-    printDeviceVector(dim.tot(), out);
-    return 0;
-}

tests/test_rtinference/rtinference.cpp

@@ -0,0 +1,34 @@
+#include<iostream>
+#include "tkdnn.h"
+#include <stdlib.h>     /* srand, rand */
+
+int main(int argc, char *argv[]) {
+
+    if(argc < 2 || !fileExist(argv[1]))
+        FatalError("unable to read serialRT file");
+
+    //always same test 
+    srand (0); 
+
+    //convert network to tensorRT
+    tkDNN::NetworkRT netRT(NULL, argv[1]);
+
+    dnnType *input = new float[netRT.input_dim.tot()];
+    dnnType *output = new float[netRT.input_dim.tot()];
+
+    printCenteredTitle(" TENSORRT inference ", '=', 30); 
+    for(int i=0; i<100; i++) {
+        for(int j=0; j<netRT.input_dim.tot(); j++)
+            input[j] = ((float) rand() / (RAND_MAX));
+        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
+    }
+    
+    return 0;
+}

tests/yolo/yolo.cfg

@@ -0,0 +1,258 @@
+[net]
+# Testing
+#batch=1
+#subdivisions=1
+# Training
+ batch=32
+ subdivisions=8
+width=608
+height=608
+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 = 500200
+policy=steps
+steps=400000,450000
+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=425
+activation=linear
+
+
+[region]
+anchors =  0.57273, 0.677385, 1.87446, 2.06253, 3.33843, 5.47434, 7.88282, 3.52778, 9.77052, 9.16828
+bias_match=1
+classes=80
+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/yolo.cpp

@@ -0,0 +1,150 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin  = "../tests/yolo/layers/input.bin";
+const char *c0_bin     = "../tests/yolo/layers/c0.bin";
+const char *c2_bin     = "../tests/yolo/layers/c2.bin";
+const char *c4_bin     = "../tests/yolo/layers/c4.bin";
+const char *c5_bin     = "../tests/yolo/layers/c5.bin";
+const char *c6_bin     = "../tests/yolo/layers/c6.bin";
+const char *c8_bin     = "../tests/yolo/layers/c8.bin";
+const char *c9_bin     = "../tests/yolo/layers/c9.bin";
+const char *c10_bin    = "../tests/yolo/layers/c10.bin";
+const char *c12_bin    = "../tests/yolo/layers/c12.bin";
+const char *c13_bin    = "../tests/yolo/layers/c13.bin";
+const char *c14_bin    = "../tests/yolo/layers/c14.bin";
+const char *c15_bin    = "../tests/yolo/layers/c15.bin";
+const char *c16_bin    = "../tests/yolo/layers/c16.bin";
+const char *c18_bin    = "../tests/yolo/layers/c18.bin";
+const char *c19_bin    = "../tests/yolo/layers/c19.bin";
+const char *c20_bin    = "../tests/yolo/layers/c20.bin";
+const char *c21_bin    = "../tests/yolo/layers/c21.bin";
+const char *c22_bin    = "../tests/yolo/layers/c22.bin";
+const char *c23_bin    = "../tests/yolo/layers/c23.bin";
+const char *c24_bin    = "../tests/yolo/layers/c24.bin";
+const char *c26_bin    = "../tests/yolo/layers/c26.bin";
+const char *c29_bin    = "../tests/yolo/layers/c29.bin";
+const char *c30_bin    = "../tests/yolo/layers/c30.bin";
+const char *g31_bin    = "../tests/yolo/layers/g31.bin";
+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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    tkDNN::Layer *m28_layers[2] = { &r27, &a24 };
+    tkDNN::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);
+    
+    tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 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
+    tkDNN::NetworkRT netRT(&net, "yolo.rt");
+
+    dnnType *out_data, *out_data2; // cudnn output, tensorRT output
+
+    tkDNN::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        out_data = net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+    
+    tkDNN::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_224/yolo_224.cfg

@@ -0,0 +1,258 @@
+[net]
+# Testing
+#batch=1
+#subdivisions=1
+# Training
+ batch=64
+ subdivisions=16
+width=224
+height=224
+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 = 500200
+policy=steps
+steps=400000,450000
+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=425
+activation=linear
+
+
+[region]
+anchors =  0.57273, 0.677385, 1.87446, 2.06253, 3.33843, 5.47434, 7.88282, 3.52778, 9.77052, 9.16828
+bias_match=1
+classes=80
+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_224/yolo_224.cpp

@@ -0,0 +1,150 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin  = "../tests/yolo_224/layers/input.bin";
+const char *c0_bin     = "../tests/yolo_224/layers/c0.bin";
+const char *c2_bin     = "../tests/yolo_224/layers/c2.bin";
+const char *c4_bin     = "../tests/yolo_224/layers/c4.bin";
+const char *c5_bin     = "../tests/yolo_224/layers/c5.bin";
+const char *c6_bin     = "../tests/yolo_224/layers/c6.bin";
+const char *c8_bin     = "../tests/yolo_224/layers/c8.bin";
+const char *c9_bin     = "../tests/yolo_224/layers/c9.bin";
+const char *c10_bin    = "../tests/yolo_224/layers/c10.bin";
+const char *c12_bin    = "../tests/yolo_224/layers/c12.bin";
+const char *c13_bin    = "../tests/yolo_224/layers/c13.bin";
+const char *c14_bin    = "../tests/yolo_224/layers/c14.bin";
+const char *c15_bin    = "../tests/yolo_224/layers/c15.bin";
+const char *c16_bin    = "../tests/yolo_224/layers/c16.bin";
+const char *c18_bin    = "../tests/yolo_224/layers/c18.bin";
+const char *c19_bin    = "../tests/yolo_224/layers/c19.bin";
+const char *c20_bin    = "../tests/yolo_224/layers/c20.bin";
+const char *c21_bin    = "../tests/yolo_224/layers/c21.bin";
+const char *c22_bin    = "../tests/yolo_224/layers/c22.bin";
+const char *c23_bin    = "../tests/yolo_224/layers/c23.bin";
+const char *c24_bin    = "../tests/yolo_224/layers/c24.bin";
+const char *c26_bin    = "../tests/yolo_224/layers/c26.bin";
+const char *c29_bin    = "../tests/yolo_224/layers/c29.bin";
+const char *c30_bin    = "../tests/yolo_224/layers/c30.bin";
+const char *g31_bin    = "../tests/yolo_224/layers/g31.bin";
+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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    tkDNN::Layer *m28_layers[2] = { &r27, &a24 };
+    tkDNN::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);
+    
+    tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 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
+    tkDNN::NetworkRT netRT(&net, "yolo_224.rt");
+
+    dnnType *out_data, *out_data2; // cudnn output, tensorRT output
+
+    tkDNN::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        out_data = net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+    
+    tkDNN::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_relu/yolo_relu.cfg

@@ -0,0 +1,258 @@
+[net]
+# Testing
+#batch=1
+#subdivisions=1
+# Training
+ batch=64
+ subdivisions=16
+width=608
+height=608
+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 = 500200
+policy=steps
+steps=400000,450000
+scales=.1,.1
+
+[convolutional]
+batch_normalize=1
+filters=32
+size=3
+stride=1
+pad=1
+activation=relu
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=3
+stride=1
+pad=1
+activation=relu
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=64
+size=1
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=3
+stride=1
+pad=1
+activation=relu
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=128
+size=1
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=3
+stride=1
+pad=1
+activation=relu
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=256
+size=1
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=3
+stride=1
+pad=1
+activation=relu
+
+[maxpool]
+size=2
+stride=2
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=512
+size=1
+stride=1
+pad=1
+activation=relu
+
+[convolutional]
+batch_normalize=1
+filters=1024
+size=3
+stride=1
+pad=1
+activation=relu
+
+
+#######
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=relu
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=relu
+
+[route]
+layers=-9
+
+[convolutional]
+batch_normalize=1
+size=1
+stride=1
+pad=1
+filters=64
+activation=relu
+
+[reorg]
+stride=2
+
+[route]
+layers=-1,-4
+
+[convolutional]
+batch_normalize=1
+size=3
+stride=1
+pad=1
+filters=1024
+activation=relu
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=425
+activation=linear
+
+
+[region]
+anchors =  0.57273, 0.677385, 1.87446, 2.06253, 3.33843, 5.47434, 7.88282, 3.52778, 9.77052, 9.16828
+bias_match=1
+classes=80
+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_relu/yolo_relu.cpp

@@ -0,0 +1,151 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin  = "../tests/yolo_relu/layers/input.bin";
+const char *c0_bin     = "../tests/yolo_relu/layers/c0.bin";
+const char *c2_bin     = "../tests/yolo_relu/layers/c2.bin";
+const char *c4_bin     = "../tests/yolo_relu/layers/c4.bin";
+const char *c5_bin     = "../tests/yolo_relu/layers/c5.bin";
+const char *c6_bin     = "../tests/yolo_relu/layers/c6.bin";
+const char *c8_bin     = "../tests/yolo_relu/layers/c8.bin";
+const char *c9_bin     = "../tests/yolo_relu/layers/c9.bin";
+const char *c10_bin    = "../tests/yolo_relu/layers/c10.bin";
+const char *c12_bin    = "../tests/yolo_relu/layers/c12.bin";
+const char *c13_bin    = "../tests/yolo_relu/layers/c13.bin";
+const char *c14_bin    = "../tests/yolo_relu/layers/c14.bin";
+const char *c15_bin    = "../tests/yolo_relu/layers/c15.bin";
+const char *c16_bin    = "../tests/yolo_relu/layers/c16.bin";
+const char *c18_bin    = "../tests/yolo_relu/layers/c18.bin";
+const char *c19_bin    = "../tests/yolo_relu/layers/c19.bin";
+const char *c20_bin    = "../tests/yolo_relu/layers/c20.bin";
+const char *c21_bin    = "../tests/yolo_relu/layers/c21.bin";
+const char *c22_bin    = "../tests/yolo_relu/layers/c22.bin";
+const char *c23_bin    = "../tests/yolo_relu/layers/c23.bin";
+const char *c24_bin    = "../tests/yolo_relu/layers/c24.bin";
+const char *c26_bin    = "../tests/yolo_relu/layers/c26.bin";
+const char *c29_bin    = "../tests/yolo_relu/layers/c29.bin";
+const char *c30_bin    = "../tests/yolo_relu/layers/c30.bin";
+const char *g31_bin    = "../tests/yolo_relu/layers/g31.bin";
+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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    tkDNN::Layer *m28_layers[2] = { &r27, &a24 };
+    tkDNN::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);
+    
+    tkDNN::RegionInterpret rI(dim, g31.output_dim, 80, 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
+    tkDNN::NetworkRT netRT(&net, "yolo_relu.rt");
+
+    dnnType *out_data, *out_data2; // cudnn output, tensorRT output
+
+    tkDNN::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        out_data = net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+    
+    tkDNN::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, 608, 608);
+    rI.showImageResult(input_h);
+
+    return 0;
+}

tests/yolo_tiny/tiny-yolo.cfg

@@ -0,0 +1,139 @@
+[net]
+ Training
+ batch=64
+ subdivisions=8
+# Testing
+# batch=1
+# subdivisions=1
+width=416
+height=416
+channels=3
+momentum=0.9
+decay=0.0005
+angle=0
+saturation = 1.5
+exposure = 1.5
+hue=.1
+
+learning_rate=0.001
+burn_in=1000
+max_batches = 500200
+policy=steps
+steps=400000,450000
+scales=.1,.1
+
+[convolutional]
+batch_normalize=1
+filters=16
+size=3
+stride=1
+pad=1
+activation=leaky
+
+[maxpool]
+size=2
+stride=2
+
+[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
+
+[maxpool]
+size=2
+stride=2
+
+[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
+
+#[maxpool]
+#size=2
+#stride=1
+
+[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=512
+activation=leaky
+
+[convolutional]
+size=1
+stride=1
+pad=1
+filters=425
+activation=linear
+
+[region]
+anchors =  0.57273, 0.677385, 1.87446, 2.06253, 3.33843, 5.47434, 7.88282, 3.52778, 9.77052, 9.16828
+bias_match=1
+classes=80
+coords=4
+num=5
+softmax=1
+jitter=.2
+rescore=0
+
+object_scale=5
+noobject_scale=1
+class_scale=1
+coord_scale=1
+
+absolute=1
+thresh = .6
+random=1

tests/yolo_tiny/yolo_tiny.cpp

@@ -0,0 +1,93 @@
+#include<iostream>
+#include "tkdnn.h"
+
+const char *input_bin  = "../tests/yolo_tiny/layers/input.bin";
+const char *c0_bin     = "../tests/yolo_tiny/layers/c0.bin";
+const char *c2_bin     = "../tests/yolo_tiny/layers/c2.bin";
+const char *c4_bin     = "../tests/yolo_tiny/layers/c4.bin";
+const char *c5_bin     = "../tests/yolo_tiny/layers/c5.bin";
+const char *c6_bin     = "../tests/yolo_tiny/layers/c6.bin";
+const char *c8_bin     = "../tests/yolo_tiny/layers/c8.bin";
+const char *c10_bin    = "../tests/yolo_tiny/layers/c10.bin";
+const char *c11_bin    = "../tests/yolo_tiny/layers/c11.bin";
+const char *c12_bin    = "../tests/yolo_tiny/layers/c12.bin";
+const char *c13_bin    = "../tests/yolo_tiny/layers/c13.bin";
+const char *g14_bin    = "../tests/yolo_tiny/layers/g14.bin";
+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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    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);
+
+    tkDNN::Conv2d     c10(&net, 512, 3, 3, 1, 1, 1, 1, c10_bin, true);
+    tkDNN::Activation a10(&net, tkDNN::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);
+
+    // Load input
+    dnnType *data;
+    dnnType *input_h;
+    readBinaryFile(input_bin, dim.tot(), &input_h, &data);
+
+    //print network model
+    net.print();
+
+    //convert network to tensorRT
+    tkDNN::NetworkRT netRT(&net, "yolo_tiny.rt");
+
+    dnnType *out_data, *out_data2; // cudnn output, tensorRT output
+
+    tkDNN::dataDim_t dim1 = dim; //input dim
+    printCenteredTitle(" CUDNN inference ", '=', 30); {
+        dim1.print();
+        TIMER_START
+        out_data = net.infer(dim1, data);    
+        TIMER_STOP
+        dim1.print();   
+    }
+ 
+    tkDNN::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);
+    return 0;
+}