From a2db98670a90532c4b797d1acf7c8fa5193e5339 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Wed, 4 Dec 2019 17:01:40 +0000 Subject: [PATCH] Add Yolov3 (COCO80) and Yolov3-tiny (COCO80), TensorRT for tiny not working Signed-off-by: Micaela Verucchi --- CMakeLists.txt | 6 + src/NetworkRT.cpp | 7 +- src/Pooling.cpp | 22 +- src/Yolo.cpp | 8 +- src/utils.cpp | 2 + tests/yolo3/yolo3.cpp | 92 ++++ tests/yolo3/yolov3.cfg | 789 +++++++++++++++++++++++++++++++ tests/yolo3_tiny/yolo3_tiny.cpp | 122 +++++ tests/yolo3_tiny/yolov3-tiny.cfg | 182 +++++++ 9 files changed, 1222 insertions(+), 8 deletions(-) create mode 100644 tests/yolo3/yolo3.cpp create mode 100644 tests/yolo3/yolov3.cfg create mode 100644 tests/yolo3_tiny/yolo3_tiny.cpp create mode 100644 tests/yolo3_tiny/yolov3-tiny.cfg diff --git a/CMakeLists.txt b/CMakeLists.txt index 7c5a301..f50b96a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -86,6 +86,12 @@ target_link_libraries(test_yolo_berkeley tkDNN) add_executable(test_yolo3_coco4 tests/yolo3_coco4/yolo3_coco4.cpp) target_link_libraries(test_yolo3_coco4 tkDNN) +add_executable(test_yolo3 tests/yolo3/yolo3.cpp) +target_link_libraries(test_yolo3 tkDNN) + +add_executable(test_yolo3_tiny tests/yolo3_tiny/yolo3_tiny.cpp) +target_link_libraries(test_yolo3_tiny tkDNN) + add_executable(test_yolo3_berkeley tests/yolo3_berkeley/yolo3_berkeley.cpp) target_link_libraries(test_yolo3_berkeley tkDNN) diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index a1ed830..7cfa816 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -209,7 +209,9 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Dense *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { - //std::cout<<"convert conv2D\n"; + std::cout<<"convert conv2D\n"; + printf("%d %d %d %d %d\n", l->kernelH, l->kernelW, l->inputs, l->outputs, l->batchnorm); + void *data_b, *bias_b, *power_b, *mean_b, *variance_b, *scales_b; if(dtRT == DataType::kHALF) { @@ -274,7 +276,8 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { } ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { - //std::cout<<"convert Pooling\n"; + std::cout<<"convert Pooling\n"; + // printf("%d %d\n", l->winW, l->winH); PoolingType ptype; if(l->pool_mode == tkdnnPoolingMode_t::POOLING_MAX) ptype = PoolingType::kMAX; diff --git a/src/Pooling.cpp b/src/Pooling.cpp index 98bbd42..ce84926 100644 --- a/src/Pooling.cpp +++ b/src/Pooling.cpp @@ -26,6 +26,9 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, int w = input_dim.w; int l = input_dim.l; + printf("before: %d %d\n", h, w); + + poolOn3d = false; if(l > 1) { @@ -38,6 +41,8 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, n = l; } + + checkCUDNN( cudnnSetPooling2dDescriptor(poolingDesc, cudnnPoolingMode_t(pool_mode), CUDNN_NOT_PROPAGATE_NAN, winH, winW, paddingH, paddingW, strideH, strideW) ); @@ -45,12 +50,23 @@ Pooling::Pooling( Network *net, int winH, int winW, int strideH, int strideW, net->tensorFormat, net->dataType, n, c, h, w) ); //get out dim - checkCUDNN( cudnnGetPooling2dForwardOutputDim(poolingDesc, srcTensorDesc, &n, &c, &h, &w)); - //h = (h + winH*this->paddingH)/strideH; - //w = (w + winW*this->paddingW)/strideW; + // checkCUDNN( cudnnGetPooling2dForwardOutputDim(poolingDesc, srcTensorDesc, &n, &c, &h, &w)); + + //compute w and h as in darknet + + int padH = paddingH == 0? winH -1 : paddingH; + int padW = paddingW == 0? winW -1 : paddingW; + + h = (h + padH - winH)/strideH +1; + w = (w + padW - winW)/strideW +1; + + // h = (h + winH*this->paddingH)/strideH; + // w = (w + winW*this->paddingW)/strideW; checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc, net->tensorFormat, net->dataType, n, c, h, w) ); + + printf("after: %d %d\n", h, w); output_dim.n = n; output_dim.c = c; diff --git a/src/Yolo.cpp b/src/Yolo.cpp index babef38..e01ed57 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -15,14 +15,16 @@ Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights) : Layer(net) { this->classes = classes; - this->num = num; + this->num = 3; // load anchors if(fname_weights != "") { int seek = 0; - readBinaryFile(fname_weights, num, &mask_h, &mask_d, seek); - seek += num; + readBinaryFile(fname_weights, 3, &mask_h, &mask_d, seek); + seek += 3; readBinaryFile(fname_weights, 3*num*2, &bias_h, &bias_d, seek); + for(int i=0; i<3*num*2; i++) + printf("%f\n", bias_h[i]); } // init default classes name diff --git a/src/utils.cpp b/src/utils.cpp index 4dac886..250cadc 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -38,7 +38,9 @@ void readBinaryFile(std::string fname, int size, dnnType** data_h, dnnType** dat dataFile.seekg(seek * sizeof(dnnType), dataFile.cur); } + // printf("data_h %d size_b %d\n", *data_h,size_b); if (!dataFile.read((char *) *data_h, size_b)) { + error_s << "Error reading file " << fname; FatalError(error_s.str()); } diff --git a/tests/yolo3/yolo3.cpp b/tests/yolo3/yolo3.cpp new file mode 100644 index 0000000..a3772e5 --- /dev/null +++ b/tests/yolo3/yolo3.cpp @@ -0,0 +1,92 @@ +#include +#include +#include "tkdnn.h" + +int main() { + + // Network layout + tk::dnn::dataDim_t dim(1, 3, 416, 416, 1); + tk::dnn::Network net(dim); + + // create yolo3 model + std::string bin_path = "../tests/yolo3"; + int classes = 80; + tk::dnn::Yolo *yolo [3]; + #include "models/Yolo3.h" + + // fill classes names + for(int i=0; i<3; i++) { + yolo[i]->classesNames = {"person" , "bicycle" , "car" , "motorbike" , "aeroplane" , "bus" , "train" , "truck" , "boat" , "traffic light" , "fire hydrant" , "stop sign" , "parking meter" , "bench" , "bird" , "cat" , "dog" , "horse" , "sheep" , "cow" , "elephant" , "bear" , "zebra" , "giraffe" , "backpack" , "umbrella" , "handbag" , "tie" , "suitcase" , "frisbee" , "skis" , "snowboard" , "sports ball" , "kite" , "baseball bat" , "baseball glove" , "skateboard" , "surfboard" , "tennis racket" , "bottle" , "wine glass" , "cup" , "fork" , "knife" , "spoon" , "bowl" , "banana" , "apple" , "sandwich" , "orange" , "broccoli" , "carrot" , "hot dog" , "pizza" , "donut" , "cake" , "chair" , "sofa" , "pottedplant" , "bed" , "diningtable" , "toilet" , "tvmonitor" , "laptop" , "mouse" , "remote" , "keyboard" , "cell phone" , "microwave" , "oven" , "toaster" , "sink" , "refrigerator" , "book" , "clock" , "vase" , "scissors" , "teddy bear" , "hair drier" , "toothbrush"}; + } + + // Load input + dnnType *data; + dnnType *input_h; + readBinaryFile(input_bin, dim.tot(), &input_h, &data); + + //print network model + net.print(); + + //convert network to tensorRT + tk::dnn::NetworkRT netRT(&net, "yolo3.rt"); + + // the network have 3 outputs + tk::dnn::dataDim_t out_dim[3]; + for(int i=0; i<3; i++) out_dim[i] = yolo[i]->output_dim; + dnnType *cudnn_out[3], *rt_out[3]; + + tk::dnn::dataDim_t dim1 = dim; //input dim + printCenteredTitle(" CUDNN inference ", '=', 30); { + dim1.print(); + TIMER_START + net.infer(dim1, data); + TIMER_STOP + dim1.print(); + } + for(int i=0; i<3; i++) cudnn_out[i] = yolo[i]->dstData; + + printCenteredTitle(" compute detections ", '=', 30); + TIMER_START + int ndets = 0; + tk::dnn::Yolo::detection *dets = tk::dnn::Yolo::allocateDetections(tk::dnn::Yolo::MAX_DETECTIONS, classes); + for(int i=0; i<3; i++) yolo[i]->computeDetections(dets, ndets, net.input_dim.w, net.input_dim.h, 0.5); + tk::dnn::Yolo::mergeDetections(dets, ndets, classes); + + for(int j=0; j 0) + cl = c; + } + std::cout< +#include "tkdnn.h" + +const char *input_bin = "../tests/yolo3_tiny/layers/input.bin"; +const char *c0_bin = "../tests/yolo3_tiny/layers/c0.bin"; +const char *c2_bin = "../tests/yolo3_tiny/layers/c2.bin"; +const char *c4_bin = "../tests/yolo3_tiny/layers/c4.bin"; +const char *c6_bin = "../tests/yolo3_tiny/layers/c6.bin"; +const char *c8_bin = "../tests/yolo3_tiny/layers/c8.bin"; +const char *c10_bin = "../tests/yolo3_tiny/layers/c10.bin"; +const char *c12_bin = "../tests/yolo3_tiny/layers/c12.bin"; +const char *c13_bin = "../tests/yolo3_tiny/layers/c13.bin"; +const char *c14_bin = "../tests/yolo3_tiny/layers/c14.bin"; +const char *c15_bin = "../tests/yolo3_tiny/layers/c15.bin"; +const char *c18_bin = "../tests/yolo3_tiny/layers/c18.bin"; +const char *c21_bin = "../tests/yolo3_tiny/layers/c21.bin"; +const char *c22_bin = "../tests/yolo3_tiny/layers/c22.bin"; +const char *g16_bin = "../tests/yolo3_tiny/layers/g16.bin"; +const char *g23_bin = "../tests/yolo3_tiny/layers/g23.bin"; +// const char *output_bin = "../tests/yolo3_tiny/layers/output.bin"; + +const char *output_bin = "../tests/yolo3_tiny/debug/layer23_out.bin"; + +int main() { + + int classes = 80; + + // Network layout + tk::dnn::dataDim_t dim(1, 3, 416, 416, 1); + tk::dnn::Network net(dim); + + + tk::dnn::Conv2d c0 (&net, 16, 3, 3, 1, 1, 1, 1, c0_bin, true); + tk::dnn::Activation a0 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Pooling p1 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX); + + tk::dnn::Conv2d c2 (&net, 32, 3, 3, 1, 1, 1, 1, c2_bin, true); + tk::dnn::Activation a2 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Pooling p3 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX); + + tk::dnn::Conv2d c4 (&net, 64, 3, 3, 1, 1, 1, 1, c4_bin, true); + tk::dnn::Activation a4 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Pooling p5 (&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX); + + tk::dnn::Conv2d c6 (&net, 128, 3, 3, 1, 1, 1, 1, c6_bin, true); + tk::dnn::Activation a6 (&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Pooling p7(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX); + + tk::dnn::Conv2d c8(&net, 256, 3, 3, 1, 1, 1, 1, c8_bin, true); + tk::dnn::Activation a8(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Pooling p9(&net, 2, 2, 2, 2, tk::dnn::POOLING_MAX); + + tk::dnn::Conv2d c10(&net, 512, 3, 3, 1, 1, 1, 1, c10_bin, true); + tk::dnn::Activation a10(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Pooling p11(&net, 2, 2, 1, 1, tk::dnn::POOLING_MAX); + + tk::dnn::Conv2d c12(&net, 1024, 3, 3, 1, 1, 1, 1, c12_bin, true); + tk::dnn::Activation a12(&net, tk::dnn::ACTIVATION_LEAKY); + + tk::dnn::Conv2d c13(&net, 256, 1, 1, 1, 1, 0, 0, c13_bin, true); + tk::dnn::Activation a13(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c14(&net, 512, 3, 3, 1, 1, 1, 1, c14_bin, true); + tk::dnn::Activation a14(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c15(&net, 255, 1, 1, 1, 1, 0, 0, c15_bin, false); + + tk::dnn::Yolo yolo0 (&net, classes, 2, g16_bin); + + tk::dnn::Layer *m17_layers[1] = { &a13 }; + tk::dnn::Route m17 (&net, m17_layers, 1); + tk::dnn::Conv2d c18(&net, 128, 1, 1, 1, 1, 0, 0, c18_bin, true); + tk::dnn::Activation a18(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Upsample u19 (&net, 2); + + tk::dnn::Layer *m20_layers[2] = { &u19, &a8 }; + tk::dnn::Route m20 (&net, m20_layers, 2); + + tk::dnn::Conv2d c21(&net, 256, 3, 3, 1, 1, 1, 1, c21_bin, true); + tk::dnn::Activation a21(&net, tk::dnn::ACTIVATION_LEAKY); + tk::dnn::Conv2d c22(&net, 255, 1, 1, 1, 1, 0, 0, c22_bin, false); + + tk::dnn::Yolo yolo1 (&net, classes, 2, g23_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 + tk::dnn::NetworkRT netRT(&net, "yolo3_tiny.rt"); + + dnnType *out_data, *out_data2; // cudnn output, tensorRT output + + tk::dnn::dataDim_t dim1 = dim; //input dim + printCenteredTitle(" CUDNN inference ", '=', 30); { + dim1.print(); + TIMER_START + out_data = net.infer(dim1, data); + TIMER_STOP + dim1.print(); + } + + // tk::dnn::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; +} diff --git a/tests/yolo3_tiny/yolov3-tiny.cfg b/tests/yolo3_tiny/yolov3-tiny.cfg new file mode 100644 index 0000000..cfca3cf --- /dev/null +++ b/tests/yolo3_tiny/yolov3-tiny.cfg @@ -0,0 +1,182 @@ +[net] +# Testing +batch=1 +subdivisions=1 +# Training +# batch=64 +# subdivisions=2 +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 +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] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + + + +[yolo] +mask = 3,4,5 +anchors = 10,14, 23,27, 37,58, 81,82, 135,169, 344,319 +classes=80 +num=6 +jitter=.3 +ignore_thresh = .7 +truth_thresh = 1 +random=1 + +[route] +layers = -4 + +[convolutional] +batch_normalize=1 +filters=128 +size=1 +stride=1 +pad=1 +activation=leaky + +[upsample] +stride=2 + +[route] +layers = -1, 8 + +[convolutional] +batch_normalize=1 +filters=256 +size=3 +stride=1 +pad=1 +activation=leaky + +[convolutional] +size=1 +stride=1 +pad=1 +filters=255 +activation=linear + +[yolo] +mask = 0,1,2 +anchors = 10,14, 23,27, 37,58, 81,82, 135,169, 344,319 +classes=80 +num=6 +jitter=.3 +ignore_thresh = .7 +truth_thresh = 1 +random=1