YOLO v5 前后处理GPU加速部署_yolov5 gpu 后处理 batch

        深度学习模型部署中复杂的前后处理往往部署cpu中，这一方面极大地增加模型处理耗时，另一方面挤兑了cpu计算资源，导致整个软件平台卡死，本博客针对C#编写界面软件调用C++模型接口情形，分别从C#与C++间图片传输优化、前处理GPU加速、后处理GPU加速三个角度进行优化。

测试服务器：window10、Intel(R) Xeon(R) Gold 5115 CPU @ 2.40GHz、RTX2080 Ti

项目工程可以从github或者gitee上克隆下载：

git clone https://github.com/fyf2022/yolov5_inference.git
git clone https://gitee.com/fyf2022/yolov5_inference.git

1、C#与C++间图片传输优化

方式1：该项目利用数组将C#读取的图片传递给C++，避免编解码方式优化传递速度。

Mat image = new Mat(img_path);
int image_Cols = image.Cols;
int image_Rows = image.Rows;
var image_data = new byte[image.Total() * 3];//这里必须乘以通道数，不然数组越界，也可以用w*h*c
Marshal.Copy(image.Data, image_data1, 0, image_data.Length);

方式2：通常的编解码传输数据到C++接口中的方式耗时：

// 配置图片数据
Mat image = new Mat(image_path);
int max_image_length = image.Cols > image.Rows ? image.Cols : image.Rows;
Mat max_image = Mat.Zeros(new Size(max_image_length, max_image_length), MatType.CV_8UC3);
Rect roi = new Rect(0, 0, image.Cols, image.Rows);
image.CopyTo(new Mat(max_image, roi));
byte[] image_data = new byte[2048 * 2048 * 3];
//存储byte的长度
ulong image_size = new ulong();
image_data = max_image.ImEncode(".bmp");
image_size = Convert.ToUInt64(image_data.Length);

以5472*3648*3的高维图片进行测试，方式1耗时仅为64.4ms，而方式2需要440ms。

2、前处理GPU加速

本项目分别实现cpu、gpu版本前处理，C++核心代码如下：

