基于OpenCL的yolov3 后处理优化

基于OpenCL的yolov3 后处理代码性能优化

背景:
目前对Yolov3的整体架构不是很了解,在网上也看了很多教程,对我来说,讲的比较复杂,可能是对很多的术语没有概念。不过Yolov3代码都是C编写的, 对我来说只要是C代码,就不是问题,针对问题,以及解决方法,改代码就好。
我拿到的yolov3的代码可能不是标准的yolov3 , 不过大致处理逻辑是不变的。
yolov3 检测部分有两个部分比较耗时:
第一个耗时点在,网络数据是以uint8类型保存的,那么如果想对数据进行处理,那么一定要讲数据转化成float32 类型的数据。在CPU处理进行转换逻辑中,是以for循环的形式一个一个处理的,那么当数据很大的时候,比如典型的网络数据1313255、2616255、5252255,这么多数据在for循环中,一个一个进行数据类型转换显然是非常耗时的,在我们的平台上验证:耗时为80ms。
另外一个耗时点在,数据转换成float型后,在进行数据处理的过程中(专业的叫法我确实不太清楚),主要的功能是,找到数据的坐标位置,对接下来的数据进行exp指数运算,也是通过for 循环一个一个调用exp进行运算的,CPU 做指数运算,效率更低。测试结果为: 114ms

那么这些处理明显是可以放在GPU 进行并行运算,时间便可以大大缩短,在极短的时间内,几毫米内,就完成了运算。

经过对代码的逻辑分析,这两部分完全可以放在一起,只需要一个kernel便可以完成运算。

#define TENSOR_NUM 3
	/*tensor 1*/
	vsi_nn_tensor_t *tensor[TENSOR_NUM];
	uint8_t *tensor_data[TENSOR_NUM];		//tensor 原始数据内存 uint8 : buff
	float *convert_result[TENSOR_NUM];		//原始数据 uint8->float	: buf
	float *calculation_result[TENSOR_NUM];	//做exp运算结果      : buf
	uint32_t data_num[TENSOR_NUM] = {1,1,1};	//tensor 数据大小  //BUG sloved:此处没有初始化,导致下面算size异常,内存分配失败 
	for(int i=0; i<TENSOR_NUM; i++)
	{
		tensor[i] = vsi_nn_GetTensor(graph, graph->output.tensors[i]);
		for(int j=0; j<tensor[i]->attr.dim_num; j++)
		{
			data_num[i] *= tensor[i]->attr.size[j];			
		}
		tensor_data[i] = (uint8_t *)vsi_nn_ConvertTensorToData(graph, tensor[i]);
		/*内存分配*/ /*TODO: 释放*/
		calculation_result[i] = (float*)malloc(sizeof(float) * data_num[i] );
		convert_result[i] = (float *)malloc(sizeof(float) * data_num[i]);
		
	}
	/*opencl init*/
	pcl_controller controller = (pcl_controller)malloc(sizeof(cl_controller));
	memset(controller,0,sizeof(cl_controller));
	cl_controller_init(controller);
	/*opencl run*/
	for(int i=0; i<TENSOR_NUM; i++)
	{
		cl_controller_set(controller, 								//opencl 控制器
						  tensor_data[i],							// buffer
						  convert_result[i],				
						  calculation_result[i],			 
						  data_num[i], 								//数据number
						  tensor[i]->attr.dtype.zero_point, 		//用于数据类型转换的参数
						  tensor[i]->attr.dtype.scale);				//同上
		cl_controller_run(controller, data_num[i]);					//opencl kernel run
		cl_controller_get(controller,								//opencl 控制器
						  convert_result[i] ,						//数据类型转换结果
						  calculation_result[i],					//exp结果
						  data_num[i]);								//数据number
	}
	/*opencl clean*/
	cl_controller_clean(controller);
	status = show_result(graph,image_name,convert_result,calculation_result);
	/*free buf*/
	for(int i=0;i<3;i++)
	{
		free(calculation_result[i]);
		free(convert_result[i]);
		vsi_nn_Free(tensor_data[i]);
	}
__kernel void fun_exp(__global unsigned char* data ,\
					  __global float* convert_data,\
					  __global float* result,\
					  int zero_point,\
					  float scale)
{
	int gid = get_global_id(0);

	float x = ((float)data[gid] - zero_point) * scale ;
	convert_data[gid] = x;
	result[gid]=1.f/(1.f + exp(-x));
	
}

kernel 的功能很简单,将uint8的数据转化为float型保存在内存,然后指数运算的结果保存在林外一个数组。在主机端进行读取。

由于解除yolov3 和opencl 时间不长, 代码写的比较垃圾,不过效果确实达到了, 原来处理完上面两部分代码耗时就需要200ms, 进过opencl 的优化, 只需要24ms 就能处理完, 这不是GPU的运算时间,GPU几毫秒以内就完成了,这里是CPU 中对结果进行memcpy等操作的耗时。我也有写memcpy的kernel,但是实际效果不好,就没有写进代码。

效果: 200ms -> 24ms 8倍!

官方有GPU版本的yolov3 不过是以CUDA 实现的。胜称整体效率是CPU的500倍。我这边没有测试环境,就没有进行测试。我也找到了一版完全用opencl 实现的yolov3 , 没有运行环境。 可以提供给大家。

  • 3
    点赞
  • 3
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 1
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

疯狂的蕉尼基

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值