NetworkRT (deallocations to be done)

CMakeLists.txt

@@ -15,7 +15,7 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11")
 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/Route.cpp src/Reorg.cpp src/Region.cpp src/Network.cpp src/utils.cpp)
+    src/Route.cpp src/Reorg.cpp src/Region.cpp src/Network.cpp src/utils.cpp src/NetworkRT.cpp)
 target_link_libraries(tkDNN kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer)
 
 add_executable(test_simple tests/simple/test_simple.cpp)

include/Layer.h

@@ -7,6 +7,19 @@
 
 namespace tkDNN {
 
+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
 */
@@ -15,6 +28,7 @@ class Layer {
 public:
     Layer(Network *net);
     virtual ~Layer();
+    virtual layerType_t getLayerType() = 0;
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData) {
         std::cout<<"No infer action for this layer\n";
@@ -63,6 +77,7 @@ class Dense : public LayerWgs {
 public:
     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);
 };
@@ -84,6 +99,7 @@ class Activation : public Layer {
 public:
     Activation(Network *net, int act_mode); 
     virtual ~Activation();
+    virtual layerType_t getLayerType() { return LAYER_ACTIVATION; };
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData);
 
@@ -103,10 +119,11 @@ public:
                 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);
 
-    int kernelH, kernelW, strideH, strideW;
+    int kernelH, kernelW, strideH, strideW, paddingH, paddingW;
 
 protected:
     cudnnFilterDescriptor_t filterDesc;
@@ -128,6 +145,7 @@ class Flatten : public Layer {
 public:
     Flatten(Network *net); 
     virtual ~Flatten();
+    virtual layerType_t getLayerType() { return LAYER_FLATTEN; };
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData);
 };
@@ -142,6 +160,7 @@ class MulAdd : public Layer {
 public:
     MulAdd(Network *net, value_type mul, value_type add); 
     virtual ~MulAdd();
+    virtual layerType_t getLayerType() { return LAYER_MULADD; };
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData);
 
@@ -168,18 +187,19 @@ typedef enum {
 class Pooling : public Layer {
 
 public:
+    int winH, winW;
+    int strideH, strideW;
+
     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);
 
 protected:
 
     cudnnPoolingDescriptor_t poolingDesc;
-
-    int winH, winW;
-    int strideH, strideW;
     tkdnnPoolingMode_t pool_mode;
     value_type *tmpInputData, *tmpOutputData;
     bool poolOn3d;
@@ -193,6 +213,7 @@ class Softmax : public Layer {
 public:
     Softmax(Network *net); 
     virtual ~Softmax();
+    virtual layerType_t getLayerType() { return LAYER_SOFTMAX; };
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData);
 };
@@ -206,6 +227,7 @@ class Route : public Layer {
 public:
     Route(Network *net, Layer **layers, int layers_n); 
     virtual ~Route();
+    virtual layerType_t getLayerType() { return LAYER_ROUTE; };
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData);
 
@@ -224,6 +246,7 @@ class Reorg : public Layer {
 public:
     Reorg(Network *net, int stride);
     virtual ~Reorg();
+    virtual layerType_t getLayerType() { return LAYER_REORG; };
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData);
 
@@ -240,6 +263,7 @@ class Region : public Layer {
 public:
     Region(Network *net, int classes, int coords, int num, float thresh);
     virtual ~Region();
+    virtual layerType_t getLayerType() { return LAYER_REGION; };
 
     virtual value_type* infer(dataDim_t &dim, value_type* srcData);
 

include/NetworkRT.h

@@ -0,0 +1,45 @@
+#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::INetworkDefinition *networkRT; 
+    
+    nvinfer1::ICudaEngine *engineRT;
+    nvinfer1::IExecutionContext *contextRT;
+    void* buffersRT[2];
+    int buf_input_idx, buf_output_idx;
+
+    dataDim_t output_dim;
+    value_type *output;
+    cudaStream_t stream;
+
+    NetworkRT(Network *net);
+    virtual ~NetworkRT();
+
+    /**
+        Do inferece
+    */
+    value_type* infer(dataDim_t &dim, value_type* data);
+
+    nvinfer1::ITensor* convert_layer(nvinfer1::ITensor *input, Layer *l);
+    nvinfer1::ITensor* convert_layer(nvinfer1::ITensor *input, Conv2d *l);
+    nvinfer1::ITensor* convert_layer(nvinfer1::ITensor *input, Activation *l);
+    nvinfer1::ITensor* convert_layer(nvinfer1::ITensor *input, Dense *l);
+    nvinfer1::ITensor* convert_layer(nvinfer1::ITensor *input, Pooling *l);
+    nvinfer1::ITensor* convert_layer(nvinfer1::ITensor *input, Softmax *l);
+};
+
+
+}
+#endif //NETWORKRT_H

include/tkdnn.h

@@ -3,6 +3,7 @@
 */
 #include "Network.h"
 #include "Layer.h"