if (BN_means == 0)
{
    std::vector<float> input_data(node_data_length);
    // 将图像归一化，并放缩到指定大小
    input_image = cv::dnn::blobFromImage(input_image, 1 / 255.0, node_shape, cv::Scalar(0, 0, 0), true, false);
    // 将图片数据copy到输入流中
    memcpy(input_data.data(), input_image.ptr<float>(), node_data_length * sizeof(float));
 
    // 创建cuda流
    cudaStreamCreate(&p->stream);
 
    // 将输入数据由内存到GPU显存
    cudaMemcpyAsync(p->data_buffer[node_index], input_data.data(), node_data_length * sizeof(float), cudaMemcpyHostToDevice, p->stream);
 
    std::vector<float>().swap(input_data);
    input_image.release();
}
else if (BN_means == 1)
{
    std::vector<float> input_data(node_data_length);
    cv::cvtColor(input_image, input_image, cv::COLOR_BGR2RGB); // 将图片通道由 BGR 转为 RGB
    // 对输入图片按照tensor输入要求进行缩放
    cv::resize(input_image, input_image, node_shape, 0, 0, cv::INTER_CUBIC);
    // 图像数据归一化
    input_image.convertTo(input_image, CV_32FC3, 1.0f / 255.0);
    std::vector<cv::Mat> input_channels(3);
    cv::split(input_image, input_channels);
    //std::vector<float> input_data(node_dim.d[2] * node_dim.d[3] * 3);
    auto data = input_data.data();
    int channelLength = INPUT_W * INPUT_H;
    for (int i = 0; i < 3; ++i)
    {
        memcpy(data, input_channels[i].data, channelLength * sizeof(float));
        data += channelLength;
    }
    // 创建cuda流
    cudaStreamCreate(&p->stream);
 
    // 将输入数据由内存到GPU显存
    cudaMemcpyAsync(p->data_buffer[node_index], input_data.data(), node_data_length * sizeof(float), cudaMemcpyHostToDevice, p->stream);
 
    std::vector<float>().swap(input_data);
    input_image.release();
}
else
{
    Npp8u *gpu_img_resize_buf;
    Npp32f *gpu_data_buf;
    Npp32f *gpu_data_planes;
    NppiSize dstSize = { INPUT_W, INPUT_H };;
    NppiRect dstROI = { 0, 0, INPUT_W, INPUT_H };;
    Npp32f m_scale[3] = { 0.00392157, 0.00392157, 0.00392157 };
    int aDstOrder[3] = { 2, 1, 0 };
    CHECK(cudaMalloc(&gpu_img_resize_buf, INPUT_W * INPUT_H * INPUT_C * sizeof(uchar)));
    CHECK(cudaMalloc(&gpu_data_buf, INPUT_W * INPUT_H * INPUT_C * sizeof(float)));
    CHECK(cudaMalloc(&gpu_data_planes, INPUT_W * INPUT_H * INPUT_C * sizeof(float)));
 
    NppiSize srcSize = { img_w, img_h };
    NppiRect srcROI = { 0, 0, img_w, img_h };
    Npp8u *gpu_img_buf;
    CHECK(cudaMalloc(&gpu_img_buf, img_w * img_h * img_c * sizeof(uchar)));
    Npp32f* dst_planes[3] = { gpu_data_planes, gpu_data_planes + INPUT_W * INPUT_H, gpu_data_planes + INPUT_W * INPUT_H * 2 };
 
    cudaMemcpy(gpu_img_buf, img_data, img_w*img_h * img_c, cudaMemcpyHostToDevice);
    nppiResize_8u_C3R(gpu_img_buf, img_w * img_c, srcSize, srcROI, gpu_img_resize_buf, INPUT_W * INPUT_C, dstSize, dstROI, NPPI_INTER_LINEAR);
    nppiSwapChannels_8u_C3IR(gpu_img_resize_buf, INPUT_W * INPUT_C, dstSize, aDstOrder);
    nppiConvert_8u32f_C3R(gpu_img_resize_buf, INPUT_W * INPUT_C, gpu_data_buf, INPUT_W * INPUT_C * sizeof(float), dstSize);
    nppiMulC_32f_C3IR(m_scale, gpu_data_buf, INPUT_W * INPUT_C * sizeof(float), dstSize);
    nppiCopy_32f_C3P3R(gpu_data_buf, INPUT_W * INPUT_C * sizeof(float), dst_planes, INPUT_W * sizeof(float), dstSize);
 
    cudaStreamCreate(&p->stream);
    // 将输入数据由内存到GPU显存
    cudaMemcpyAsync(p->data_buffer[node_index], gpu_data_planes, INPUT_C * INPUT_W * INPUT_H * sizeof(float), cudaMemcpyDeviceToDevice, p->stream);
    CHECK(cudaFree(gpu_img_buf));
    CHECK(cudaFree(gpu_img_resize_buf));
    CHECK(cudaFree(gpu_data_buf));
    CHECK(cudaFree(gpu_data_planes));
}

该项目提供了C#调用C++不同版本前处理的接口，BNFlag.Normal来调用cv::dnn::blobFromImage的cpu前处理操作，BNFlag.fyf_cpu来调用自己实现的cpu版本C++前处理，MMSFlag.fyf_gpu来调用gpu版本前处理。

nvinfer.load_image_data(input_node_name, image_data, image_Cols, image_Rows, BNFlag.Normal);
nvinfer.load_image_data(input_node_name, image_data, image_Cols, image_Rows, BNFlag.fyf_cpu);
nvinfer.load_image_data(input_node_name, image_data, image_Cols, image_Rows, BNFlag.fyf_gpu);

下表给出了前处理耗时对比，从测试结果看gpu前处理优势明显，且在模型输入越大优势越明显：

3、后处理GPU加速

本项目分别实现cpu、gpu版本后处理，C++核心代码如下：

