From 3d940a9fa2ca485f35f7d9bb4ed47f61a4ce0a7b Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 21 Apr 2020 19:41:40 +0000 Subject: [PATCH] batch seems ok in yolo3_berkely layers to be checked: DeformableConvRT FlattenConcatRT ReshapeRT RouteRT (dont know why but seems working) --- README.md | 5 +++++ include/tkDNN/NetworkRT.h | 8 +++++++ include/tkDNN/pluginsRT/ActivationLeakyRT.h | 2 +- .../tkDNN/pluginsRT/ActivationReLUCeilingRT.h | 2 +- include/tkDNN/pluginsRT/ActivationSigmoidRT.h | 2 +- .../tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h | 2 +- include/tkDNN/pluginsRT/RegionRT.h | 10 ++++----- include/tkDNN/pluginsRT/YoloRT.h | 8 +++---- src/NetworkRT.cpp | 12 ++++++----- tests/test_rtinference/rtinference.cpp | 21 ++++++++++++++++--- 10 files changed, 51 insertions(+), 21 deletions(-) diff --git a/README.md b/README.md index c8a547c..b8c9770 100644 --- a/README.md +++ b/README.md @@ -163,6 +163,11 @@ N.b. The test will be slower: this is due to the INT8 calibration, which may tak N.b. INT8 calibration requires TensorRT version greater than or equal to 6.0 +### BatchSize bigger than 1 +``` +export TKDNN_BATCHSIZE=2 +``` + ## mAP demo To compute mAP, precision, recall and f1score, run the map_demo. diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 1152234..bcb6b89 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -63,6 +63,7 @@ public: const static int MAX_BUFFERS_RT = 10; void* buffersRT[MAX_BUFFERS_RT]; + dataDim_t buffersDIM[MAX_BUFFERS_RT]; int buf_input_idx, buf_output_idx; dataDim_t input_dim, output_dim; @@ -81,6 +82,13 @@ public: return 0; } + int getBuffersN() { + if(engineRT != nullptr) + return engineRT->getNbBindings(); + else + return 0; + } + /** Do inferece */ diff --git a/include/tkDNN/pluginsRT/ActivationLeakyRT.h b/include/tkDNN/pluginsRT/ActivationLeakyRT.h index 1957e7c..d3f66fb 100644 --- a/include/tkDNN/pluginsRT/ActivationLeakyRT.h +++ b/include/tkDNN/pluginsRT/ActivationLeakyRT.h @@ -42,7 +42,7 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { activationLEAKYForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), size, stream); + reinterpret_cast(outputs[0]), batchSize*size, stream); return 0; } diff --git a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h index 0074520..286f22e 100644 --- a/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h +++ b/include/tkDNN/pluginsRT/ActivationReLUCeilingRT.h @@ -41,7 +41,7 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { activationReLUCeilingForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), size, ceiling, stream); + reinterpret_cast(outputs[0]), batchSize*size, ceiling, stream); return 0; } diff --git a/include/tkDNN/pluginsRT/ActivationSigmoidRT.h b/include/tkDNN/pluginsRT/ActivationSigmoidRT.h index 4435f08..1d47136 100644 --- a/include/tkDNN/pluginsRT/ActivationSigmoidRT.h +++ b/include/tkDNN/pluginsRT/ActivationSigmoidRT.h @@ -42,7 +42,7 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { activationSIGMOIDForward((dnnType*)reinterpret_cast(inputs[0]), - reinterpret_cast(outputs[0]), size, stream); + reinterpret_cast(outputs[0]), batchSize*size, stream); return 0; } diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h index 2efa6e2..74f88cd 100644 --- a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h +++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h @@ -45,7 +45,7 @@ public: //std::cout<n<<" "<c<<" "<h<<" "<w<<" "<stride_H<<" "<stride_W<<" "<winSize<<" "<padding<(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); - MaxPoolingForward(srcData, dstData, this->n, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding); + MaxPoolingForward(srcData, dstData, batchSize, this->c, this->h, this->w, this->stride_H, this->stride_W, this->winSize, this->padding); return 0; } diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h index 6331525..176d676 100644 --- a/include/tkDNN/pluginsRT/RegionRT.h +++ b/include/tkDNN/pluginsRT/RegionRT.h @@ -50,16 +50,16 @@ public: for (int b = 0; b < batchSize; ++b){ for(int n = 0; n < num; ++n){ - int index = entry_index(b, n*w*h, 0, batchSize); + int index = entry_index(b, n*w*h, 0); activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); - index = entry_index(b, n*w*h, coords, batchSize); + index = entry_index(b, n*w*h, coords); activationLOGISTICForward(srcData + index, dstData + index, w*h, stream); } } //softmax start - int index = entry_index(0, 0, coords + 1, batchSize); + int index = entry_index(0, 0, coords + 1); softmaxForward( srcData + index, classes, batchSize*num, (batchSize*c*h*w)/num, w*h, 1, w*h, 1, dstData + index, stream); @@ -85,10 +85,10 @@ public: int c, h, w; int classes, coords, num; - int entry_index(int batch, int location, int entry, int batchSize) { + int entry_index(int batch, int location, int entry) { int n = location / (w*h); int loc = location % (w*h); - return batch*c*h*w*batchSize + n*w*h*(coords+classes+1) + entry*w*h + loc; + return batch*c*h*w + n*w*h*(coords+classes+1) + entry*w*h + loc; } }; diff --git a/include/tkDNN/pluginsRT/YoloRT.h b/include/tkDNN/pluginsRT/YoloRT.h index a8b3b7d..66b590b 100644 --- a/include/tkDNN/pluginsRT/YoloRT.h +++ b/include/tkDNN/pluginsRT/YoloRT.h @@ -62,10 +62,10 @@ public: for (int b = 0; b < batchSize; ++b){ for(int n = 0; n < n_masks; ++n){ - int index = entry_index(b, n*w*h, 0, batchSize); + int index = entry_index(b, n*w*h, 0); activationLOGISTICForward(srcData + index, dstData + index, 2*w*h, stream); - index = entry_index(b, n*w*h, 4, batchSize); + index = entry_index(b, n*w*h, 4); activationLOGISTICForward(srcData + index, dstData + index, (1+classes)*w*h, stream); } } @@ -109,10 +109,10 @@ public: dnnType *mask; dnnType *bias; - int entry_index(int batch, int location, int entry, int batchSize) { + int entry_index(int batch, int location, int entry) { int n = location / (w*h); int loc = location % (w*h); - return batch*c*h*w*batchSize + n*w*h*(4+classes+1) + entry*w*h + loc; + return batch*c*h*w + n*w*h*(4+classes+1) + entry*w*h + loc; } }; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 6ed6efe..dc02163 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -182,6 +182,8 @@ NetworkRT::NetworkRT(Network *net, const char *name) { // create GPU buffers and a stream for(int i=0; igetNbBindings(); i++) { Dims dim = engineRT->getBindingDimensions(i); + buffersDIM[i] = dataDim_t(1, dim.d[0], dim.d[1], dim.d[2]); + std::cout<<"RtBuffer "<getMaxBatchSize()*dim.d[0]*dim.d[1]*dim.d[2]*sizeof(dnnType))); } checkCuda(cudaMalloc(&output, engineRT->getMaxBatchSize()*output_dim.tot()*sizeof(dnnType))); @@ -326,7 +328,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { lRT = (ILayer*) lRTconv; Dims d = lRTconv->getOutput(0)->getDimensions(); - std::cout<<"DECONV: "<preconv); checkNULL(preconv); @@ -541,7 +543,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { inputs[0] = input; inputs[1] = preconv->getOutput(0); - std::cout<<"New plugin DEFORMABLE\n"; + //std::cout<<"New plugin DEFORMABLE\n"; IPlugin *plugin = new DeformableConvRT(l->chunk_dim, l->kernelH, l->kernelW, l->strideH, l->strideW, l->paddingH, l->paddingW, l->deformableGroup, l->input_dim.n, l->input_dim.c, l->input_dim.h, l->input_dim.w, l->output_dim.n, l->output_dim.c, l->output_dim.h, l->output_dim.w, l); @@ -568,7 +570,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { Weights power{dtRT, power_b, l->outputs}; Weights shift{dtRT, mean_b, l->outputs}; Weights scale{dtRT, variance_b, l->outputs}; - std::cout<getNbOutputs()<getNbOutputs()<addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, shift, scale, power); @@ -628,7 +630,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa const char * buf = reinterpret_cast(serialData); std::string name(layerName); - std::cout<