用CUDA实现yolov5后处理核函数_decode + nms
·
目录
1 yolov5s_postprocess_simple.h
2 yolov5s_postprocess_simple.hip
1 yolov5s_postprocess_simple.h
// yolov5s_postprocess_simple.h
#pragma once
#include <hip/hip_runtime.h>
#include <vector>
#include <iostream>
#include <algorithm>
// 常量定义
constexpr int MAX_CLASSES = 80;
constexpr int OBJECT_TOPK = 1000; // 单张图最大输出框数量
constexpr int NUM_BOX_ELEMENT = 7; // left, top, right, bottom, confidence, class_id, keep_flag
constexpr int OUT_ELEM_SIZE = 6; // 输出最终保留的6个值 (left, top, right, bottom, confidence, class_id)
inline void checkHipError(hipError_t err) {
if (err != hipSuccess) {
std::cerr << "HIP Error: " << hipGetErrorString(err) << std::endl;
exit(-1);
}
}
// Bounding box结构体
struct Box {
float left; // 左上角 x
float top; // 左上角 y
float right; // 右下角 x
float bottom; // 右下角 y
float confidence; // 置信度
int class_id; // 类别
Box() : left(0), top(0), right(0), bottom(0), confidence(0), class_id(-1) {}
Box(float l, float t, float r, float b, float conf, int cls)
: left(l), top(t), right(r), bottom(b), confidence(conf), class_id(cls) {}
};
// 后处理类
class Yolov5sPostProcessSimple {
public:
Yolov5sPostProcessSimple();
~Yolov5sPostProcessSimple();
// 初始化：分配GPU/CPU内存和流
bool init(int max_batch = 1);
// 销毁资源
void destroy();
// // 执行GPU解码 + Fast NMS
// std::vector<Box> process(
// float* predict, // host pointer to [rows, cols] tensor
// int rows, // number of boxes (e.g., 25200)
// int cols, // 5 + num_classes
// float confidence_threshold = 0.25f,
// float nms_threshold = 0.45f
// );
// 执行GPU解码 + Fast NMS
std::vector<Box> process(
float* predict // host pointer to [rows, cols] tensor
);
// 是否已初始化
bool is_initialized() const { return initialized_; }
private:
// GPU 流
hipStream_t stream_;
// 显存指针
float* predict_device_; // 输入预测结果显存
float* output_device_; // 输出结果显存 [count + boxes]
float* output_host_; // 主机端映射内存用于同步
// 缓存参数
int max_objects_ = 1000; // 最大检测框数（默认1000）
size_t max_input_bytes_; // predict最大占用字节
bool initialized_;
int rows_ = 25200;
int cols_ = 85;
float confidence_threshold_ = 0.25;
float nms_threshold_ = 0.45;
};
2 yolov5s_postprocess_simple.hip
#include "hip/hip_runtime.h"
#include "yolov5s_postprocess_simple.h"
#include <hip/hip_runtime.h>
#include <iostream>
using namespace std;
static __device__ float box_iou(
float aleft, float atop, float aright, float abottom,
float bleft, float btop, float bright, float bbottom)
{
float cleft = max(aleft, bleft);
float ctop = max(atop, btop);
float cright = min(aright, bright);
float cbottom = min(abottom, bbottom);
float c_area = max(cright - cleft, 0.0f) * max(cbottom - ctop, 0.0f);
if (c_area == 0.0f)
return 0.0f;
float a_area = max(0.0f, aright - aleft) * max(0.0f, abottom - atop);
float b_area = max(0.0f, bright - bleft) * max(0.0f, bbottom - btop);
return c_area / (a_area + b_area - c_area);
}
static __global__ void __launch_bounds__(1024) fast_nms_kernel(float *bboxes, int max_objects, float threshold, int NUM_BOX_ELEMENT)
{
/*
参数解析:
bboxes：存储了所有待处理的检测框信息的一维数组；
max_objects：最大的输出检测框数量； 案例设置的是1000, 预计一张图不会超过1000个bounding box
threshold：用于判断两个检测框是否重叠的 IOU 阈值；
NUM_BOX_ELEMENT：每个检测框存储的元素个数
一般包含: left, top, right, bottom, confidence, class, keepflag
*/
// 计算position, 超过count不用进行下面计算了
int position = blockDim.x * blockIdx.x + threadIdx.x;
int count = min((int)*bboxes, max_objects);
if (position > count){
return;
}
/*
重叠度高, 并且类别相同，然后是condience小于另外一个, 就删掉他
极端情况下会有误删, 如果测试cpu map的时候, 只能采用cpu nms
日常推理的时候, 则可以使用这个NMS
left, top, right, bottom, confidence, class, keepflag
*/
// 这里计算出来当前的指针, 在bboxes上
float *pcurrent = bboxes + 1 + position * NUM_BOX_ELEMENT;
// 便利每一个bbox
for (int i = 0; i < count; ++i){
float *pitem = bboxes + 1 + i * NUM_BOX_ELEMENT;
// NMS计算需要保证类别必须相同
//if (i == position || pcurrent[5] ){//这一行错误，结果会有重叠框。
if (i == position || pcurrent[5] != pitem[5]){
continue;
}
// 判断置信度大小, 如果比pcurrent大，干掉pcurrent
if (pitem[4] > pcurrent[4]){
// 如果两个一样大，保留编号小的那个
if (pitem[4] == pcurrent[4] && i < position){
continue;
}
// 拿前面四个信息计算IOU
float iou = box_iou(
pcurrent[0], pcurrent[1], pcurrent[2], pcurrent[3],
pitem[0], pitem[1], pitem[2], pitem[3]);
if (iou > threshold){
pcurrent[6] = 0; // 这里pitem跟pcurrent重合度高而且达到阈值
return;
}
}
}
}
static __global__ void __launch_bounds__(1024) decode_kernel(
float *predict, int num_bboxes, int num_classes, float confidence_threshold,
float *invert_affine_matrix, float *parray, int max_objects, int NUM_BOX_ELEMENT)
{
// 确保有足够的thread, 每一个thread处理一个bounding box
// 如果threadId超过了bounding box的数量, 这样就不会进行后续处理, 每个预测框都敲好被处理了一次
int position = blockDim.x * blockIdx.x + threadIdx.x;
if (position >= num_bboxes){
return;
}
/*
predict是n x 85tensor输出的首地址
pitem 就是每行的指针, pitem[0] - pitem[3] 是位置信息, pitem[4]是objness
*/
float *pitem = predict + (num_classes + 5) * position;
float objectness = pitem[4];
if (objectness < confidence_threshold){
return;
}
// 从这个元素开始都是confidence
float *class_confidence = pitem + 5;
// 这里是第一个condience, 取到数值
float confidence = *class_confidence++;
// for循环判断是哪个类别
int label = 0;
for (int i = 1; i < num_classes; i++, ++class_confidence)
{
if (*class_confidence > confidence)
{
// 如果大了, 就更新class_confidence
confidence = *class_confidence;
label = i; // 取到label
}
}
/*
上面的最后算出来的condifence是class_confidence只是条件概率
当前bounding box的 confidence(置信度) = objectness(物体概率) x class_confidence(条件概率)
最后拿来计算置信度的confidence是最大的class_confidence
*/
confidence *= objectness;
if (confidence < confidence_threshold){
return;
}
/*
这里是恢复boudingbox的操作, 需要先取出来中心点(cx, cy), width, height
*/
float cx = *pitem++;
float cy = *pitem++;
float width = *pitem++;
float height = *pitem++;
float left = cx - width * 0.5f;
float top = cy - height * 0.5f;
float right = cx + width * 0.5f;
float bottom = cy + height * 0.5f;
// affine_project(invert_affine_matrix, left, top, &left, &top);
// affine_project(invert_affine_matrix, right, bottom, &right, &bottom);
/*
atomicAdd()简介:
int atomicAdd(int* address, int val);
这个函数执行的操作是将指定地址 address 处的值与 val 相加，并将结果写回 address 处。这个操作是原子性的，即不会受到并发写入的干扰，保证了数据的正确性。
使用 atomicAdd 函数可以保证多个线程在对同一个内存地址进行写操作时，不会发生数据覆盖的问题。
由于每个线程都会在输出中写入一个bounding box，因此需要使用原子操作确保每个线程写入的位置唯一
*/
/*
[count, box1, box2, box3]
因为GPU解码是多线程的, 所以需要用count记录已经处理了多少个bounding box。
CPU单线程不需要, GPU需要确保不会将一个检测框重复输出或者漏掉。
atomicAdd -> count +=1 返回 old_count
这里是对parray(output_device第一个值+1)
*/
int index = atomicAdd(parray, 1);
// 如果超过了1000, 这个线程就没必要处理后面的boxes
if (index >= max_objects)
return;
// left, top, right, bottom, confidence, class, keepflag
float *pout_item = parray + 1 + index * NUM_BOX_ELEMENT;
*pout_item++ = left;
*pout_item++ = top;
*pout_item++ = right;
*pout_item++ = bottom;
*pout_item++ = confidence;
*pout_item++ = label;
*pout_item++ = 1; // 1 = keep, 0 = ignore
}
/*
decode_kernel_invoker(
predict_device, rows, cols - 5, confidence_threshold,
nms_threshold, nullptr, output_device, max_objects, NUM_BOX_ELEMENT, stream);
*/
void decode_kernel_invoker(
float* predict, int num_bboxes, int num_classes, float confidence_threshold,
float nms_threshold, float* invert_affine_matrix, float* parray, int max_objects, int NUM_BOX_ELEMENT, hipStream_t stream){
/*
参数解析:
predict: 预测结果, 这个就是data, 未处理未过滤的predict
num_bboxes: 在预测结果的（n x num_classes+ 5） tensor中, 多少行就是多少个box
num_classes: 类别数量
confidence_threshold: 置信度阈值
nms_threshold: nms阈值
invert_affine_matrix: 逆矩阵的指针
parray: 输出结果数组
max_objects: 最大数量框, 这边设置的是1000, 只是拿来确保有足够的内存
NUM_BOX_ELEMENT: Box的element, left, top, right, bottom, confidence, class, keepflag 一共7个
stream： 流
*/
// 这里是确保有足够的线程去处理每一个box, 也就是每一个预测结果，所以用num_boxxes
// 确保每个block的线程不超过512
auto block = num_bboxes > 512 ? 512 : num_bboxes;
auto grid = (num_bboxes + block - 1) / block;
/* 如果核函数有波浪线，没关系，他是正常的，你只是看不顺眼罢了 */
decode_kernel<<<grid, block, 0, stream>>>(
predict, num_bboxes, num_classes, confidence_threshold,
invert_affine_matrix, parray, max_objects, NUM_BOX_ELEMENT
);
// 这里是针对每张图的框，确保每个狂都能被线程处理
// 同样确保每个block的线程不超过512
block = max_objects > 512 ? 512 : max_objects;
grid = (max_objects + block - 1) / block;
fast_nms_kernel<<<grid, block, 0, stream>>>(parray, max_objects, nms_threshold, NUM_BOX_ELEMENT);
}
#if 0
vector<Box> gpu_decode(float *predict, int rows, int cols, float confidence_threshold, float nms_threshold)
{
/*
参数解析:
predict 是指向data的指针
rows 是输出data的行数
cols是是5(位置信息) + 类别
confidence_threshold 置信度阈值
nms_threshold nms阈值
*/
vector<Box> box_result; // 用来保存解码后的边界框信息
// 创建一个hipStream_t类型的流, stream
hipStream_t stream = nullptr;
checkHipError(hipStreamCreate(&stream));
// 定义三个float指针变量
float *predict_device = nullptr;
float *output_device = nullptr;
float *output_host = nullptr;
/*表示输出数量不确定的数组，用[count, box1, box2, box3]的方式，此时需要有最大数量限制
分配一块足够大的内存, 第一个元素是count, 这里我们设置每个图像不超过1000个bounding box
*/
int max_objects = 1000;
// left, top, right, bottom, confidence, class, keepflag 一共7个
int NUM_BOX_ELEMENT = 7;
// 分配内存到GPU上
checkHipError(hipMalloc(&predict_device, rows * cols * sizeof(float)));
// 分配一块足够大的内存, 第一个元素是count
checkHipError(hipMalloc(&output_device, sizeof(float) + max_objects * NUM_BOX_ELEMENT * sizeof(float)));
// 分配CPU内存
checkHipError(hipHostMalloc(&output_host, sizeof(float) + max_objects * NUM_BOX_ELEMENT * sizeof(float)));
// 把predict的内存复制到predict_device
checkHipError(hipMemcpyAsync(predict_device, predict, rows * cols * sizeof(float), hipMemcpyHostToDevice, stream));
// 调用kernel函数, 在GPU进行解码
decode_kernel_invoker(
predict_device, rows, cols - 5, confidence_threshold,
nms_threshold, nullptr, output_device, max_objects, NUM_BOX_ELEMENT, stream);
// 使用hipMemcpyAsync将解码好的output_device拷贝到CPU上面去
checkHipError(hipMemcpyAsync(output_host, output_device,
sizeof(float) + max_objects * NUM_BOX_ELEMENT * sizeof(float),
hipMemcpyDeviceToHost, stream));
printf("===============sizeof(int) is %d, sizeof(float) is %d,\n", sizeof(int), sizeof(float));
// 等待异步拷贝完成
checkHipError(hipStreamSynchronize(stream));
// 这里解码完了之后, 这个是框框的个数
int num_boxes = min((int)output_host[0], max_objects);
// 遍历每一个框框
for (int i = 0; i < num_boxes; i++)
{
// 找到第一个地址
float *ptr = output_host + 1 + NUM_BOX_ELEMENT * i;
int keep_flag = ptr[6]; // 最后一个位置就是flag的值
if (keep_flag) // True
{
box_result.emplace_back(
ptr[0], ptr[1], ptr[2], ptr[3], ptr[4], (int)ptr[5]);
}
}
// 销毁先前创建的CUDA流对象, 释放流对象占用的内存空间
checkHipError(hipStreamDestroy(stream));
checkHipError(hipFree(predict_device)); // 释放GPU上分配预测结果的内存
checkHipError(hipFree(output_device)); // 释放GPU上分配解码输出的内存
checkHipError(hipHostFree(output_host)); // 释放在主机上分配的输出结果缓冲区的内存
return box_result;
}
#endif
// =========================================
// 类方法实现
// =========================================
Yolov5sPostProcessSimple::Yolov5sPostProcessSimple()
: stream_(nullptr)
, predict_device_(nullptr)
, output_device_(nullptr)
, output_host_(nullptr)
, max_objects_(OBJECT_TOPK)
, max_input_bytes_(0)
, initialized_(false)
{}
Yolov5sPostProcessSimple::~Yolov5sPostProcessSimple() {
destroy();
}
bool Yolov5sPostProcessSimple::init(int max_batch) {
if (initialized_) return true;
checkHipError(hipStreamCreate(&stream_));
// 假设最大输入为 25200 x 85 (YOLOv5s)
// 可根据实际模型调整
size_t max_predict_size = rows_ * cols_ * sizeof(float);
checkHipError(hipMalloc(&predict_device_, max_predict_size));
max_input_bytes_ = max_predict_size;
size_t output_size = sizeof(float) + max_objects_ * NUM_BOX_ELEMENT * sizeof(float);
checkHipError(hipMalloc(&output_device_, output_size));
checkHipError(hipHostMalloc(&output_host_, output_size));
initialized_ = true;
return true;
}
void Yolov5sPostProcessSimple::destroy() {
if (stream_) {
checkHipError(hipStreamDestroy(stream_));
stream_ = nullptr;
}
if (predict_device_) {
checkHipError(hipFree(predict_device_));
predict_device_ = nullptr;
}
if (output_device_) {
checkHipError(hipFree(output_device_));
output_device_ = nullptr;
}
if (output_host_) {
checkHipError(hipHostFree(output_host_));
output_host_ = nullptr;
}
max_input_bytes_ = 0;
initialized_ = false;
}
// std::vector<Box> Yolov5sPostProcessSimple::process(
// float* predict,
// int rows,
// int cols,
// float confidence_threshold,
// float nms_threshold)
std::vector<Box> Yolov5sPostProcessSimple::process(float* predict)
{
std::vector<Box> box_result;
if (!initialized_) {
fprintf(stderr, "Error: Yolov5sPostProcessSimple not initialized!\n");
return box_result;
}
int num_classes = cols_ - 5;
size_t input_bytes = rows_ * cols_ * sizeof(float);
// 检查输入大小是否超出预分配
if (input_bytes > max_input_bytes_) {
fprintf(stderr, "Input size too large: %zu > %zu\n", input_bytes, max_input_bytes_);
return box_result;
}
// 复制输入到GPU
checkHipError(hipMemcpyAsync(predict_device_, predict, input_bytes,
hipMemcpyHostToDevice, stream_));
// 重置输出计数器
float init_count = 0.0f;
checkHipError(hipMemcpyAsync(output_device_, &init_count, sizeof(float),
hipMemcpyHostToDevice, stream_));
// 调用核函数
decode_kernel_invoker(
predict_device_, rows_, num_classes, confidence_threshold_,
nms_threshold_, nullptr, output_device_, max_objects_, NUM_BOX_ELEMENT, stream_
);
// 拷贝结果回主机
size_t output_size = sizeof(float) + max_objects_ * NUM_BOX_ELEMENT * sizeof(float);
checkHipError(hipMemcpyAsync(output_host_, output_device_, output_size,
hipMemcpyDeviceToHost, stream_));
// 同步流
checkHipError(hipStreamSynchronize(stream_));
// 解析结果
int num_boxes = (int)output_host_[0];
num_boxes = min(num_boxes, max_objects_);
for (int i = 0; i < num_boxes; ++i) {
float* ptr = output_host_ + 1 + i * NUM_BOX_ELEMENT;
if ((int)ptr[6] == 1) { // keep flag
box_result.emplace_back(ptr[0], ptr[1], ptr[2], ptr[3], ptr[4], (int)ptr[5]);
}
}
return box_result;
}3 demo
void Postprocess(
const std::unique_ptr<tc::InferResult>& result,
const std::vector<std::string>& filenames,
const size_t batch_size,
const std::string& output_name,
const int maxBatchSize)
{
if (!result->RequestStatus().IsOk()) {
std::cerr << "inference failed with error: " << result->RequestStatus() << std::endl;
exit(1);
}
if (filenames.size() != batch_size) {
std::cerr << "expected " << batch_size << " filenames, got " << filenames.size() << std::endl;
exit(1);
}
// 获取输出 shape
std::vector<int64_t> shape;
tc::Error err = result->Shape(output_name, &shape);
if (!err.IsOk() || shape.size() != 3) {
std::cerr << "unable to get proper shape for " << output_name << std::endl;
exit(1);
}
int batch = static_cast<int>(shape[0]);
int numProposal = static_cast<int>(shape[1]);
int numOut = static_cast<int>(shape[2]);
// 获取输出数据
const uint8_t* result_data;
size_t output_count;
err = result->RawData(output_name, &result_data, &output_count);
if (!err.IsOk()) {
std::cerr << "unable to get data for " << output_name << std::endl;
exit(1);
}
const float* pdata = reinterpret_cast<const float*>(result_data);
if (batch * numProposal * numOut != output_count / sizeof(float)) {
std::cerr << "data size mismatch" << std::endl;
exit(1);
}
// 申请 pinned host buffer 并拷贝
void* pinned_input = nullptr;
hipHostAlloc(&pinned_input, output_count, hipHostMallocDefault);
memcpy(pinned_input, pdata, output_count);
const void* input_ptrs[] = {pinned_input};
// // 创建后处理类
// auto yolov5sPostProcess = std::make_shared<Yolov5sPostProcess>();
// // 初始化
// auto t1 = std::chrono::high_resolution_clock::now();
// yolov5sPostProcess->initialize(batch, numProposal, numOut);
// auto t2 = std::chrono::high_resolution_clock::now();
// std::cout << "<<<<<< initialize time = "
// << std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count() * 1000
// << " ms" << std::endl;
// // 申请输出 buffer
// void* outputs = nullptr;
// hipHostAlloc(&outputs, batch * OBJECT_TOPK * OUT_ELEM_SIZE * sizeof(float), hipHostMallocDefault);
// auto t3 = std::chrono::high_resolution_clock::now();
// std::cout << "<<<<<< hipHostAlloc time = "
// << std::chrono::duration_cast<std::chrono::duration<double>>(t3 - t2).count() * 1000
// << " ms" << std::endl;
// // 调用 GPU 核函数进行 decode + NMS
// printf("numProposal is %d,numOut is %d\n", numProposal, numOut);
// auto box_result = yolov5sPostProcess->gpu_decode(
// reinterpret_cast<float*>(pinned_input),
// numProposal,
// numOut - 5 // num_classes
// );
// auto box_result = gpu_decode(
// reinterpret_cast<float*>(pinned_input),
// numProposal,
// numOut); // num_classes
;
auto yolov5sPostProcessSimple_p = std::make_shared<Yolov5sPostProcessSimple>();
yolov5sPostProcessSimple_p->init();
auto box_result = yolov5sPostProcessSimple_p->process(
reinterpret_cast<float*>(pinned_input)
);
// auto t4 = std::chrono::high_resolution_clock::now();
// std::cout << "<<<<<< decode + NMS time = "
// << std::chrono::duration_cast<std::chrono::duration<double>>(t4 - t3).count() * 1000
// << " ms" << std::endl;
std::cout << "===== Decoded Boxes =====" << std::endl;
for (size_t i = 0; i < 50; ++i) {
const Box& box = box_result[i];
std::cout << "Box " << i
<< ": [" << box.left << ", " << box.top
<< ", " << box.right << ", " << box.bottom << "]"
<< ", confidence = " << box.confidence
<< ", class_id = " << box.class_id
<< std::endl;
}
std::cout << "Total boxes: " << box_result.size() << std::endl;
// // 绘制结果并保存
// 绘制结果并保存
for (int b = 0; b < batch; ++b) {
const std::string& filename = filenames[b];
cv::Mat srcImage = cv::imread(filename, 1);
if (srcImage.empty()) {
std::cerr << "Failed to load image: " << filename << std::endl;
continue;
}
// 遍历检测框
for (size_t i = 0; i < box_result.size(); i++) {
const Box& box = box_result[i];
//if (box.confidence < 0.5f) continue; // 过滤低置信度
cv::Point pt1((int)box.left, (int)box.top);
cv::Point pt2((int)box.right, (int)box.bottom);
// 画矩形框
cv::rectangle(srcImage, pt1, pt2, cv::Scalar(0, 255, 255), 2);
// 显示类别与置信度
std::string label = "ID:" + std::to_string(box.class_id) +
" conf:" + cv::format("%.2f", box.confidence);
int baseline = 0;
cv::Size label_size = cv::getTextSize(label, cv::FONT_HERSHEY_SIMPLEX, 0.5, 1, &baseline);
cv::Point label_origin(pt1.x, pt1.y - 5);
// 背景方块
// cv::rectangle(srcImage,
// cv::Rect(label_origin.x,
// label_origin.y - label_size.height,
// label_size.width,
// label_size.height + baseline),
// cv::Scalar(0, 255, 255));
// 文本
cv::putText(srcImage, label, label_origin,
cv::FONT_HERSHEY_SIMPLEX, 0.5, cv::Scalar(0, 0, 0), 1);
}
std::string out_filename = "result_" + std::to_string(b) + ".jpg";
cv::imwrite(out_filename, srcImage);
std::cout << "Saved result image: " << out_filename << std::endl;
}
// 释放资源
//yolov5sPostProcess->terminate();
//hipHostFree(outputs);
hipHostFree(pinned_input);
// auto t5 = std::chrono::high_resolution_clock::now();
// std::cout << "<<<<<< total terminate/free time = "
// << std::chrono::duration_cast<std::chrono::duration<double>>(t5 - t4).count() * 1000
// << " ms" << std::endl;
}
}4 说明
这其实是参考https://blog.csdn.net/bobchen1017/article/details/129664903 进行修改的，做了如下修改。
4.1 封装成类
4.2 修改nms错误
if (i == position || pcurrent[5] )这一行是错误的，要改成if (i == position || pcurrent[5] != pitem[5])，也就是类别不同的时候不做nms，只有同类别的才做nms，这一行如果不修改，那么最终结果有很多重复的框。
static __global__ void fast_nms_kernel(float *bboxes, int max_objects, float threshold, int NUM_BOX_ELEMENT)
{
/*
参数解析:
bboxes：存储了所有待处理的检测框信息的一维数组；
max_objects：最大的输出检测框数量； 案例设置的是1000, 预计一张图不会超过1000个bounding box
threshold：用于判断两个检测框是否重叠的 IOU 阈值；
NUM_BOX_ELEMENT：每个检测框存储的元素个数
一般包含: left, top, right, bottom, confidence, class, keepflag
*/
// 计算position, 超过count不用进行下面计算了
int position = blockDim.x * blockIdx.x + threadIdx.x;
int count = min((int)*bboxes, max_objects);
if (position > count){
return;
}
/*
重叠度高, 并且类别相同，然后是condience小于另外一个, 就删掉他
极端情况下会有误删, 如果测试cpu map的时候, 只能采用cpu nms
日常推理的时候, 则可以使用这个NMS
left, top, right, bottom, confidence, class, keepflag
*/
// 这里计算出来当前的指针, 在bboxes上
float *pcurrent = bboxes + 1 + position * NUM_BOX_ELEMENT;
// 便利每一个bbox
for (int i = 0; i < count; ++i){
float *pitem = bboxes + 1 + i * NUM_BOX_ELEMENT;
// NMS计算需要保证类别必须相同
//if (i == position || pcurrent[5] ){
if (i == position || pcurrent[5] != pitem[5]){
continue;
}
// 判断置信度大小, 如果比pcurrent大，干掉pcurrent
if (pitem[4] > pcurrent[4]){
// 如果两个一样大，保留编号小的那个
if (pitem[4] == pcurrent[4] && i < position){
continue;
}
// 拿前面四个信息计算IOU
float iou = box_iou(
pcurrent[0], pcurrent[1], pcurrent[2], pcurrent[3],
pitem[0], pitem[1], pitem[2], pitem[3]);
if (iou > threshold){
pcurrent[6] = 0; // 这里pitem跟pcurrent重合度高而且达到阈值
return;
}
}
}
}4.3 修改sizeof错误
// 使用hipMemcpyAsync将解码好的output_device拷贝到CPU上面去
checkHipError(hipMemcpyAsync(output_host, output_device,
sizeof(float) + max_objects * NUM_BOX_ELEMENT * sizeof(float),
hipMemcpyDeviceToHost, stream));这个地方最开始是sizeof(int)， 但其实这里应该改成sizeof(float)，只不过这里用sizeof(int)最终结果也是对的，因为他们两个是一样的，都是4.
5 结果图
更多推荐
5
0- 0
已为社区贡献12条内容
所有评论(0)