diff --git a/README.md b/README.md index c2f017e..457e2a5 100644 --- a/README.md +++ b/README.md @@ -23,6 +23,7 @@ NVIDIA DeepStream SDK 6.2 / 6.1.1 / 6.1 / 6.0.1 / 6.0 configuration for YOLO mod * Models benchmarks * **Support for Darknet YOLO models (YOLOv4, etc) using cfg and weights conversion with GPU post-processing** * **Support for YOLO-NAS, PPYOLOE+, PPYOLOE, DAMO-YOLO, YOLOX, YOLOR, YOLOv8, YOLOv7, YOLOv6 and YOLOv5 using ONNX conversion with GPU post-processing** +* **Add GPU bbox parser (it is slightly slower than CPU bbox parser on V100 GPU tests)** ## @@ -153,7 +154,7 @@ sample = 1920x1080 video - Eval ``` -nms-iou-threshold = 0.6 (Darknet) / 0.65 (YOLOv5, YOLOv6, YOLOv7, YOLOR and YOLOX) / 0.7 (Paddle, YOLO-NAS, YOLOv8 and YOLOv7-u6) +nms-iou-threshold = 0.6 (Darknet) / 0.65 (YOLOv5, YOLOv6, YOLOv7, YOLOR and YOLOX) / 0.7 (Paddle, YOLO-NAS, DAMO-YOLO, YOLOv8 and YOLOv7-u6) pre-cluster-threshold = 0.001 topk = 300 ``` @@ -172,7 +173,11 @@ topk = 300 **NOTE**: ** = The YOLOv4 is trained with the trainvalno5k set, so the mAP is high on val2017 test. -**NOTE**: The V100 GPU decoder seems to max out at 625-635 FPS on DeepStream even using lighter models. +**NOTE**: star = DAMO-YOLO model trained with distillation. + +**NOTE**: The V100 GPU decoder max out at 625-635 FPS on DeepStream even using lighter models. + +**NOTE**: The GPU bbox parser is a bit slower than CPU bbox parser on V100 GPU tests. | DeepStream | Precision | Resolution | IoU=0.5:0.95 | IoU=0.5 | IoU=0.75 | FPS
(without display) | |:------------------:|:---------:|:----------:|:------------:|:-------:|:--------:|:--------------------------:| @@ -184,6 +189,14 @@ topk = 300 | PP-YOLOE+_m | FP16 | 640 | 0.483 | 0.658 | 0.528 | 264.39 | | PP-YOLOE+_s | FP16 | 640 | 0.424 | 0.594 | 0.464 | 476.13 | | PP-YOLOE-s (400) | FP16 | 640 | 0.423 | 0.589 | 0.463 | 461.23 | +| DAMO-YOLO-L star | FP16 | 640 | 0.502 | 0.674 | 0.551 | 176.93 | +| DAMO-YOLO-M star | FP16 | 640 | 0.485 | 0.656 | 0.530 | 242.24 | +| DAMO-YOLO-S star | FP16 | 640 | 0.460 | 0.631 | 0.502 | 385.09 | +| DAMO-YOLO-S | FP16 | 640 | 0.445 | 0.611 | 0.486 | 378.68 | +| DAMO-YOLO-T star | FP16 | 640 | 0.419 | 0.586 | 0.455 | 492.24 | +| DAMO-YOLO-Nl | FP16 | 416 | 0.392 | 0.559 | 0.423 | 483.73 | +| DAMO-YOLO-Nm | FP16 | 416 | 0.371 | 0.532 | 0.402 | 555.94 | +| DAMO-YOLO-Ns | FP16 | 416 | 0.312 | 0.460 | 0.335 | 627.67 | | YOLOX-x | FP16 | 640 | 0.447 | 0.616 | 0.483 | 125.40 | | YOLOX-l | FP16 | 640 | 0.430 | 0.598 | 0.466 | 193.10 | | YOLOX-m | FP16 | 640 | 0.397 | 0.566 | 0.431 | 298.61 | diff --git a/nvdsinfer_custom_impl_Yolo/nvdsparsebbox_Yolo_cuda.cu b/nvdsinfer_custom_impl_Yolo/nvdsparsebbox_Yolo_cuda.cu new file mode 100644 index 0000000..ffff5ee --- /dev/null +++ b/nvdsinfer_custom_impl_Yolo/nvdsparsebbox_Yolo_cuda.cu @@ -0,0 +1,190 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a + * copy of this software and associated documentation files (the "Software"), + * to deal in the Software without restriction, including without limitation + * the rights to use, copy, modify, merge, publish, distribute, sublicense, + * and/or sell copies of the Software, and to permit persons to whom the + * Software is furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL + * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER + * DEALINGS IN THE SOFTWARE. + * + * Edited by Marcos Luciano + * https://www.github.com/marcoslucianops + */ + +#include +#include +#include + +#include "nvdsinfer_custom_impl.h" + +extern "C" bool +NvDsInferParseYolo_cuda(std::vector const& outputLayersInfo, NvDsInferNetworkInfo const& networkInfo, + NvDsInferParseDetectionParams const& detectionParams, std::vector& objectList); + +extern "C" bool +NvDsInferParseYoloE_cuda(std::vector const& outputLayersInfo, NvDsInferNetworkInfo const& networkInfo, + NvDsInferParseDetectionParams const& detectionParams, std::vector& objectList); + +__global__ void decodeTensorYolo_cuda(NvDsInferParseObjectInfo *binfo, float* input, int outputSize, int netW, int netH, + float minPreclusterThreshold) +{ + int x_id = blockIdx.x * blockDim.x + threadIdx.x; + + if (x_id >= outputSize) + return; + + float maxProb = input[x_id * 6 + 4]; + int maxIndex = (int) input[x_id * 6 + 5]; + + if (maxProb < minPreclusterThreshold) { + binfo[x_id].detectionConfidence = 0.0; + return; + } + + float bxc = input[x_id * 6 + 0]; + float byc = input[x_id * 6 + 1]; + float bw = input[x_id * 6 + 2]; + float bh = input[x_id * 6 + 3]; + + float x0 = bxc - bw / 2; + float y0 = byc - bh / 2; + float x1 = x0 + bw; + float y1 = y0 + bh; + + x0 = fminf(float(netW), fmaxf(float(0.0), x0)); + y0 = fminf(float(netH), fmaxf(float(0.0), y0)); + x1 = fminf(float(netW), fmaxf(float(0.0), x1)); + y1 = fminf(float(netH), fmaxf(float(0.0), y1)); + + binfo[x_id].left = x0; + binfo[x_id].top = y0; + binfo[x_id].width = fminf(float(netW), fmaxf(float(0.0), x1 - x0)); + binfo[x_id].height = fminf(float(netH), fmaxf(float(0.0), y1 - y0)); + binfo[x_id].detectionConfidence = maxProb; + binfo[x_id].classId = maxIndex; +} + +__global__ void decodeTensorYoloE_cuda(NvDsInferParseObjectInfo *binfo, float* input, int outputSize, int netW, int netH, + float minPreclusterThreshold) +{ + int x_id = blockIdx.x * blockDim.x + threadIdx.x; + + if (x_id >= outputSize) + return; + + float maxProb = input[x_id * 6 + 4]; + int maxIndex = (int) input[x_id * 6 + 5]; + + if (maxProb < minPreclusterThreshold) { + binfo[x_id].detectionConfidence = 0.0; + return; + } + + float x0 = input[x_id * 6 + 0]; + float y0 = input[x_id * 6 + 1]; + float x1 = input[x_id * 6 + 2]; + float y1 = input[x_id * 6 + 3]; + + x0 = fminf(float(netW), fmaxf(float(0.0), x0)); + y0 = fminf(float(netH), fmaxf(float(0.0), y0)); + x1 = fminf(float(netW), fmaxf(float(0.0), x1)); + y1 = fminf(float(netH), fmaxf(float(0.0), y1)); + + binfo[x_id].left = x0; + binfo[x_id].top = y0; + binfo[x_id].width = fminf(float(netW), fmaxf(float(0.0), x1 - x0)); + binfo[x_id].height = fminf(float(netH), fmaxf(float(0.0), y1 - y0)); + binfo[x_id].detectionConfidence = maxProb; + binfo[x_id].classId = maxIndex; +} + +static bool NvDsInferParseCustomYolo_cuda(std::vector const& outputLayersInfo, + NvDsInferNetworkInfo const& networkInfo, NvDsInferParseDetectionParams const& detectionParams, + std::vector& objectList) +{ + if (outputLayersInfo.empty()) { + std::cerr << "ERROR: Could not find output layer in bbox parsing" << std::endl; + return false; + } + + const NvDsInferLayerInfo &layer = outputLayersInfo[0]; + + const int outputSize = layer.inferDims.d[0]; + + thrust::device_vector objects(outputSize); + + float minPreclusterThreshold = *(std::min_element(detectionParams.perClassPreclusterThreshold.begin(), + detectionParams.perClassPreclusterThreshold.end())); + + int threads_per_block = 1024; + int number_of_blocks = ((outputSize - 1) / threads_per_block) + 1; + + decodeTensorYolo_cuda<<>>( + thrust::raw_pointer_cast(objects.data()), (float*) layer.buffer, outputSize, networkInfo.width, networkInfo.height, + minPreclusterThreshold); + + objectList.resize(outputSize); + thrust::copy(objects.begin(), objects.end(), objectList.begin()); + + return true; +} + +static bool NvDsInferParseCustomYoloE_cuda(std::vector const& outputLayersInfo, + NvDsInferNetworkInfo const& networkInfo, NvDsInferParseDetectionParams const& detectionParams, + std::vector& objectList) +{ + if (outputLayersInfo.empty()) { + std::cerr << "ERROR: Could not find output layer in bbox parsing" << std::endl; + return false; + } + + const NvDsInferLayerInfo &layer = outputLayersInfo[0]; + + const int outputSize = layer.inferDims.d[0]; + + thrust::device_vector objects(outputSize); + + float minPreclusterThreshold = *(std::min_element(detectionParams.perClassPreclusterThreshold.begin(), + detectionParams.perClassPreclusterThreshold.end())); + + int threads_per_block = 1024; + int number_of_blocks = ((outputSize - 1) / threads_per_block) + 1; + + decodeTensorYoloE_cuda<<>>( + thrust::raw_pointer_cast(objects.data()), (float*) layer.buffer, outputSize, networkInfo.width, networkInfo.height, + minPreclusterThreshold); + + objectList.resize(outputSize); + thrust::copy(objects.begin(), objects.end(), objectList.begin()); + + return true; +} + +extern "C" bool +NvDsInferParseYolo_cuda(std::vector const& outputLayersInfo, NvDsInferNetworkInfo const& networkInfo, + NvDsInferParseDetectionParams const& detectionParams, std::vector& objectList) +{ + return NvDsInferParseCustomYolo_cuda(outputLayersInfo, networkInfo, detectionParams, objectList); +} + +extern "C" bool +NvDsInferParseYoloE_cuda(std::vector const& outputLayersInfo, NvDsInferNetworkInfo const& networkInfo, + NvDsInferParseDetectionParams const& detectionParams, std::vector& objectList) +{ + return NvDsInferParseCustomYoloE_cuda(outputLayersInfo, networkInfo, detectionParams, objectList); +} + +CHECK_CUSTOM_PARSE_FUNC_PROTOTYPE(NvDsInferParseYolo_cuda); +CHECK_CUSTOM_PARSE_FUNC_PROTOTYPE(NvDsInferParseYoloE_cuda);