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

Move GPU post-processing separately to cuda_postprocess.cu (#1331)

Browse Source 
* Move GPU post-processing separately to cuda_postprocess.cu

* Update postprocess.cpp

* Update postprocess.cu

---------

Co-authored-by: wang-xinyu <wangxinyu_es@163.com>
Co-authored-by: Wang Xinyu <shaywxy@gmail.com>
This commit is contained in:
lindsayshuo 2023-07-18 13:23:56 +08:00 committed by GitHub
parent e9a972e332
commit ae3bd5e6b2
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
9 changed files with 596 additions and 610 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
yolov8/include/model.h
View File

@ -3,17 +3,17 @@
#include <string>
#include <assert.h>
nvinfer1::IHostMemory* buildEngineYolov8n(const int& batchsize, nvinfer1::IBuilder* builder,
nvinfer1::IHostMemory* buildEngineYolov8n(nvinfer1::IBuilder* builder,
nvinfer1::IBuilderConfig* config, nvinfer1::DataType dt, const std::string& wts_path);
nvinfer1::IHostMemory* buildEngineYolov8s(const int& batchsize, nvinfer1::IBuilder* builder,
nvinfer1::IHostMemory* buildEngineYolov8s(nvinfer1::IBuilder* builder,
nvinfer1::IBuilderConfig* config, nvinfer1::DataType dt, const std::string& wts_path);
nvinfer1::IHostMemory* buildEngineYolov8m(const int& batchsize, nvinfer1::IBuilder* builder,
nvinfer1::IHostMemory* buildEngineYolov8m(nvinfer1::IBuilder* builder,
nvinfer1::IBuilderConfig* config, nvinfer1::DataType dt, const std::string& wts_path);
nvinfer1::IHostMemory* buildEngineYolov8l(const int& batchsize, nvinfer1::IBuilder* builder,
nvinfer1::IHostMemory* buildEngineYolov8l(nvinfer1::IBuilder* builder,
nvinfer1::IBuilderConfig* config, nvinfer1::DataType dt, const std::string& wts_path);
nvinfer1::IHostMemory* buildEngineYolov8x(const int& batchsize, nvinfer1::IBuilder* builder,
nvinfer1::IHostMemory* buildEngineYolov8x(nvinfer1::IBuilder* builder,
nvinfer1::IBuilderConfig* config, nvinfer1::DataType dt, const std::string& wts_path);

12
yolov8/include/postprocess.h
View File

@ -1,6 +1,7 @@
#pragma once
#include "types.h"
#include "NvInfer.h"
#include <opencv2/opencv.hpp>
cv::Rect get_rect(cv::Mat& img, float bbox[4]);
@ -9,6 +10,13 @@ void nms(std::vector<Detection>& res, float *output, float conf_thresh, float nm
void batch_nms(std::vector<std::vector<Detection>>& batch_res, float *output, int batch_size, int output_size, float conf_thresh, float nms_thresh = 0.5);
void draw_bbox(std::vector<cv::Mat>& img_batch, std::vector<std::vector<Detection>>& res_batch);
void draw_bbox(std::vector<cv::Mat> &img_batch, std::vector<std::vector<Detection>> &res_batch);
void batch_process(std::vector<std::vector<Detection>> &res_batch, const float* decode_ptr_host, int batch_size, int bbox_element, const std::vector<cv::Mat>& img_batch);
void process_decode_ptr_host(std::vector<Detection> &res, const float* decode_ptr_host, int bbox_element, cv::Mat& img, int count);
void cuda_decode(float* predict, int num_bboxes, float confidence_threshold,float* parray,int max_objects, cudaStream_t stream);
void cuda_nms(float* parray, float nms_threshold, int max_objects, cudaStream_t stream);
void batch_process(std::vector<std::vector<Detection>>& res_batch, const float* decode_ptr_host, int batch_size, int bbox_element, const std::vector<cv::Mat>& img_batch);

9
yolov8/include/preprocess.h
View File

@ -6,12 +6,6 @@
#include <map>
struct AffineMatrix {
float value[6];
};
const int bbox_element = sizeof(AffineMatrix) / sizeof(float)+1; // left, top, right, bottom, confidence, class, keepflag
void cuda_preprocess_init(int max_image_size);
void cuda_preprocess_destroy();
@ -20,6 +14,3 @@ void cuda_preprocess(uint8_t *src, int src_width, int src_height, float *dst, in
void cuda_batch_preprocess(std::vector<cv::Mat> &img_batch, float *dst, int dst_width, int dst_height, cudaStream_t stream);
void cuda_decode(float* predict, int num_bboxes, float confidence_threshold,float* parray,int max_objects, cudaStream_t stream);
void cuda_nms(float* parray, float nms_threshold, int max_objects, cudaStream_t stream);

5
yolov8/include/types.h
View File

@ -8,3 +8,8 @@ struct alignas(float) Detection {
float class_id;
};
struct AffineMatrix {
float value[6];
};
const int bbox_element = sizeof(AffineMatrix) / sizeof(float)+1; // left, top, right, bottom, confidence, class, keepflag

29
yolov8/main.cpp
View File

@ -1,32 +1,33 @@
#include <iostream>
#include <fstream>
#include <opencv2/opencv.hpp>
#include "model.h"
#include "utils.h"
#include "preprocess.h"
#include "postprocess.h"
#include <iostream>
#include <opencv2/opencv.hpp>
#include "cuda_utils.h"
#include <fstream>
#include "logging.h"
Logger gLogger;
using namespace nvinfer1;
const int kOutputSize = kMaxNumOutputBbox * sizeof(Detection) / sizeof(float) + 1;
void serialize_engine(const int &kBatchSize, std::string &wts_name, std::string &engine_name, std::string &sub_type) {
void serialize_engine(std::string &wts_name, std::string &engine_name, std::string &sub_type) {
IBuilder *builder = createInferBuilder(gLogger);
IBuilderConfig *config = builder->createBuilderConfig();
IHostMemory *serialized_engine = nullptr;
if (sub_type == "n") {
serialized_engine = buildEngineYolov8n(kBatchSize, builder, config, DataType::kFLOAT, wts_name);
serialized_engine = buildEngineYolov8n(builder, config, DataType::kFLOAT, wts_name);
} else if (sub_type == "s") {
serialized_engine = buildEngineYolov8s(kBatchSize, builder, config, DataType::kFLOAT, wts_name);
serialized_engine = buildEngineYolov8s(builder, config, DataType::kFLOAT, wts_name);
} else if (sub_type == "m") {
serialized_engine = buildEngineYolov8m(kBatchSize, builder, config, DataType::kFLOAT, wts_name);
serialized_engine = buildEngineYolov8m(builder, config, DataType::kFLOAT, wts_name);
} else if (sub_type == "l") {
serialized_engine = buildEngineYolov8l(kBatchSize, builder, config, DataType::kFLOAT, wts_name);
serialized_engine = buildEngineYolov8l(builder, config, DataType::kFLOAT, wts_name);
} else if (sub_type == "x") {
serialized_engine = buildEngineYolov8x(kBatchSize, builder, config, DataType::kFLOAT, wts_name);
serialized_engine = buildEngineYolov8x(builder, config, DataType::kFLOAT, wts_name);
}
assert(serialized_engine);
@ -88,12 +89,12 @@ void prepare_buffer(ICudaEngine *engine, float **input_buffer_device, float **ou
}
}
void infer(IExecutionContext &context, cudaStream_t &stream, void **buffers, float *output, int batchSize, float* decode_ptr_host, float* decode_ptr_device, int batchSize_in, int model_bboxes, std::string cuda_post_process) {
void infer(IExecutionContext &context, cudaStream_t &stream, void **buffers, float *output, int batchsize, float* decode_ptr_host, float* decode_ptr_device, int model_bboxes, std::string cuda_post_process) {
// infer on the batch asynchronously, and DMA output back to host
auto start = std::chrono::system_clock::now();
context.enqueue(batchSize, buffers, stream, nullptr);
context.enqueue(batchsize, buffers, stream, nullptr);
if (cuda_post_process == "c") {
CUDA_CHECK(cudaMemcpyAsync(output, buffers[1], batchSize * kOutputSize * sizeof(float), cudaMemcpyDeviceToHost,stream));
CUDA_CHECK(cudaMemcpyAsync(output, buffers[1], batchsize * kOutputSize * sizeof(float), cudaMemcpyDeviceToHost,stream));
auto end = std::chrono::system_clock::now();
std::cout << "inference time: " << std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count() << "ms" << std::endl;
} else if (cuda_post_process == "g") {
@ -143,7 +144,7 @@ int main(int argc, char **argv) {
// Create a model using the API directly and serialize it to a file
if (!wts_name.empty()) {
serialize_engine(kBatchSize, wts_name, engine_name, sub_type);
serialize_engine(wts_name, engine_name, sub_type);
return 0;
}
@ -185,7 +186,7 @@ int main(int argc, char **argv) {
// Preprocess
cuda_batch_preprocess(img_batch, device_buffers[0], kInputW, kInputH, stream);
// Run inference
infer(*context, stream, (void **)device_buffers, output_buffer_host, kBatchSize, decode_ptr_host, decode_ptr_device, img_batch.size(), model_bboxes, cuda_post_process);
infer(*context, stream, (void **)device_buffers, output_buffer_host, kBatchSize, decode_ptr_host, decode_ptr_device, model_bboxes, cuda_post_process);
std::vector<std::vector<Detection>> res_batch;
if (cuda_post_process == "c") {
// NMS

952
yolov8/src/model.cpp
View File

File diff suppressed because it is too large Load Diff

2
yolov8/src/postprocess.cpp
View File

@ -82,7 +82,6 @@ void batch_nms(std::vector<std::vector<Detection>> &res_batch, float *output, in
}
}
void process_decode_ptr_host(std::vector<Detection> &res, const float* decode_ptr_host, int bbox_element, cv::Mat& img, int count) {
Detection det;
for (int i = 0; i < count; i++) {
@ -110,7 +109,6 @@ void batch_process(std::vector<std::vector<Detection>> &res_batch, const float*
}
}
void draw_bbox(std::vector<cv::Mat> &img_batch, std::vector<std::vector<Detection>> &res_batch) {
for (size_t i = 0; i < img_batch.size(); i++) {
auto &res = res_batch[i];

98
yolov8/src/postprocess.cu Normal file
View File

@ -0,0 +1,98 @@
//
// Created by lindsay on 23-7-17.
//
#include "types.h"
#include "postprocess.h"
static __global__ void
decode_kernel(float *predict, int num_bboxes, float confidence_threshold, float *parray, int max_objects) {
float count = predict[0];
int position = (blockDim.x * blockIdx.x + threadIdx.x);
if (position >= count)
return;
float *pitem = predict + 1 + position * 6;
int index = atomicAdd(parray, 1);
if (index >= max_objects)
return;
float confidence = pitem[4];
if (confidence < confidence_threshold)
return;
float left = pitem[0];
float top = pitem[1];
float right = pitem[2];
float bottom = pitem[3];
float label = pitem[5];
float *pout_item = parray + 1 + index * bbox_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
}
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 nms_kernel(float *bboxes, int max_objects, float threshold) {
int position = (blockDim.x * blockIdx.x + threadIdx.x);
int count = bboxes[0];
// float count = 0.0f;
if (position >= count)
return;
float *pcurrent = bboxes + 1 + position * bbox_element;
for (int i = 1; i < count; ++i) {
float *pitem = bboxes + 1 + i * bbox_element;
if (i == position || pcurrent[5] != pitem[5]) continue;
if (pitem[4] >= pcurrent[4]) {
if (pitem[4] == pcurrent[4] && i < position)
continue;
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;
return;
}
}
}
}
void cuda_decode(float *predict, int num_bboxes, float confidence_threshold, float *parray, int max_objects,
cudaStream_t stream) {
int block = 256;
int grid = ceil(num_bboxes / (float) block);
decode_kernel << <
grid, block, 0, stream >> > ((float *) predict, num_bboxes, confidence_threshold, parray, max_objects);
}
void cuda_nms(float *parray, float nms_threshold, int max_objects, cudaStream_t stream) {
int block = max_objects < 256 ? max_objects : 256;
int grid = ceil(max_objects / (float) block);
nms_kernel << < grid, block, 0, stream >> > (parray, max_objects, nms_threshold);
}

89
yolov8/src/preprocess.cu
View File

@ -87,82 +87,7 @@ warpaffine_kernel(uint8_t *src, int src_line_size, int src_width, int src_height
*pdst_c2 = c2;
}
static __global__ void
decode_kernel(float *predict, int num_bboxes, float confidence_threshold, float *parray, int max_objects) {
float count = predict[0];
int position = (blockDim.x * blockIdx.x + threadIdx.x);
if (position >= count)
return;
float *pitem = predict + 1 + position * 6;
int index = atomicAdd(parray, 1);
if (index >= max_objects)
return;
float confidence = pitem[4];
if (confidence < confidence_threshold)
return;
float left = pitem[0];
float top = pitem[1];
float right = pitem[2];
float bottom = pitem[3];
float label = pitem[5];
float *pout_item = parray + 1 + index * bbox_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
}
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 nms_kernel(float *bboxes, int max_objects, float threshold) {
int position = (blockDim.x * blockIdx.x + threadIdx.x);
int count = bboxes[0];
// float count = 0.0f;
if (position >= count)
return;
float *pcurrent = bboxes + 1 + position * bbox_element;
for (int i = 1; i < count; ++i) {
float *pitem = bboxes + 1 + i * bbox_element;
if (i == position || pcurrent[5] != pitem[5]) continue;
if (pitem[4] >= pcurrent[4]) {
if (pitem[4] == pcurrent[4] && i < position)
continue;
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;
return;
}
}
}
}
void cuda_preprocess(uint8_t *src, int src_width, int src_height, float *dst, int dst_width, int dst_height,
@ -210,21 +135,7 @@ void cuda_batch_preprocess(std::vector<cv::Mat> &img_batch,
}
void cuda_decode(float *predict, int num_bboxes, float confidence_threshold, float *parray, int max_objects,
cudaStream_t stream) {
int block = 256;
int grid = ceil(num_bboxes / (float) block);
decode_kernel << <
grid, block, 0, stream >> > ((float *) predict, num_bboxes, confidence_threshold, parray, max_objects);
}
void cuda_nms(float *parray, float nms_threshold, int max_objects, cudaStream_t stream) {
int block = max_objects < 256 ? max_objects : 256;
int grid = ceil(max_objects / (float) block);
nms_kernel << < grid, block, 0, stream >> > (parray, max_objects, nms_threshold);
}
void cuda_preprocess_init(int max_image_size) {