基于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]);
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
__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));
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
kernel 的功能很简单,将uint8的数据转化为float型保存在内存,然后指数运算的结果保存在林外一个数组。在主机端进行读取。
由于解除yolov3 和opencl 时间不长, 代码写的比较垃圾,不过效果确实达到了, 原来处理完上面两部分代码耗时就需要200ms, 进过opencl 的优化, 只需要24ms 就能处理完, 这不是GPU的运算时间,GPU几毫秒以内就完成了,这里是CPU 中对结果进行memcpy等操作的耗时。我也有写memcpy的kernel,但是实际效果不好,就没有写进代码。
效果: 200ms -> 24ms 8倍!
官方有GPU版本的yolov3 不过是以CUDA 实现的。胜称整体效率是CPU的500倍。我这边没有测试环境,就没有进行测试。我也找到了一版完全用opencl 实现的yolov3 , 没有运行环境。 可以提供给大家。
————————————————
版权声明:本文为CSDN博主「疯狂的蕉尼基」的原创文章,遵循CC 4.0 BY-SA版权协议,转载请附上原文出处链接及本声明。
原文链接:https://blog.csdn.net/qq_38505858/article/details/121978416