#include "DetectionPostProcessPlugin.hpp" namespace MNN { template __global__ void decodeBoxes_kernel(const int count, const float* boxesPtr, const CenterSizeEncoding* anchorsPtr, BoxCornerEncoding* decodeBoxesPtr, const CenterSizeEncoding& scaleValues, int numBoxes, int boxCoordNum, int anchorsCoordNum, int numAnchors1) { CUDA_KERNEL_LOOP(idx, count) { const int boxIndex = idx * boxCoordNum; CenterSizeEncoding boxCenterSize = *reinterpret_cast(boxesPtr + boxIndex); CenterSizeEncoding anchor = anchorsPtr[idx]; float ycenter = boxCenterSize.y / scaleValues.y * anchor.h + anchor.y; float xcenter = boxCenterSize.x / scaleValues.x * anchor.w + anchor.x; float halfh = 0.5f * static_cast(exp(boxCenterSize.h / scaleValues.h)) * anchor.h; float halfw = 0.5f * static_cast(exp(boxCenterSize.w / scaleValues.w)) * anchor.w; auto& curBox = decodeBoxesPtr[idx]; curBox.ymin = ycenter - halfh; curBox.xmin = xcenter - halfw; curBox.ymax = ycenter + halfh; curBox.xmax = xcenter + halfw; } } template __global__ void maxScores_kernel(const int count, const float* scoresStartPtr, int numClassWithBackground, int labelOffset, int* sortedClassIndicesPtr, int numClasses, int numCategoriesPerAnchor, float* maxScores){ CUDA_KERNEL_LOOP(idx, count) { const auto boxScores = scoresStartPtr + idx * numClassWithBackground + labelOffset; int* classIndices = sortedClassIndicesPtr + idx * numClasses; // iota(classIndices, numClasses, 0); int data = 0; for(int i = 0; i < numClasses; i++){ classIndices[i] = data; data += 1; } // std::partial_sort(classIndices, classIndices + numCategoriesPerAnchor, classIndices + numClasses, // [&boxScores](const int i, const int j) { return boxScores[i] > boxScores[j]; }); int score = classIndices[0]; for(int i = 0; i < numClasses; i++){ score = max(classIndices[i], score); } maxScores[idx] = boxScores[score]; } } template __global__ void copy_candidate(const int count, Candidate* candidatePtr, const float* score){ CUDA_KERNEL_LOOP(idx, count) { int index = 0; for(int i = 0; i < count; i++){ if(score[idx] < score[i]){ index++; } } candidatePtr[idx].index = index; candidatePtr[idx].boxIndex = idx; candidatePtr[idx].score = score[idx]; } } __device__ __forceinline__ float IOU(const float* boxes, int i, int j) { const float yMinI = min(boxes[i * 4 + 0], boxes[i * 4 + 2]); const float xMinI = min(boxes[i * 4 + 1], boxes[i * 4 + 3]); const float yMaxI = max(boxes[i * 4 + 0], boxes[i * 4 + 2]); const float xMaxI = max(boxes[i * 4 + 1], boxes[i * 4 + 3]); const float yMinJ = min(boxes[j * 4 + 0], boxes[j * 4 + 2]); const float xMinJ = min(boxes[j * 4 + 1], boxes[j * 4 + 3]); const float yMaxJ = max(boxes[j * 4 + 0], boxes[j * 4 + 2]); const float xMaxJ = max(boxes[j * 4 + 1], boxes[j * 4 + 3]); const float areaI = (yMaxI - yMinI) * (xMaxI - xMinI); const float areaJ = (yMaxJ - yMinJ) * (xMaxJ - xMinJ); if (areaI <= 0 || areaJ <= 0) return 0.0; const float intersectionYMin = max(yMinI, yMinJ); const float intersectionXMin = max(xMinI, xMinJ); const float intersectionYMax = min(yMaxI, yMaxJ); const float intersectionXMax = min(xMaxI, xMaxJ); const float intersectionArea = max(intersectionYMax - intersectionYMin, 0.0) * max(intersectionXMax - intersectionXMin, 0.0); return intersectionArea / (areaI + areaJ - intersectionArea); } template __global__ void nms_kernel(const int count, int numBoxes, float scoreThreshold, float iouThreshold, Candidate* candidatePtr, int* selectedSize, float* decodedBoxesPtr, int* selectedPtr){ CUDA_KERNEL_LOOP(idx, count) { int boxIndex = 0; float originalScore = 0; for(int i = 0; i < numBoxes; i++){ if(candidatePtr[i].index == idx){ boxIndex = candidatePtr[i].boxIndex; originalScore = candidatePtr[i].score; } } if(originalScore <= scoreThreshold){ return; } bool shouldSelect = true; for (int j = (selectedSize[0] - 1); j >= 0; --j) { float iou = IOU(decodedBoxesPtr, boxIndex, selectedPtr[j]); if (iou == 0.0) { continue; } if (iou > iouThreshold) { shouldSelect = false; } } if (shouldSelect) { selectedPtr[selectedSize[0]] = boxIndex; atomicAdd(selectedSize, 1); } } } template __global__ void set_output(const int count, const BoxCornerEncoding* decodedBoxesPtr, BoxCornerEncoding* detectionBoxesPtr, float* detectionClassesPtr, float* detectionScoresPtr, float* numDetectionsPtr, const float* scoresStartPtr, int numClassWithBackground, int labelOffset, int* sortedClassIndicesPtr, int numClasses, int numCategoriesPerAnchor, int* selectedPtr){ CUDA_KERNEL_LOOP(index, count) { int selectedIndex = selectedPtr[index]; const float* boxScores = scoresStartPtr + selectedIndex * numClassWithBackground + labelOffset; const int* classIndices = sortedClassIndicesPtr + selectedIndex * numClasses; for (int col = 0; col < numCategoriesPerAnchor; ++col) { int boxOffset = numCategoriesPerAnchor * numDetectionsPtr[0] + col; detectionBoxesPtr[boxOffset] = decodedBoxesPtr[selectedIndex]; detectionClassesPtr[boxOffset] = classIndices[col]; detectionScoresPtr[boxOffset] = boxScores[classIndices[col]]; atomicAdd(numDetectionsPtr, 1); } } } void DetectionPostProcessPlugin::decodeBoxes(nvinfer1::DataType dataType, const int count, const void *const * inputs, const void *const * outputs, const void * scaleValues, void * decodeBoxes, int numBoxes, int boxCoordNum, int anchorsCoordNum, int numAnchors1) { auto boxesEncoding = inputs[0]; auto anchors = inputs[2]; if (dataType == nvinfer1::DataType::kFLOAT){ return decodeBoxes_kernel<<>>(count, (float*)boxesEncoding, reinterpret_cast(anchors), reinterpret_cast(decodeBoxes), *reinterpret_cast(scaleValues), numBoxes, boxCoordNum, anchorsCoordNum, numAnchors1); }else{ return decodeBoxes_kernel<__half><<>>(count, (float*)boxesEncoding, reinterpret_cast(anchors), reinterpret_cast(decodeBoxes), *reinterpret_cast(scaleValues), numBoxes, boxCoordNum, anchorsCoordNum, numAnchors1); } } void DetectionPostProcessPlugin::maxScores(nvinfer1::DataType dataType, const int count, const void *const * inputs, const void *const * outputs, int numClassWithBackground, int* sortedClassIndicesPtr, int numClasses, float* maxScores, int maxClassesPerAnchor) { auto classPredictions = inputs[1]; const int labelOffset = numClassWithBackground - numClasses; int numCategoriesPerAnchor = std::min(maxClassesPerAnchor, numClasses); maxScores_kernel<<>>(count, (const float*)classPredictions, numClassWithBackground, labelOffset, sortedClassIndicesPtr, numClasses, numCategoriesPerAnchor, maxScores); } void DetectionPostProcessPlugin::NMSSingleClasss(float* decodedBoxesPtr, const float* scoresPtr, int maxDetections, float iouThreshold, float scoreThreshold, int* selectedPtr, int* selectedSize, int numBoxes, int outputNum, Candidate* candidate, Candidate* mCandidatePriorityQueue){ copy_candidate<<>>(numBoxes, candidate, scoresPtr); nms_kernel<<>>(outputNum, numBoxes, scoreThreshold, iouThreshold, candidate, selectedSize, decodedBoxesPtr, selectedPtr); } void DetectionPostProcessPlugin::setOutput(const int selectSize, const BoxCornerEncoding* decodedBoxesPtr, BoxCornerEncoding* detectionBoxesPtr, float* detectionClassesPtr, float* detectionScoresPtr, float* numDetectionsPtr, const float* scoresStartPtr, int numClassWithBackground, int labelOffset, int* sortedClassIndicesPtr, int numClasses, int numCategoriesPerAnchor, int* selectedPtr){ set_output<<>>(selectSize, decodedBoxesPtr, detectionBoxesPtr, detectionClassesPtr, detectionScoresPtr, numDetectionsPtr, scoresStartPtr, numClassWithBackground, labelOffset, sortedClassIndicesPtr, numClasses, numCategoriesPerAnchor, selectedPtr); } }; // namespace MNN