用OpenCL实现HEVC中ME模块的测试数据分析
来源:互联网 发布:ubuntu ftp 不能上传 编辑:程序博客网 时间:2024/06/08 16:02
使用opencl来实现编码算法中运动搜索模块!
下面测试数据时在GTX570上的测试结果:
LCU为32x32, 100帧720P, CPU上纯C算法使用搜索时间是67s, GPU上是0.915s
LCU为16x16, CPU 是76.8s, GPU上是1.6s
LCU为8x8, CPU 是82.5s, GPU上是4.2s
同样的程序, CPU改为SSE实现, GPU做一个小的改动, 使用缩减算法! 结果如下:
从上面数据可以看出, sse 比C语言快5倍左右, 新的GPU kenel快了20%左右, 其中LCU为8x8的快了好几倍!
综合看来
OPENCL实现 比C语言实现接近100倍的级别, 比SSE快了接近20倍左右!
另外提一句, 如果OPENCL不适用__local 内存的话, 会慢一半!
下面贴出部分代码供参考:
#define SearchRange 16#define Edge_SIZE_T 48//32x32 version of kernel__kernel void opencl_me_32x32(const __global short* p_ref, __global short* p_cur, __global int* outputBuf, __local int* local_refBuf, __local int* local_curBuf, __local int* mv_cost){int searchrange = SearchRange;int edeg = Edge_SIZE_T;int width = get_global_size(0);int height = get_global_size(1);int block_w = get_local_size(0);int block_h = get_local_size(1);int local_x = get_local_id(0);int local_y = get_local_id(1);int lcu_x = get_group_id(0);int lcu_y = get_group_id(1);int stride = width + 2 * edeg;int lcu_adr_offset = edeg * stride + edeg;int local_refBuf_stride = block_w + 2 * searchrange;//LCU blcok adrlcu_adr_offset += lcu_y * stride * block_h + lcu_x * block_w;int ref_lcu_adr_offset = lcu_adr_offset - searchrange - searchrange * stride;//thread adrint global_thread_adr_offset = local_y * stride + local_x;int thread_adr_offset = local_y * local_refBuf_stride + local_x;local_curBuf[local_y * block_w + local_x]= p_cur[lcu_adr_offset + global_thread_adr_offset];local_refBuf[thread_adr_offset]= p_ref[ref_lcu_adr_offset + global_thread_adr_offset ];local_refBuf[thread_adr_offset + block_w]= p_ref[ref_lcu_adr_offset + global_thread_adr_offset + block_w];local_refBuf[thread_adr_offset + local_refBuf_stride * block_h]= p_ref[ref_lcu_adr_offset + global_thread_adr_offset + stride * block_h];local_refBuf[thread_adr_offset + local_refBuf_stride * block_h + block_w]= p_ref[ref_lcu_adr_offset + global_thread_adr_offset + stride * block_h + block_w];
barrier(CLK_LOCAL_MEM_FENCE); { int i; int uiSum = 0; for( int i = 0; i < block_h; i++ ) {
计算sad
}
{
比较最小SAD 保存bestcost
}
if((local_y ==0) && (local_x == 0)) { int best_sad = mv_cost[local_y*2*SearchRange*3 + local_x*3 + 2]; int best_mvx = mv_cost[local_y*2*SearchRange*3 + local_x*3 + 0]; int best_mvy = mv_cost[local_y*2*SearchRange*3 + local_x*3 + 1]; outputBuf[(lcu_y * get_num_groups(0) + lcu_x)*3 + 0] = best_mvx; outputBuf[(lcu_y * get_num_groups(0) + lcu_x)*3 + 1] = best_mvy; outputBuf[(lcu_y * get_num_groups(0) + lcu_x)*3 + 2] = best_sad; //printf("\nxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx = %d, y = %d, sad = %d",best_mvx, best_mvy, best_sad); }
- 用OpenCL实现HEVC中ME模块的测试数据分析
- 用OpenCL实现HEVC中ME模块的测试数据分析
- HEVC中比较耗时的模块
- OpenCV中OpenCL模块函数
- HEVC ME之 整像素TZ搜索代码分析
- HEVC ME之 整像素TZ搜索代码分析
- HEVC ME之 H.264论文综述与代码实现
- HEVC/H.265 的未来必须是使用并行处理(OpenCL?) OpenCV和OpenCL区别
- HEVC/H.265 的未来必须是使用并行处理(OpenCL?) OpenCV和OpenCL区别
- HEVC/H.265 的未来必须是使用并行处理(OpenCL?) OpenCV和OpenCL区别
- 【HEVC学习与研究】41、HEVC帧内编码的原理和实现(中)
- 【HEVC学习与研究】46、HEVC参考代码中SAO的实现
- 使用opencl实现编码算法中运动搜索模块速度对比
- Nginx 推送模块性能的测试数据
- HEVC CU 级别的分析
- HEVC CU 级别的分析
- 使用OpenCV的OpenCL(ocl)模块
- 使用OpenCV的OpenCL(ocl)模块
- 线段树之Poj3264
- 一步一步学习Ubuntu之安装篇
- Android中Activity启动模式详解
- Android 4.2 API 变更
- 让VS调试器帮你格式化显示自定义数据
- 用OpenCL实现HEVC中ME模块的测试数据分析
- win7浏览器IE9升级到IE10后无法启动的问题
- 在线升级uboot,内核和文件系统
- Android 数据持久化
- get/post本质区别
- Android异步加载图像(含线程池,缓存方法)
- centos install redis
- NYOJ - 郁闷的C小加(二)
- oozie上传实例解读