haotian/duan8
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Attach the decode kernel to the same CUDA stream

Browse Source 
fix to load engine file in DeepStream 6.1.1
This commit is contained in:
LiberiFatali 2022-10-08 21:41:35 +07:00 committed by GitHub
parent dfc437daa0
commit 569cc4a438
1 changed files with 2 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
retinaface/decode.cu
View File

@ -176,14 +176,14 @@ namespace nvinfer1
totalCount += decodeplugin::INPUT_H / 16 * decodeplugin::INPUT_W / 16 * 2 * sizeof(decodeplugin::Detection) / sizeof(float); totalCount += decodeplugin::INPUT_H / 16 * decodeplugin::INPUT_W / 16 * 2 * sizeof(decodeplugin::Detection) / sizeof(float);
totalCount += decodeplugin::INPUT_H / 32 * decodeplugin::INPUT_W / 32 * 2 * sizeof(decodeplugin::Detection) / sizeof(float); totalCount += decodeplugin::INPUT_H / 32 * decodeplugin::INPUT_W / 32 * 2 * sizeof(decodeplugin::Detection) / sizeof(float);
for(int idx = 0 ; idx < batchSize; ++idx) { for(int idx = 0 ; idx < batchSize; ++idx) {
cudaMemset(output + idx * totalCount, 0, sizeof(float)); cudaMemsetAsync(output + idx * totalCount, 0, sizeof(float), stream);
} }
for (unsigned int i = 0; i < 3; ++i) for (unsigned int i = 0; i < 3; ++i)
{ {
num_elem = batchSize * decodeplugin::INPUT_H / base_step * decodeplugin::INPUT_W / base_step; num_elem = batchSize * decodeplugin::INPUT_H / base_step * decodeplugin::INPUT_W / base_step;
thread_count = (num_elem < thread_count_) ? num_elem : thread_count_; thread_count = (num_elem < thread_count_) ? num_elem : thread_count_;
CalDetection<<< (num_elem + thread_count - 1) / thread_count, thread_count>>> CalDetection<<< (num_elem + thread_count - 1) / thread_count, thread_count, 0, stream>>>
(inputs[i], output, num_elem, base_step, base_anchor, totalCount); (inputs[i], output, num_elem, base_step, base_anchor, totalCount);
base_step *= 2; base_step *= 2;
base_anchor *= 4; base_anchor *= 4;