diff --git a/CMakeLists.txt b/CMakeLists.txt index 4e5e4e6..0ebdea5 100644 --- a/CMakeLists.txt +++ b/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) diff --git a/include/Layer.h b/include/Layer.h index 8a15711..06b28f6 100644 --- a/include/Layer.h +++ b/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); diff --git a/include/NetworkRT.h b/include/NetworkRT.h new file mode 100644 index 0000000..520bf65 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/include/tkdnn.h b/include/tkdnn.h index a008d91..372ca48 100644 --- a/include/tkdnn.h +++ b/include/tkdnn.h @@ -3,6 +3,7 @@ */ #include "Network.h" #include "Layer.h" +#include "NetworkRT.h" namespace tkDNN { diff --git a/include/utils.h b/include/utils.h index 98a3fea..e724dc2 100644 --- a/include/utils.h +++ b/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); diff --git a/src/Conv2d.cpp b/src/Conv2d.cpp index 1821194..e553baf 100644 --- a/src/Conv2d.cpp +++ b/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) ); diff --git a/src/Network.cpp b/src/Network.cpp index a901fcd..601099a 100644 --- a/src/Network.cpp +++ b/src/Network.cpp @@ -37,6 +37,7 @@ value_type* Network::infer(dataDim_t &dim, value_type* data) { for(int i=0; iinfer(dim, data); + checkCuda(cudaDeviceSynchronize()); return data; } diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp new file mode 100644 index 0000000..23ed901 --- /dev/null +++ b/src/NetworkRT.cpp @@ -0,0 +1,201 @@ +#include +#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; inum_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 = "< output index = "<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; ioutputs; i++) power_h[i] = 1.0f; + + //convert mean + for(int i=0; ioutputs; i++) + l->mean_h[i] = l->mean_h[i] / -sqrt(eps + l->variance_h[i]); + + //convert variance + for(int i=0; ioutputs; 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); +} + +} \ No newline at end of file diff --git a/src/utils.cpp b/src/utils.cpp index 9733a79..006fbea 100644 --- a/src/utils.cpp +++ b/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 0.0001) { diffs += 1; - printf("%d\n", i); + if(diffs < 10) + printf("%f %f\n", data_h[i], correct_h[i]); } } diff --git a/tests/mnist/test_mnist.cpp b/tests/mnist/test_mnist.cpp index ba9e216..f85cfcd 100644 --- a/tests/mnist/test_mnist.cpp +++ b/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); - dim.print(); //print initial dimension - - TIMER_START + value_type *out_data, *out_data2; - // Inference - data = net.infer(dim, data); - - TIMER_STOP - dim.print(); + 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======= RESULT =======\n"; - printDeviceVector(dim.tot(), data); + //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"; + std::cout<<"Diffs: "<