+#include "NetworkRT.h"
 
 namespace tkDNN {
 

include/utils.h

@@ -62,6 +62,14 @@
     }                                                                  \
 }
 
+#define checkNULL(ptr) {                                               \
+    std::stringstream _error;                                          \
+    if (ptr == nullptr) {                                              \
+      _error << "Null pointer";                                        \
+      FatalError(_error.str());                                        \
+    }                                                                  \
+}
+
 void readBinaryFile(const char* fname, int size, value_type** data_h, value_type** data_d, int seek = 0);
 int  checkResult(int size, value_type *data_d, value_type *correct_d, bool device = true);
 void printDeviceVector(int size, value_type* vec_d, bool device = true);

src/Conv2d.cpp

@@ -15,6 +15,8 @@ Conv2d::Conv2d( Network *net, int out_ch, int kernelH, int kernelW,
     this->kernelW = kernelW;
     this->strideH = strideH;
     this->strideW = strideW;
+    this->paddingH = paddingH;
+    this->paddingW = paddingW;
 
     checkCUDNN( cudnnCreateFilterDescriptor(&filterDesc) );
     checkCUDNN( cudnnCreateConvolutionDescriptor(&convDesc) );

src/Network.cpp

@@ -37,6 +37,7 @@ value_type* Network::infer(dataDim_t &dim, value_type* data) {
     for(int i=0; i<num_layers; i++)
         data = layers[i]->infer(dim, data);
 
+    checkCuda(cudaDeviceSynchronize());
     return data;
 }
 

src/NetworkRT.cpp

@@ -0,0 +1,201 @@
+#include <iostream>
+#include "NvInfer.h"
+
+#include "NetworkRT.h"
+
+using namespace nvinfer1;
+
+// Logger for info/warning/errors
+class Logger : public ILogger			
+{
+	void log(Severity severity, const char* msg) override
+	{
+		std::cout <<"TENSORRT: "<< msg << std::endl;
+	}
+} loggerRT;
+
+namespace tkDNN {
+
+NetworkRT::NetworkRT(Network *net) {
+
+    builderRT = createInferBuilder(loggerRT);
+    networkRT = builderRT->createNetwork();
+    dtRT = DataType::kFLOAT;
+
+    //add input layer
+    dataDim_t dim = net->layers[0]->input_dim;
+    ITensor *input = networkRT->addInput("data", dtRT, 
+                     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];
+        input = convert_layer(input, l);
+    }
+    if(input == NULL)
+        FatalError("conversion failed");
+    output_dim = net->layers[net->num_layers-1]->output_dim;
+
+    //build tensorRT
+    input->setName("out");
+	networkRT->markOutput(*input);
+
+	// Build the engine
+	builderRT->setMaxBatchSize(1);
+	builderRT->setMaxWorkspaceSize(1 << 20);
+
+    std::cout<<"BUILD cuda engine\n";
+	engineRT = builderRT->buildCudaEngine(*networkRT);
+	// we don't need the network any more
+	//networkRT->destroy();
+
+    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";
+
+	// create GPU buffers and a stream
+    checkCuda(cudaMalloc(&buffersRT[buf_input_idx],  dim.tot()*sizeof(value_type)));
+    checkCuda(cudaMalloc(&buffersRT[buf_output_idx], output_dim.tot()*sizeof(value_type)));
+    checkCuda(cudaMalloc(&output, output_dim.tot()*sizeof(value_type)));
+	checkCuda(cudaStreamCreate(&stream));
+}
+
+NetworkRT::~NetworkRT() {
+
+}
+
+value_type* NetworkRT::infer(dataDim_t &dim, value_type* data) {
+
+    checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, 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;
+}
+
+ITensor* 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);
+
+    FatalError("Layer not implemented in tensorRT");
+    return NULL;
+}
+
+ITensor* NetworkRT::convert_layer(ITensor *input, Dense *l) {
+    std::cout<<"convert Dense\n";
+
+    Weights w { dtRT, l->data_h, l->inputs*l->outputs};
+    Weights b = { dtRT, l->bias_h, l->outputs};
+    IFullyConnectedLayer *lRT = networkRT->addFullyConnected(*input, l->outputs, w, b);
+
+    checkNULL(lRT);
+    return lRT->getOutput(0);
+}
+
+ITensor* NetworkRT::convert_layer(ITensor *input, Conv2d *l) {
+    std::cout<<"convert conv2D\n";
+
+    Weights w { dtRT, l->data_h, l->inputs*l->outputs*l->kernelH*l->kernelW};
+    Weights b;
+    if(!l->batchnorm)
+        b = { dtRT, l->bias_h, 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) {
+        float eps = CUDNN_BN_MIN_EPSILON;
+
+        //make power array of ones
+        value_type *power_h = new value_type[l->outputs];
+        for(int i=0; i<l->outputs; i++) power_h[i] = 1.0f;
+
+        //convert mean
+        for(int i=0; i<l->outputs; i++)
+            l->mean_h[i] = l->mean_h[i] / -sqrt(eps + l->variance_h[i]); 
+        
+        //convert variance
+        for(int i=0; i<l->outputs; i++)
+            l->variance_h[i] = 1.0f / sqrt(eps + l->variance_h[i]); 
+
+        Weights power{dtRT, power_h, l->outputs};
+        Weights shift{dtRT, l->mean_h, l->outputs};
+        Weights scale{dtRT, l->variance_h, l->outputs};
+        IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, 
+                    shift, scale, power);
+        checkNULL(lRT2);
+
+        Weights shift2{dtRT, l->bias_h, l->outputs};
+        Weights scale2{dtRT, l->scales_h, l->outputs};
+        IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, 
+                    shift2, scale2, power);
+        checkNULL(lRT3);
+
+        return lRT3->getOutput(0);
+    }
+
+    return lRT->getOutput(0);
+}
+
+ITensor* 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->getOutput(0);
+}
+
+ITensor* NetworkRT::convert_layer(ITensor *input, Activation *l) {
+    std::cout<<"convert Activation\n";
+
+    IActivationLayer *lRT = networkRT->addActivation(*input, ActivationType::kRELU);
+    checkNULL(lRT);
+
+    return lRT->getOutput(0);
+}
+
+ITensor* NetworkRT::convert_layer(ITensor *input, Softmax *l) {
+    std::cout<<"convert Activation\n";
+
+    ISoftMaxLayer *lRT = networkRT->addSoftMax(*input);
+    checkNULL(lRT);
+
+    return lRT->getOutput(0);
+}
+
+}

