From 44b2bce3ff98b407e91bb2abfdb308b683913624 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Thu, 5 Dec 2019 20:43:48 +0000 Subject: [PATCH] Yolo3_tiny CUDNN works, TensorRT doesn't. Add n_masks to Yolo layer. Signed-off-by: Micaela Verucchi --- include/tkDNN/Layer.h | 4 +-- include/tkDNN/pluginsRT/YoloRT.h | 22 +++++++------- src/NetworkRT.cpp | 50 ++++++++++++++++++++++++++------ src/Yolo.cpp | 15 +++++----- tests/yolo3_tiny/yolo3_tiny.cpp | 22 +++++++------- 5 files changed, 74 insertions(+), 39 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 31dc4d2..d001497 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -379,11 +379,11 @@ public: int sort_class; }; - Yolo(Network *net, int classes, int num, std::string fname_weights); + Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks=3); virtual ~Yolo(); virtual layerType_t getLayerType() { return LAYER_YOLO; }; - int classes, num; + int classes, num, n_masks; dnnType *mask_h, *mask_d; //anchors dnnType *bias_h, *bias_d; //anchors std::vector classesNames; diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index dab7c18..a8b3b7d 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -8,16 +8,17 @@ class YoloRT : public IPlugin { public: - YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr) { + YoloRT(int classes, int num, tk::dnn::Yolo *yolo = nullptr, int n_masks=3) { this->classes = classes; this->num = num; + this->n_masks = n_masks; - mask = new dnnType[num]; - bias = new dnnType[num*3*2]; + mask = new dnnType[n_masks]; + bias = new dnnType[num*n_masks*2]; if(yolo != nullptr) { - memcpy(mask, yolo->mask_h, sizeof(dnnType)*num); - memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*3*2); + memcpy(mask, yolo->mask_h, sizeof(dnnType)*n_masks); + memcpy(bias, yolo->bias_h, sizeof(dnnType)*num*n_masks*2); classesNames = yolo->classesNames; } } @@ -60,7 +61,7 @@ public: checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); for (int b = 0; b < batchSize; ++b){ - for(int n = 0; n < num; ++n){ + for(int n = 0; n < n_masks; ++n){ int index = entry_index(b, n*w*h, 0, batchSize); activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); @@ -75,19 +76,20 @@ public: virtual size_t getSerializationSize() override { - return 5*sizeof(int) + num*sizeof(dnnType) + num*3*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char); + return 6*sizeof(int) + n_masks*sizeof(dnnType) + num*n_masks*2*sizeof(dnnType) + YOLORT_CLASSNAME_W*classes*sizeof(char); } virtual void serialize(void* buffer) override { char *buf = reinterpret_cast(buffer); tk::dnn::writeBUF(buf, classes); tk::dnn::writeBUF(buf, num); + tk::dnn::writeBUF(buf, n_masks); tk::dnn::writeBUF(buf, c); tk::dnn::writeBUF(buf, h); tk::dnn::writeBUF(buf, w); - for(int i=0; i classesNames; dnnType *mask; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 7cfa816..4500be5 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -15,9 +15,9 @@ using namespace nvinfer1; // Logger for info/warning/errors class Logger : public ILogger { void log(Severity severity, const char* msg) override { -#ifdef DEBUG +// #ifdef DEBUG std::cout <<"TENSORRT LOG: "<< msg << std::endl; -#endif +// #endif } } loggerRT; @@ -253,14 +253,18 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { lRTconv->setStride(DimsHW{l->strideH, l->strideW}); lRTconv->setPadding(DimsHW{l->paddingH, l->paddingW}); lRT = (ILayer*) lRTconv; + } + checkNULL(lRT); if(l->batchnorm) { Weights power{dtRT, power_b, l->outputs}; Weights shift{dtRT, mean_b, l->outputs}; Weights scale{dtRT, variance_b, l->outputs}; + std::cout<getNbOutputs()<addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, shift, scale, power); + checkNULL(lRT2); Weights shift2{dtRT, bias_b, l->outputs}; @@ -277,7 +281,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { std::cout<<"convert Pooling\n"; - // printf("%d %d\n", l->winW, l->winH); + printf("%d %d %d %d %d %d %d %d %d %d %d %d (layer)\n", l->input_dim.h, l->input_dim.w, l->output_dim.h, l->output_dim.w, l->winW, l->winH, l->strideH, l->strideW, l->paddingH, l->paddingW, l->pool_mode, tkdnnPoolingMode_t::POOLING_MAX) ; PoolingType ptype; if(l->pool_mode == tkdnnPoolingMode_t::POOLING_MAX) ptype = PoolingType::kMAX; @@ -287,8 +291,27 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Pooling *l) { IPoolingLayer *lRT = networkRT->addPooling(*input, ptype, DimsHW{l->winH, l->winW}); checkNULL(lRT); - lRT->setStride(DimsHW{l->strideH, l->strideW}); - lRT->setPadding(DimsHW{l->paddingH, l->paddingW}); + + // if (l->input_dim.h == 13 && l->output_dim.h == 13) + // { + // lRT->setPadding(DimsHW{7, 7}); + // lRT->setStride(DimsHW{2, 2}); + // } + // else + // { + lRT->setPadding(DimsHW{l->paddingH, l->paddingW}); + lRT->setStride(DimsHW{l->strideH, l->strideW}); + // } + + // IResizeLayer *lRT = networkRT->addResize(*lRT->getOutput(0)); + // checkNULL(lRT); + // lRT->setOutputDimensions(l->output_dim); + + ITensor *t = lRT->getOutput(0); + for(int j=0; jgetDimensions().nbDims; j++) { + std::cout<getDimensions().d[j]<<" "; + } + std::cout<<" (TensorRT)\n"; return lRT; } @@ -324,12 +347,19 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Softmax *l) { } ILayer* NetworkRT::convert_layer(ITensor *input, Route *l) { - //std::cout<<"convert route\n"; + std::cout<<"convert route\n"; + + ITensor **tens = new ITensor*[l->layers_n]; for(int i=0; ilayers_n; i++) { tens[i] = tensors[l->layers[i]]; + for(int j=0; jgetDimensions().nbDims; j++) { + std::cout<getDimensions().d[j]<<" "; + } + std::cout<<"\n"; } + IConcatenationLayer *lRT = networkRT->addConcatenation(tens, l->layers_n); //IPlugin *plugin = new RouteRT(); //IPluginLayer *lRT = networkRT->addPlugin(tens, l->layers_n, *plugin); @@ -474,13 +504,15 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa if(name.find("Yolo") == 0) { YoloRT *r = new YoloRT(readBUF(buf), //classes - readBUF(buf)); //num + readBUF(buf), //num + nullptr, + readBUF(buf)); //n_masks r->c = readBUF(buf); r->h = readBUF(buf); r->w = readBUF(buf); - for(int i=0; inum; i++) + for(int i=0; in_masks; i++) r->mask[i] = readBUF(buf); - for(int i=0; i<3*2*r->num; i++) + for(int i=0; in_masks*2*r->num; i++) r->bias[i] = readBUF(buf); // save classes names diff --git a/src/Yolo.cpp b/src/Yolo.cpp index e01ed57..ae4694c 100644 --- a/src/Yolo.cpp +++ b/src/Yolo.cpp @@ -11,19 +11,20 @@ namespace tk { namespace dnn { -Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights) : +Yolo::Yolo(Network *net, int classes, int num, std::string fname_weights, int n_masks) : Layer(net) { this->classes = classes; - this->num = 3; + this->num = num; + this->n_masks = n_masks; // load anchors if(fname_weights != "") { int seek = 0; - 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++) + readBinaryFile(fname_weights, n_masks, &mask_h, &mask_d, seek); + seek += n_masks; + readBinaryFile(fname_weights, n_masks*num*2, &bias_h, &bias_d, seek); + for(int i=0; i