From a26ef98d2d201ce6bed158c8340546d1bd0b0908 Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Wed, 30 Aug 2017 09:12:46 +0000 Subject: [PATCH] yolo alternatives --- .gitignore | 1 + CMakeLists.txt | 12 +- demo/detection/detection.cpp | 65 +++-- src/Region.cpp | 3 +- tests/yolo/yolo.cfg | 258 ++++++++++++++++++ tests/yolo_224/yolo_224.cfg | 258 ++++++++++++++++++ tests/yolo_224/yolo_224.cpp | 150 ++++++++++ tests/yolo_relu/yolo_relu.cfg | 258 ++++++++++++++++++ tests/yolo_relu/yolo_relu.cpp | 150 ++++++++++ tests/yolo_tiny/tiny-yolo.cfg | 139 ++++++++++ .../yolo-tiny.cpp => yolo_tiny/yolo_tiny.cpp} | 28 +- 11 files changed, 1288 insertions(+), 34 deletions(-) create mode 100644 tests/yolo/yolo.cfg create mode 100644 tests/yolo_224/yolo_224.cfg create mode 100644 tests/yolo_224/yolo_224.cpp create mode 100644 tests/yolo_relu/yolo_relu.cfg create mode 100644 tests/yolo_relu/yolo_relu.cpp create mode 100644 tests/yolo_tiny/tiny-yolo.cfg rename tests/{yolo-tiny/yolo-tiny.cpp => yolo_tiny/yolo_tiny.cpp} (76%) diff --git a/.gitignore b/.gitignore index 38d5bae..446d490 100644 --- a/.gitignore +++ b/.gitignore @@ -7,3 +7,4 @@ build/ *.caffemodel *.h5 *.tar.gz +*.weights diff --git a/CMakeLists.txt b/CMakeLists.txt index bdc30e4..a8f4744 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -57,12 +57,22 @@ 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) +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) diff --git a/demo/detection/detection.cpp b/demo/detection/detection.cpp index 98751af..f0b78c7 100644 --- a/demo/detection/detection.cpp +++ b/demo/detection/detection.cpp @@ -1,6 +1,7 @@ #include #include "tkdnn.h" #include /* srand, rand */ +#include #include #include @@ -84,35 +85,61 @@ void compute_image( cv::Mat imageORIG, } +int print_usage() { + std::cout<<"usage: ./detection net.rt validation_list.txt [-t ] [-s]\n" + <<" -t: set thresh value\n -s: show images as compute\n\n" + <<"> validation_list.txt format: \n" + <<" path/to/image.jpg path/to/label.txt\n" + <<"> label.txt format: \n" + <<" \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[]) { - if(argc < 3) { - std::cout<<"usage: "< validation_list.txt format: \n" - <<" path/to/image.jpg path/to/label.txt\n" - <<"> label.txt format: \n" - <<" \n" - <<" x and y are the box center, " - <<"all values are relative to the image size\n\n"; - return 1; + //params + char *tensor_path = NULL; + char *imageset_path = NULL; + float thresh = 0.3f; + bool show = false; + + //parse params + int c; + while ((c = getopt (argc, argv, "t:s")) != -1) { + switch(c) { + case 't': thresh = atof(optarg); break; + case 's': show = true; break; + case '?': + return print_usage(); + default: return print_usage(); + } } - if(!fileExist(argv[1])) - FatalError("unable to read serialRT file"); + 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, argv[1]); - tkDNN::RegionInterpret rI(netRT.input_dim, netRT.output_dim, 80, 4, 5, 0.3f, reg_bias); + 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(argv[2]); + std::ifstream imageset(imageset_path); if(!imageset.is_open()) FatalError("could not read imageset"); - float mAP = 0; int processed_images; @@ -193,9 +220,11 @@ int main(int argc, char *argv[]) { <<", mAP: "< +#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; +} diff --git a/tests/yolo_relu/yolo_relu.cfg b/tests/yolo_relu/yolo_relu.cfg new file mode 100644 index 0000000..0abab4e --- /dev/null +++ b/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 diff --git a/tests/yolo_relu/yolo_relu.cpp b/tests/yolo_relu/yolo_relu.cpp new file mode 100644 index 0000000..ecc4862 --- /dev/null +++ b/tests/yolo_relu/yolo_relu.cpp @@ -0,0 +1,150 @@ +#include +#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; +} diff --git a/tests/yolo_tiny/tiny-yolo.cfg b/tests/yolo_tiny/tiny-yolo.cfg new file mode 100644 index 0000000..630a209 --- /dev/null +++ b/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 diff --git a/tests/yolo-tiny/yolo-tiny.cpp b/tests/yolo_tiny/yolo_tiny.cpp similarity index 76% rename from tests/yolo-tiny/yolo-tiny.cpp rename to tests/yolo_tiny/yolo_tiny.cpp index 465c601..fc2703d 100644 --- a/tests/yolo-tiny/yolo-tiny.cpp +++ b/tests/yolo_tiny/yolo_tiny.cpp @@ -1,19 +1,19 @@ #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 *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"; +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() { @@ -60,7 +60,7 @@ int main() { net.print(); //convert network to tensorRT - tkDNN::NetworkRT netRT(&net, "yolo-tiny.rt"); + tkDNN::NetworkRT netRT(&net, "yolo_tiny.rt"); dnnType *out_data, *out_data2; // cudnn output, tensorRT output