src/utils.cpp

@@ -69,7 +69,8 @@ int checkResult(int size, value_type *data_d, value_type *correct_d, bool device
     for(int i=0; i<size; i++) {
         if(fabs(data_h[i] - correct_h[i]) > 0.0001) {
             diffs += 1;
-            printf("%d\n", i);
+            if(diffs < 10)
+                printf("%f %f\n", data_h[i], correct_h[i]);
         }
     }
 

tests/mnist/test_mnist.cpp

@@ -21,33 +21,52 @@ int main() {
     tkDNN::Activation l5(&net, CUDNN_ACTIVATION_RELU);
     tkDNN::Dense      l6(&net, 10, d3_bin);
     tkDNN::Softmax    l7(&net);
- 
+
+    tkDNN::NetworkRT netRT(&net);
+
     // Load input
     value_type *data;
     value_type *input_h;
     readBinaryFile(input_bin, dim.tot(), &input_h, &data);
 
-    printDeviceVector(dim.tot(), data);
+    value_type *out_data, *out_data2;
-    dim.print(); //print initial dimension
-    
-    TIMER_START
 
-    // Inference
+    std::cout<<"CUDNN inference:\n"; {
-    data = net.infer(dim, data);
+        dim.print(); //print initial dimension  
-    
+        TIMER_START
-    TIMER_STOP
+        out_data = net.infer(dim, data);    
-    dim.print();   
+        TIMER_STOP
+        dim.print();   
+    }
 
     // Print result
-    std::cout<<"\n======= RESULT =======\n";
+    //std::cout<<"\n======= CUDNN RESULT =======\n";
-    printDeviceVector(dim.tot(), data);
+    //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";
+    std::cout<<"Diffs: "<<checkResult(dim.tot(), out_data, out_data2)<<"\n";
+
+ /*
     // 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/yolo/yolo.cpp

@@ -104,22 +104,30 @@ int main() {
     value_type *input_h;
     readBinaryFile(input_bin, dim.tot(), &input_h, &data);
 
-    dim.print(); //print initial dimension
+    tkDNN::NetworkRT netRT(&net);
-    
-    TIMER_START
 
-    // Inference
+    value_type *out_data, *out_data2;
-    data = net.infer(dim, data);
+
-    
+    std::cout<<"CUDNN inference:\n"; {
-    TIMER_STOP
+        dim.print(); //print initial dimension  
-    dim.print();   
+        TIMER_START
+        out_data = net.infer(dim, data);    
+        TIMER_STOP
+        dim.print();   
+    }
  
-    // Print real test
+    tkDNN::dataDim_t dim2(1, 3, 608, 608, 1);
-    std::cout<<"\n==== CHECK RESULT ====\n";
+
-    value_type *out;
+    std::cout<<"TENSORRT inference:\n"; {
-    value_type *out_h;
+        dim2.print();
-    readBinaryFile(output_bin, dim.tot(), &out_h, &out);
+        TIMER_START
-    int diff = checkResult(dim.tot(), data, out);
+        out_data2 = netRT.infer(dim2, data);
-    printf("Output diffs: %d\n", diff);
+        TIMER_STOP
+        dim2.print();
+    }
+
+    std::cout<<"\n======= CHECK RESULT =======\n";
+    std::cout<<"Diffs: "<<checkResult(dim.tot(), out_data, out_data2)<<"\n";
+
     return 0;
 }