diff --git a/README.md b/README.md index 1c50c63..851427e 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/Network.h b/include/tkDNN/Network.h index eb17ea8..0e8fef8 100644 --- a/include/tkDNN/Network.h +++ b/include/tkDNN/Network.h @@ -62,6 +62,7 @@ public: dataDim_t getOutputDim(); bool fp16, dla, int8; + int maxBatchSize; bool dontLoadWeights; std::string fileImgList; std::string fileLabelList; diff --git a/include/tkDNN/NetworkRT.h b/include/tkDNN/NetworkRT.h index 64deac9..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; @@ -74,11 +75,25 @@ public: NetworkRT(Network *net, const char *name); virtual ~NetworkRT(); + int getMaxBatchSize() { + if(engineRT != nullptr) + return engineRT->getMaxBatchSize(); + else + return 0; + } + + int getBuffersN() { + if(engineRT != nullptr) + return engineRT->getNbBindings(); + else + return 0; + } + /** Do inferece */ dnnType* infer(dataDim_t &dim, dnnType* data); - void enqueue(); + void enqueue(int batchSize = 1); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Layer *l); nvinfer1::ILayer* convert_layer(nvinfer1::ITensor *input, Conv2d *l); diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index afcc168..16f07d1 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -41,7 +41,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, const int stride_h, const int stride_w, const int pad_h, const int pad_w, const int dilation_h, const int dilation_w, - const int deformable_group, + const int deformable_group, const int batch_id, const int in_n, const int in_c, const int in_h, const int in_w, const int out_n, const int out_c, const int out_h, const int out_w, const int dst_dim, cudaStream_t stream = cudaStream_t(0)); 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/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index 4804c03..bff6370 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -86,26 +86,26 @@ public: dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); // split conv2d outputs into offset to mask - checkCuda(cudaMemcpy(offset, output_conv, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); - checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); - // kernel sigmoide - activationSIGMOIDForward(mask, mask, chunk_dim); - - // deformable convolution - dcnV2CudaForward(stat, handle, - srcData, data_d, - bias2_d, ones_d1, - offset, mask, - reinterpret_cast(outputs[0]), ones_d2, - kh, kw, - sh, sw, - ph, pw, - 1, 1, - deformableGroup, - i_n, i_c, i_h, i_w, - o_n, o_c, o_h, o_w, - chunk_dim); - + for(int b=0; b(outputs[0]), ones_d2, + kh, kw, + sh, sw, + ph, pw, + 1, 1, + deformableGroup, b, + i_n, i_c, i_h, i_w, + o_n, o_c, o_h, o_w, + chunk_dim); + } return 0; } @@ -185,6 +185,11 @@ public: dnnType * offset; dnnType * mask; dnnType *ones_d2; + // dnnType *input_n; + // dnnType *offset_n; + // dnnType *mask_n; + // dnnType *output_n; + tk::dnn::DeformConv2d *defRT; }; diff --git a/include/tkDNN/pluginsRT/FlattenConcatRT.h b/include/tkDNN/pluginsRT/FlattenConcatRT.h index e6e4eb2..4ebce69 100644 --- a/include/tkDNN/pluginsRT/FlattenConcatRT.h +++ b/include/tkDNN/pluginsRT/FlattenConcatRT.h @@ -47,11 +47,15 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); - checkCuda( cudaMemcpy(dstData, srcData, rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice)); - - float const alpha(1.0); - float const beta(0.0); - checkERROR( cublasSgeam( handle, CUBLAS_OP_T, CUBLAS_OP_N, rows, cols, &alpha, srcData, cols, &beta, srcData, rows, dstData, rows )); + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*rows*cols*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + + checkERROR( cublasSetStream(handle, stream) ); + for(int i=0; in<<" "<c<<" "<h<<" "<w<<" "<stride_H<<" "<stride_W<<" "<winSize<<" "<padding<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, stream); return 0; } diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h index 6331525..f0d127e 100644 --- a/include/tkDNN/pluginsRT/RegionRT.h +++ b/include/tkDNN/pluginsRT/RegionRT.h @@ -50,18 +50,18 @@ 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, + (c*h*w)/num, w*h, 1, w*h, 1, dstData + index, stream); return 0; @@ -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/ReshapeRT.h b/include/tkDNN/pluginsRT/ReshapeRT.h index 36c8c63..97030db 100644 --- a/include/tkDNN/pluginsRT/ReshapeRT.h +++ b/include/tkDNN/pluginsRT/ReshapeRT.h @@ -40,7 +40,7 @@ public: dnnType *srcData = (dnnType*)reinterpret_cast(inputs[0]); dnnType *dstData = reinterpret_cast(outputs[0]); - checkCuda( cudaMemcpy(dstData, srcData, c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice)); + checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); return 0; } diff --git a/include/tkDNN/pluginsRT/RouteRT.h b/include/tkDNN/pluginsRT/RouteRT.h index 9abcd7f..0e94a97 100644 --- a/include/tkDNN/pluginsRT/RouteRT.h +++ b/include/tkDNN/pluginsRT/RouteRT.h @@ -3,6 +3,10 @@ class RouteRT : public IPlugin { + /** + THIS IS NOT USED ANYMORE + */ + public: RouteRT() { } diff --git a/include/tkDNN/pluginsRT/ShortcutRT.h b/include/tkDNN/pluginsRT/ShortcutRT.h index fd2c60d..3eadd3f 100644 --- a/include/tkDNN/pluginsRT/ShortcutRT.h +++ b/include/tkDNN/pluginsRT/ShortcutRT.h @@ -47,7 +47,8 @@ public: dnnType *dstData = reinterpret_cast(outputs[0]); checkCuda( cudaMemcpyAsync(dstData, srcData, batchSize*c*h*w*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - shortcutForward(srcDataBack, dstData, batchSize, c, h, w, 1, batchSize, bc, bh, bw, 1, stream); + for(int b=0; b < batchSize; ++b) + shortcutForward(srcDataBack + b*bc*bh*bw, dstData + b*c*h*w, 1, c, h, w, 1, 1, bc, bh, bw, 1, stream); return 0; } 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/scripts/test_all_tests.sh b/scripts/test_all_tests.sh index d663d01..41b2781 100644 --- a/scripts/test_all_tests.sh +++ b/scripts/test_all_tests.sh @@ -32,6 +32,14 @@ function print_output { out_file=results.log rm $out_file +function test_net { + ./test_$1 &>> $out_file + print_output $? $1 + ./test_rtinference $1*.rt $TKDNN_BATCHSIZE &>> $out_file + print_output $? "batched $1" +} + + modes=( 1 ) # only FP32 # modes=( 1 2 ) # FP32 and FP16 # modes=( 1 2 3 ) # FP32, FP16 and INT8 @@ -57,73 +65,28 @@ do echo -e "${ORANGE}Test INT8${NC}" fi + export TKDNN_BATCHSIZE=2 + echo -e "${ORANGE}Batch $TKDNN_BATCHSIZE ${NC}" + ./test_imuodom &>> $out_file - res_imuodom=$? - print_output $res_imuodom test_imuodom + print_output $? imuodom - ./test_resnet101_cnet &>> $out_file - res_resnet101_cnet=$? - print_output $res_resnet101_cnet test_resnet101_cnet - - ./test_yolo3 &>> $out_file - res_yolo3=$? - print_output $res_yolo3 test_yolo3 - - ./test_yolo3_flir &>> $out_file - res_yolo3_flir=$? - print_output $res_yolo3_flir test_yolo3_flir - - ./test_yolo3_512 &>> $out_file - res_yolo3_512=$? - print_output $res_yolo3_512 test_yolo3_512 - - ./test_yolo3_tiny &>> $out_file - res_yolo3_tiny=$? - print_output $res_yolo3_tiny test_yolo3_tiny - - ./test_csresnext50-panet-spp &>> $out_file - res_csresnext50panetspp=$? - print_output $res_csresnext50panetspp "test_csresnext50-panet-spp" - - ./test_mobilenetv2ssd &>> $out_file - res_mobilenetv2ssd=$? - print_output $res_mobilenetv2ssd test_mobilenetv2ssd - - ./test_yolo3_tiny512 &>> $out_file - res_yolo3_tiny512=$? - print_output $res_yolo3_tiny512 test_yolo3_tiny512 - - ./test_yolo_tiny &>> $out_file - res_yolo_tiny=$? - print_output $res_yolo_tiny test_yolo_tiny - - ./test_mobilenetv2ssd512 &>> $out_file - res_mobilenetv2ssd512=$? - print_output $res_mobilenetv2ssd512 test_mobilenetv2ssd512 - - ./test_mnist &>> $out_file - res_mnist=$? - print_output $res_mnist test_mnist - - ./test_yolo &>> $out_file - res_yolo=$? - print_output $res_yolo test_yolo - - ./test_yolo3_berkeley &>> $out_file - res_yolo3_berkeley=$? - print_output $res_yolo3_berkeley test_yolo3_berkeley - - ./test_yolo_voc &>> $out_file - res_yolo_voc=$? - print_output $res_yolo_voc test_yolo_voc - - ./test_dla34_cnet &>> $out_file - res_dla34_cnet=$? - print_output $res_dla34_cnet test_dla34_cnet - - ./test_yolo3_coco4 &>> $out_file - res_yolo3_coco4=$? - print_output $res_yolo3_coco4 test_yolo3_coco4 + test_net resnet101_cnet + test_net yolo3 + test_net yolo3_flir + test_net yolo3_512 + test_net yolo3_tiny + test_net csresnext50-panet-spp + test_net mobilenetv2ssd + test_net yolo3_tiny512 + test_net yolo_tiny + test_net mobilenetv2ssd512 + test_net mnist + test_net yolo + test_net yolo3_berkeley + test_net yolo_voc + test_net dla34_cnet + test_net yolo3_coco4 done diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp index 06bba7f..b161a22 100644 --- a/src/DeformConv2d.cpp +++ b/src/DeformConv2d.cpp @@ -102,13 +102,13 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { dcnV2CudaForward(stat, handle, srcData, this->data_d, this->bias2_d, ones_d1, - offset, mask, + offset, mask, dstData, ones_d2, this->kernelH, this->kernelW, this->strideH, this->strideW, this->paddingH, this->paddingW, 1, 1, - this->deformableGroup, + this->deformableGroup, 0, //batch_id for cudnn is set to 0 (no batch) preconv->input_dim.n, preconv->input_dim.c, preconv->input_dim.h, preconv->input_dim.w, this->output_dim.n, this->output_dim.c, this->output_dim.h, this->output_dim.w, chunk_dim); diff --git a/src/Network.cpp b/src/Network.cpp index 885ec1a..f199fc2 100644 --- a/src/Network.cpp +++ b/src/Network.cpp @@ -34,6 +34,10 @@ Network::Network(dataDim_t input_dim) { int8 = true; } } + maxBatchSize = 1; + if(const char* env_p = std::getenv("TKDNN_BATCHSIZE")) { + maxBatchSize = atoi(env_p); + } if(const char* env_p = std::getenv("TKDNN_CALIB_IMG_PATH")) fileImgList = env_p; diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 6a8dabe..dc02163 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -58,7 +58,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { dataDim_t dim = net->layers[0]->input_dim; dtRT = DataType::kFLOAT; - builderRT->setMaxBatchSize(1); + builderRT->setMaxBatchSize(net->maxBatchSize); builderRT->setMaxWorkspaceSize(1 << 30); if(net->fp16 && builderRT->platformHasFastFp16()) { @@ -133,6 +133,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { input->setName("out"); networkRT->markOutput(*input); + std::cout<<"Selected maxBatchSize: "<getMaxBatchSize()<<"\n"; std::cout<<"Building tensorRT cuda engine...\n"; #if NV_TENSORRT_MAJOR >= 6 engineRT = builderRT->buildEngineWithConfig(*networkRT, *configRT); @@ -181,9 +182,11 @@ NetworkRT::NetworkRT(Network *net, const char *name) { // create GPU buffers and a stream for(int i=0; igetNbBindings(); i++) { Dims dim = engineRT->getBindingDimensions(i); - checkCuda(cudaMalloc(&buffersRT[i], dim.d[0]*dim.d[1]*dim.d[2]*sizeof(dnnType))); + 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, output_dim.tot()*sizeof(dnnType))); + checkCuda(cudaMalloc(&output, engineRT->getMaxBatchSize()*output_dim.tot()*sizeof(dnnType))); checkCuda(cudaStreamCreate(&stream)); } @@ -192,19 +195,24 @@ NetworkRT::~NetworkRT() { } dnnType* NetworkRT::infer(dataDim_t &dim, dnnType* data) { + int batches = dim.n; + if(batches > getMaxBatchSize()) { + FatalError("input batch size too large"); + } - checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, input_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - contextRT->enqueue(1, buffersRT, stream, nullptr); - checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); - cudaStreamSynchronize(stream); + checkCuda(cudaMemcpyAsync(buffersRT[buf_input_idx], data, batches*input_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + contextRT->enqueue(batches, buffersRT, stream, nullptr); + checkCuda(cudaMemcpyAsync(output, buffersRT[buf_output_idx], batches*output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice, stream)); + checkCuda(cudaStreamSynchronize(stream)); dim = output_dim; + dim.n = batches; return output; } -void NetworkRT::enqueue() { - contextRT->enqueue(1, buffersRT, stream, nullptr); +void NetworkRT::enqueue(int batchSize) { + contextRT->enqueue(batchSize, buffersRT, stream, nullptr); } ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { @@ -320,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); @@ -535,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); @@ -562,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); @@ -622,7 +630,7 @@ IPlugin* PluginFactory::createPlugin(const char* layerName, const void* serialDa const char * buf = reinterpret_cast(serialData); std::string name(layerName); - std::cout< /* srand, rand */ + int main(int argc, char *argv[]) { if(argc < 2 || !fileExist(argv[1])) FatalError("unable to read serialRT file"); + int BATCH_SIZE = 1; + if(argc >2) + BATCH_SIZE = atoi(argv[2]); + //always same test srand (0); //convert network to tensorRT tk::dnn::NetworkRT netRT(NULL, argv[1]); + + tk::dnn::dataDim_t idim = netRT.input_dim; + tk::dnn::dataDim_t odim = netRT.output_dim; + idim.n = BATCH_SIZE; + odim.n = BATCH_SIZE; + dnnType *input = new float[idim.tot()]; + dnnType *output = new float[odim.tot()]; + dnnType *input_d; + checkCuda( cudaMalloc(&input_d, idim.tot()*sizeof(dnnType))); - dnnType *input = new float[netRT.input_dim.tot()]; - dnnType *output = new float[netRT.input_dim.tot()]; - + int ret_tensorrt = 0; + std::cout<<"Testing with batchsize: "<