From 6792b82f4b024389f75fa2283c47d9b15ca9cec2 Mon Sep 17 00:00:00 2001 From: qiuyunzhe Date: Thu, 18 Jun 2020 16:52:08 -0700 Subject: [PATCH 1/3] prepare anchor during initialization --- yolov4/yololayer.cu | 25 +++++++++++++++++++------ yolov4/yololayer.h | 1 + 2 files changed, 20 insertions(+), 6 deletions(-) diff --git a/yolov4/yololayer.cu b/yolov4/yololayer.cu index fd0af16..bbab043 100644 --- a/yolov4/yololayer.cu +++ b/yolov4/yololayer.cu @@ -13,6 +13,15 @@ namespace nvinfer1 mYoloKernel.push_back(yolo3); mKernelCount = mYoloKernel.size(); + + CUDA_CHECK(cudaMallocHost(&mAnchor, mKernelCount * sizeof(void*))); + size_t AnchorLen = sizeof(float)* CHECK_COUNT*2; + for(int ii = 0; ii < mKernelCount; ii ++) + { + CUDA_CHECK(cudaMalloc(&mAnchor[ii],AnchorLen)); + const auto& yolo = mYoloKernel[ii]; + CUDA_CHECK(cudaMemcpy(mAnchor[ii], yolo.anchors, AnchorLen, cudaMemcpyHostToDevice)); + } } YoloLayerPlugin::~YoloLayerPlugin() @@ -32,6 +41,15 @@ namespace nvinfer1 memcpy(mYoloKernel.data(),d,kernelSize); d += kernelSize; + CUDA_CHECK(cudaMallocHost(&mAnchor, mKernelCount * sizeof(void*))); + size_t AnchorLen = sizeof(float)* CHECK_COUNT*2; + for(int ii = 0; ii < mKernelCount; ii ++) + { + CUDA_CHECK(cudaMalloc(&mAnchor[ii],AnchorLen)); + const auto& yolo = mYoloKernel[ii]; + CUDA_CHECK(cudaMemcpy(mAnchor[ii], yolo.anchors, AnchorLen, cudaMemcpyHostToDevice)); + } + assert(d == a + length); } @@ -179,9 +197,6 @@ namespace nvinfer1 } void YoloLayerPlugin::forwardGpu(const float *const * inputs, float* output, cudaStream_t stream, int batchSize) { - void* devAnchor; - size_t AnchorLen = sizeof(float)* CHECK_COUNT*2; - CUDA_CHECK(cudaMalloc(&devAnchor,AnchorLen)); int outputElem = 1 + MAX_OUTPUT_BBOX_COUNT * sizeof(Detection) / sizeof(float); @@ -195,12 +210,10 @@ namespace nvinfer1 numElem = yolo.width*yolo.height*batchSize; if (numElem < mThreadCount) mThreadCount = numElem; - CUDA_CHECK(cudaMemcpy(devAnchor, yolo.anchors, AnchorLen, cudaMemcpyHostToDevice)); CalDetection<<< (yolo.width*yolo.height*batchSize + mThreadCount - 1) / mThreadCount, mThreadCount>>> - (inputs[i],output, numElem, yolo.width, yolo.height, (float *)devAnchor, mClassCount ,outputElem); + (inputs[i],output, numElem, yolo.width, yolo.height, (float *)mAnchor[i], mClassCount ,outputElem); } - CUDA_CHECK(cudaFree(devAnchor)); } diff --git a/yolov4/yololayer.h b/yolov4/yololayer.h index 5074636..af30e9c 100644 --- a/yolov4/yololayer.h +++ b/yolov4/yololayer.h @@ -116,6 +116,7 @@ namespace nvinfer1 int mKernelCount; std::vector mYoloKernel; int mThreadCount = 256; + void** mAnchor; const char* mPluginNamespace; }; From 65ee683c2802df977035740196be1695710e0ea1 Mon Sep 17 00:00:00 2001 From: qiuyunzhe Date: Thu, 18 Jun 2020 17:08:35 -0700 Subject: [PATCH 2/3] init --- tutorials/multi_GPU_processing.md | 5 +++++ 1 file changed, 5 insertions(+) create mode 100644 tutorials/multi_GPU_processing.md diff --git a/tutorials/multi_GPU_processing.md b/tutorials/multi_GPU_processing.md new file mode 100644 index 0000000..ed053d6 --- /dev/null +++ b/tutorials/multi_GPU_processing.md @@ -0,0 +1,5 @@ +# How to Implement Multi-GPU Processing + +Maybe you are hoping to take advantage of multiple GPU to accelerate inference speed. This is a tutrial to help you deal with it! + + From 60603fa669f54589054d13ead86ff80162bab34e Mon Sep 17 00:00:00 2001 From: qiuyunzhe Date: Thu, 18 Jun 2020 20:46:25 -0700 Subject: [PATCH 3/3] Update multi_GPU_processing.md --- tutorials/multi_GPU_processing.md | 27 ++++++++++++++++++++++++++- 1 file changed, 26 insertions(+), 1 deletion(-) diff --git a/tutorials/multi_GPU_processing.md b/tutorials/multi_GPU_processing.md index ed053d6..b4af794 100644 --- a/tutorials/multi_GPU_processing.md +++ b/tutorials/multi_GPU_processing.md @@ -1,5 +1,30 @@ # How to Implement Multi-GPU Processing -Maybe you are hoping to take advantage of multiple GPU to accelerate inference speed. This is a tutrial to help you deal with it! +Maybe you hope to take advantage of multiple GPU to make inference even faster. Here are few tips to help you deal with it! Take **YOLO V4** as an example. +## 1. Make custom plugin (i.e. YOLO layer and Mish layer for YOLO V4) running asynchronically. +To do this, we need to use CudaStream parameter in the kernels of all custom layers and use asynchronous functions. +For example, in function ` forwardGpu()` of **yololayer.cu**, you need to do the following changes to make sure that the engine will be running on a specific CudaStream. + + 1) Change `cudaMemset(output + idx*outputElem, 0, sizeof(float))` to `cudaMemsetAsync(output + idx*outputElem, 0, sizeof(float), stream)` + 2) Change `CalDetection<<< (yolo.width*yolo.height*batchSize + mThreadCount - 1) / mThreadCount, mThreadCount>>>(inputs[i],output, numElem, yolo.width, yolo.height, (float *)mAnchor[i], mClassCount ,outputElem)` to `CalDetection<<< (yolo.width*yolo.height*batchSize + mThreadCount - 1) / mThreadCount, mThreadCount, 0, stream>>>(inputs[i],output, numElem, yolo.width, yolo.height, (float *)mAnchor[i], mClassCount ,outputElem)` + + ## 2. Create an engine for each device you want to use. + + Maybe it is a good idea to create a struct to store the engine, context and buffer for each device individually. For example, + ``` + struct Plan{ + IRuntime* runtime; + ICudaEngine* engine; + IExecutionContext* context; + void buffers[2]; + cudaStream_t stream; + }; + ``` + And then use `cudaSetDevice()` to make each engine you create running on specific device. Moreover, to maximize performance, make sure that the engine file you are using to deserialize is the one tensor RT optimized for this device. + + ## 3. Use function wisely + Here are some knowledge I learned when trying to parallelize the inference. + 1) Do not use synchronized function , like `cudaFree()`, during inference. + 2) Using `cudaMallocHost()` instead of `malloc()` when allocating memory on the host side.