if (NMS_means==0)
{
	// 读取输出数据
	// 创建输出数据
	float* result = new float[node_data_length];
	// 将输出数据由GPU显存到内存
	const clock_t t0 = clock();
	cudaMemcpyAsync(result, p->data_buffer[node_index], node_data_length * sizeof(float), cudaMemcpyDeviceToHost, p->stream);
 
	cv::Mat det_output = cv::Mat(feature_dim, class_num, CV_32F, result);
	 post-process
	std::vector<cv::Rect> position_boxes;
	std::vector<int> classIds;
	std::vector<float> confidences;
	for (int i = 0; i < det_output.rows; i++) {
		float confidence = det_output.at<float>(i, 4);
		if (confidence < 0.2) {
			continue;
		}
		cv::Mat classes_scores = det_output.row(i).colRange(5, class_num);
		cv::Point classIdPoint;
		double score;
		// 获取一组数据中最大值及其位置
		minMaxLoc(classes_scores, 0, &score, 0, &classIdPoint);
		// 置信度 0～1之间
		if (score > 0.25)
		{
			float cx = det_output.at<float>(i, 0);
			float cy = det_output.at<float>(i, 1);
			float ow = det_output.at<float>(i, 2);
			float oh = det_output.at<float>(i, 3);
			int x = static_cast<int>(cx - 0.5 * ow);
			int y = static_cast<int>(cy - 0.5 * oh);
			int width = static_cast<int>(ow);
			int height = static_cast<int>(oh);
			cv::Rect box;
			box.x = x;
			box.y = y;
			box.width = width;
			box.height = height;
 
			position_boxes.push_back(box);
			classIds.push_back(classIdPoint.x);
			confidences.push_back(score);
		}
	}
	// NMS
	std::vector<int> indexes;
	cv::dnn::NMSBoxes(position_boxes, confidences, 0.25, 0.45, indexes);
 
	for (size_t i = 0; i < indexes.size(); i++) {
		int index = indexes[i];
		*output_result = classIds[index];
		output_result++;
		*output_result = confidences[index] * 1000;
		output_result++;
		*output_result = position_boxes[index].tl().x;
		output_result++;
		*output_result = position_boxes[index].tl().y;
		output_result++;
		*output_result = position_boxes[index].width;
		output_result++;
		*output_result = position_boxes[index].height;
		output_result++;
	}
	delete[] result;
	std::vector<cv::Rect>().swap(position_boxes);
	std::vector<int>().swap(classIds);
	std::vector<float>().swap(confidences);
}
else
{
	float confidence_threshold = 0.25f; 
	float nms_threshold = 0.45f;
	vector<Box> box_result;
	cudaStream_t stream = nullptr;
	checkRuntime(cudaStreamCreate(&stream));
 
	float* output_device = nullptr;
	float* output_host = nullptr;
	int max_objects = 1000;
	int NUM_BOX_ELEMENT = 7;  // left, top, right, bottom, confidence, class, keepflag
	checkRuntime(cudaMalloc(&output_device, sizeof(float) + max_objects * NUM_BOX_ELEMENT * sizeof(float)));
	//output_host = (float*)malloc(sizeof(float) + max_objects * NUM_BOX_ELEMENT * sizeof(float));
	//memset(output_host, 0, sizeof(float) + max_objects * NUM_BOX_ELEMENT * sizeof(float));
	checkRuntime(cudaMallocHost(&output_host, sizeof(float) + max_objects * NUM_BOX_ELEMENT * sizeof(float)));
 
	//checkRuntime(cudaStreamSynchronize(stream));
	cudaDeviceSynchronize();
	
	decode_kernel_invoker(
		(float *)p->data_buffer[node_index], feature_dim, class_num - 5, confidence_threshold,
		nms_threshold, nullptr, output_device, max_objects, NUM_BOX_ELEMENT, stream
	);
	
	
	checkRuntime(cudaMemcpy(output_host, output_device,
		sizeof(int) + max_objects * NUM_BOX_ELEMENT * sizeof(float),
		cudaMemcpyDeviceToHost));
	
	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];
		if (keep_flag) {
			*output_result = (int)ptr[5];
			output_result++;
			*output_result = (int)(ptr[4] * 1000);
			output_result++;
			*output_result = (int)ptr[0];
			output_result++;
			*output_result = (int)ptr[1];
			output_result++;
			*output_result = (int)(ptr[2] - ptr[0]);
			output_result++;
			*output_result = (int)(ptr[3] - ptr[1]);
			output_result++;
			// left, top, right, bottom, confidence, class, keepflag
		}
	}
	
	/*
	checkRuntime(cudaStreamDestroy(stream));
	//checkRuntime(cudaFree(output_device));
	//free(output_host);
	checkRuntime(cudaFreeHost(output_host));
	*/
}

 该项目提供了C#调用C++不同版本前处理的接口，MMSFlag.fyf_cpu来调用cpu版本C++后处理，MMSFlag.fyf_gpu来调用gpu版本后处理

int[] result_array = nvinfer.read_infer_result(output_node_name, MMSFlag.fyf_cpu);
int[] result_array = nvinfer.read_infer_result(output_node_name, MMSFlag.fyf_gpu);

下表给出了后处理耗时对比，从测试结果看gpu后处理优势明显，且在检测目标越多优势越明显：

参考

[1] GitHub - guojin-yan/Inference: At OpenVINO ™、 TensorRT, ONNX runtime, and OpenCV Dnn deployment platforms are based on C # language deployment models.

[2] CPU/GPU(CUDA)版本的 YOLOv5后处理代码-CSDN博客