diff --git a/CMakeLists.txt b/CMakeLists.txt index 0ebdea5..a133f1c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -28,4 +28,7 @@ add_executable(test_mnistRT tests/mnist/test_mnistRT.cpp) target_link_libraries(test_mnistRT tkDNN) add_executable(test_yolo tests/yolo/yolo.cpp) -target_link_libraries(test_yolo tkDNN) \ No newline at end of file +target_link_libraries(test_yolo tkDNN) + +add_executable(test_yolo_tiny tests/yolo-tiny/yolo-tiny.cpp) +target_link_libraries(test_yolo_tiny tkDNN) \ No newline at end of file diff --git a/include/Layer.h b/include/Layer.h index 0156cf4..4b5c6ae 100644 --- a/include/Layer.h +++ b/include/Layer.h @@ -190,6 +190,7 @@ class Pooling : public Layer { public: int winH, winW; int strideH, strideW; + int paddingH, paddingW; Pooling(Network *net, int winH, int winW, int strideH, int strideW, tkdnnPoolingMode_t pool_mode); diff --git a/src/Network.cpp b/src/Network.cpp index d96af55..096cb09 100644 --- a/src/Network.cpp +++ b/src/Network.cpp @@ -32,9 +32,10 @@ Network::~Network() { value_type* Network::infer(dataDim_t &dim, value_type* data) { //do infer for every layer - for(int i=0; iinfer(dim, data); - + //dim.print(); + } checkCuda(cudaDeviceSynchronize()); return data; } diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 91f1e5c..9947c63 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -45,7 +45,11 @@ NetworkRT::NetworkRT(Network *net) { } if(input == NULL) FatalError("conversion failed"); - output_dim = net->layers[net->num_layers-1]->output_dim; + output_dim = dim; + Dims oDim = input->getDimensions(); + output_dim.c = oDim.d[0]; + output_dim.h = oDim.d[1]; + output_dim.w = oDim.d[2]; //build tensorRT input->setName("out"); diff --git a/src/Pooling.cpp b/src/Pooling.cpp index fa442c6..e4cacbd 100644 --- a/src/Pooling.cpp +++ b/src/Pooling.cpp @@ -5,20 +5,18 @@ namespace tkDNN { -Pooling::Pooling( Network *net, int winH, int winW, - int strideH, int strideW, tkdnnPoolingMode_t pool_mode) : +Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, + tkdnnPoolingMode_t pool_mode) : Layer(net) { - - if(winH != strideH || winW != strideW) - FatalError("stride pooling not yet implemented"); - 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,12 +44,13 @@ Pooling::Pooling( Network *net, int winH, int winW, 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; diff --git a/tests/yolo-tiny/yolo-tiny.cpp b/tests/yolo-tiny/yolo-tiny.cpp new file mode 100644 index 0000000..a2dd201 --- /dev/null +++ b/tests/yolo-tiny/yolo-tiny.cpp @@ -0,0 +1,90 @@ +#include +#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 *c12_bin = "../tests/yolo-tiny/layers/c12.bin"; +const char *c13_bin = "../tests/yolo-tiny/layers/c13.bin"; +const char *c14_bin = "../tests/yolo-tiny/layers/c14.bin"; +const char *output_bin = "../tests/yolo-tiny/layers/outputLEL.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::Pooling p11(&net, 2, 2, 1, 1, tkDNN::POOLING_MAX); + + tkDNN::Conv2d c12(&net, 1024, 3, 3, 1, 1, 1, 1, c12_bin, true); + tkDNN::Activation a12(&net, tkDNN::ACTIVATION_LEAKY); + tkDNN::Conv2d c13(&net, 1024, 3, 3, 1, 1, 1, 1, c13_bin, true); + tkDNN::Activation a13(&net, tkDNN::ACTIVATION_LEAKY); + tkDNN::Conv2d c14(&net, 125, 1, 1, 1, 1, 0, 0, c14_bin, false); + tkDNN::Region g15(&net, 20, 4, 5, 0.6f); + + // Load input + value_type *data; + value_type *input_h; + readBinaryFile(input_bin, dim.tot(), &input_h, &data); + + //convert network to tensorRT + tkDNN::NetworkRT netRT(&net); + + value_type *out_data, *out_data2; // cudnn output, tensorRT output + + tkDNN::dataDim_t dim1 = dim; //input dim + std::cout<<"\n==== CUDNN inference =======\n"; { + dim1.print(); + TIMER_START + out_data = net.infer(dim1, data); + TIMER_STOP + dim1.print(); + } + + tkDNN::dataDim_t dim2 = dim; + std::cout<<"\n==== TENSORRT inference ====\n"; { + dim2.print(); + TIMER_START + out_data2 = netRT.infer(dim2, data); + TIMER_STOP + dim2.print(); + } + + std::cout<<"\n======= CHECK RESULT =======\n"; + value_type *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; +}