From 7c81c5a43cdc04550d1c65260c1394551307d21f Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 21 Apr 2020 19:20:44 +0200 Subject: [PATCH 01/10] batch size > 1 --- include/tkDNN/Network.h | 1 + include/tkDNN/NetworkRT.h | 9 ++++- .../tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h | 2 +- src/Network.cpp | 4 +++ src/NetworkRT.cpp | 24 +++++++++----- tests/test_rtinference/rtinference.cpp | 33 ++++++++++++------- 6 files changed, 51 insertions(+), 22 deletions(-) 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..1152234 100644 --- a/include/tkDNN/NetworkRT.h +++ b/include/tkDNN/NetworkRT.h @@ -74,11 +74,18 @@ public: NetworkRT(Network *net, const char *name); virtual ~NetworkRT(); + int getMaxBatchSize() { + if(engineRT != nullptr) + return engineRT->getMaxBatchSize(); + 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/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h index 992042b..2efa6e2 100644 --- a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h +++ b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h @@ -42,7 +42,7 @@ public: virtual int enqueue(int batchSize, const void*const * inputs, void** outputs, void* workspace, cudaStream_t stream) override { - std::cout<n<<" "<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); 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..6ed6efe 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,9 @@ 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))); + checkCuda(cudaMalloc(&buffersRT[i], engineRT->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 +193,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) { diff --git a/tests/test_rtinference/rtinference.cpp b/tests/test_rtinference/rtinference.cpp index f99d003..88aedf0 100644 --- a/tests/test_rtinference/rtinference.cpp +++ b/tests/test_rtinference/rtinference.cpp @@ -2,31 +2,42 @@ #include "tkdnn.h" #include /* 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()]; - + std::cout<<"Testing with batchsize: "< Date: Tue, 21 Apr 2020 19:50:45 +0000 Subject: [PATCH 03/10] RouteRt is not used --- include/tkDNN/pluginsRT/RouteRT.h | 4 ++++ tests/test_rtinference/rtinference.cpp | 2 +- 2 files changed, 5 insertions(+), 1 deletion(-) 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/tests/test_rtinference/rtinference.cpp b/tests/test_rtinference/rtinference.cpp index 8e9fcec..95bb20c 100644 --- a/tests/test_rtinference/rtinference.cpp +++ b/tests/test_rtinference/rtinference.cpp @@ -51,7 +51,7 @@ int main(int argc, char *argv[]) { dnnType *out_d = (dnnType*) netRT.buffersRT[o]; dnnType *out0_d = out_d; dnnType *outI_d = out_d + netRT.buffersDIM[o].tot()*b; - //ret_tensorrt |= checkResult(netRT.buffersDIM[o].tot(), outI_d, out0_d); + ret_tensorrt |= checkResult(netRT.buffersDIM[o].tot(), outI_d, out0_d); } } } From 488887992c2f84eb90ee8e76b604bda772fb7f5a Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 23 Apr 2020 00:54:03 +0200 Subject: [PATCH 04/10] test batch --- scripts/test_all_tests.sh | 93 ++++++++------------------ tests/test_rtinference/rtinference.cpp | 2 +- 2 files changed, 29 insertions(+), 66 deletions(-) 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/tests/test_rtinference/rtinference.cpp b/tests/test_rtinference/rtinference.cpp index 95bb20c..2b521ad 100644 --- a/tests/test_rtinference/rtinference.cpp +++ b/tests/test_rtinference/rtinference.cpp @@ -51,7 +51,7 @@ int main(int argc, char *argv[]) { dnnType *out_d = (dnnType*) netRT.buffersRT[o]; dnnType *out0_d = out_d; dnnType *outI_d = out_d + netRT.buffersDIM[o].tot()*b; - ret_tensorrt |= checkResult(netRT.buffersDIM[o].tot(), outI_d, out0_d); + ret_tensorrt |= checkResult(netRT.buffersDIM[o].tot(), outI_d, out0_d) == 0 ? 0 : ERROR_TENSORRT; } } } From 8cff886ee535eb5ba94f8545bbd48f1b72928d4d Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Thu, 23 Apr 2020 01:18:27 +0200 Subject: [PATCH 05/10] Flatten batch to be checked --- include/tkDNN/pluginsRT/FlattenConcatRT.h | 14 +++++++++----- tests/test_rtinference/rtinference.cpp | 3 ++- 2 files changed, 11 insertions(+), 6 deletions(-) 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; i Date: Thu, 23 Apr 2020 23:51:23 +0200 Subject: [PATCH 06/10] Reshape batch fix --- include/tkDNN/pluginsRT/ReshapeRT.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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; } From 13c9dc662081d4fe25e32ab1612efd084e488a8d Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Fri, 24 Apr 2020 00:08:06 +0200 Subject: [PATCH 07/10] RegionRT softmax fix --- include/tkDNN/pluginsRT/RegionRT.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/tkDNN/pluginsRT/RegionRT.h b/include/tkDNN/pluginsRT/RegionRT.h index 176d676..f0d127e 100644 --- a/include/tkDNN/pluginsRT/RegionRT.h +++ b/include/tkDNN/pluginsRT/RegionRT.h @@ -61,7 +61,7 @@ public: //softmax start 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; From 6e5ab031c450dee87d31a7950924ba828b4e3f01 Mon Sep 17 00:00:00 2001 From: dsapienza <177992@studenti.unimore.it> Date: Sat, 25 Apr 2020 02:47:31 +0200 Subject: [PATCH 08/10] Deformable batch works Signed-off-by: Davide Sapienza --- include/tkDNN/kernels.h | 2 +- include/tkDNN/pluginsRT/DeformableConvRT.h | 45 ++++++++++++---------- src/DeformConv2d.cpp | 4 +- src/kernels/deformable_conv.cu | 12 +++--- 4 files changed, 35 insertions(+), 28 deletions(-) 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/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/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/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index e46c579..592c538 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -241,12 +241,13 @@ 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 chunk_dim, cudaStream_t stream) { // stat and handle have be moved out to preserve 2 - 6 milliseconds every 100. + const int batch = batch_id; const int channels = in_c; const int height = in_h; const int width = in_w; @@ -265,13 +266,14 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, n, m, k, &alpha, ones, k, bias, k, - &beta, output, n); + &beta, output + batch * out_c * out_h * out_w, n); if (stat != CUBLAS_STATUS_SUCCESS) FatalError("CUBLAS initialization failed\n"); modulatedDeformableIm2colCuda(stream, - input, offset, - mask, + input + batch * channels * height * width, + offset,// + b * 2 * int((float)chunk_dim / batch), + mask,// + b * int((float)chunk_dim / batch), 1, channels, height, width, height_out, width_out, deformable_group, columns); // modulatedDeformableIm2colCudaGeneralVersion(stream, @@ -290,7 +292,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, m, k, &alpha, columns, n, weight, k, - &beta, output, n); + &beta, output + batch * out_c * out_h * out_w, n); if (stat != CUBLAS_STATUS_SUCCESS) FatalError("CUBLAS initialization failed\n"); From 89f91f568e847632b831bd2a83931bafb1887806 Mon Sep 17 00:00:00 2001 From: Micaela Verucchi Date: Tue, 28 Apr 2020 14:34:12 +0200 Subject: [PATCH 09/10] Fix shortcutRT plugin, now works with batches, fix yolo3_512 downlaod link Signed-off-by: Micaela Verucchi --- include/tkDNN/pluginsRT/ShortcutRT.h | 3 ++- tests/yolo3_512/yolo3_512.cpp | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) 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/tests/yolo3_512/yolo3_512.cpp b/tests/yolo3_512/yolo3_512.cpp index cb50410..5e796c1 100644 --- a/tests/yolo3_512/yolo3_512.cpp +++ b/tests/yolo3_512/yolo3_512.cpp @@ -10,7 +10,7 @@ int main() { // create yolo3 model std::string bin_path = "yolo3_512"; - downloadWeightsifDoNotExist("yolo3_512/layers/input.bin", bin_path, "https://cloud.hipert.unimore.it/s/e7HfScx77JEHeYb/download"); + downloadWeightsifDoNotExist("yolo3_512/layers/input.bin", bin_path, "https://cloud.hipert.unimore.it/s/RGecMeGLD4cXEWL/download"); int classes = 80; tk::dnn::Yolo *yolo [3]; #include "models/Yolo3.h" From 04bd5d7ff46270c732b76cd9656e22de930e138e Mon Sep 17 00:00:00 2001 From: Francesco Gatti Date: Tue, 28 Apr 2020 14:36:02 +0200 Subject: [PATCH 10/10] stream --- include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h b/include/tkDNN/pluginsRT/MaxPoolingFixedSizeRT.h index 74f88cd..911fca2 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, batchSize, 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; }