织梦CMS - 轻松建站从此开始!

罗索实验室

用OpenCL实现HEVC中ME模块的测试数据分析

落鹤生 发布于 2015-06-23 13:37 点击:次 
从上面数据可以看出, sse 比C语言快5倍左右, 新的GPU kenel快了20%左右, 其中LCU为8x8的快了好几倍!综合看来OPENCL实现 比C语言实现接近100倍的级别, 比SSE快了接近20倍左右!另外提一句, 如果OPENCL不适用__local 内存的话, 会慢一半!
TAG: OpenCL  HEVC  

使用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 内存的话, 会慢一半!

下面贴出部分代码供参考:

  1. #define SearchRange 16 
  2. #define Edge_SIZE_T 48 
  3. //32x32 version of kernel 
  4. __kernel void opencl_me_32x32(const __global short* p_ref, __global short* p_cur
  5. , __global int* outputBuf, __local int* local_refBuf, __local int* local_curBuf
  6. , __local int* mv_cost) 
  7.     int searchrange = SearchRange; 
  8.     int edeg = Edge_SIZE_T; 
  9.     int width = get_global_size(0); 
  10.     int height = get_global_size(1); 
  11.     int block_w = get_local_size(0); 
  12.     int block_h = get_local_size(1); 
  13.     int local_x = get_local_id(0); 
  14.     int local_y = get_local_id(1); 
  15.     int lcu_x = get_group_id(0); 
  16.     int lcu_y = get_group_id(1); 
  17.     int stride = width + 2 * edeg; 
  18.     int lcu_adr_offset = edeg * stride + edeg; 
  19.     int local_refBuf_stride = block_w + 2 * searchrange; 
  20.              
  21.     //LCU blcok adr 
  22.     lcu_adr_offset += lcu_y * stride * block_h + lcu_x * block_w; 
  23.     int ref_lcu_adr_offset = lcu_adr_offset - searchrange - searchrange * stride; 
  24.     //thread adr 
  25.     int global_thread_adr_offset = local_y * stride + local_x; 
  26.     int thread_adr_offset  = local_y * local_refBuf_stride + local_x; 
  27.  
  28.     local_curBuf[local_y * block_w + local_x]  
  29. = p_cur[lcu_adr_offset + global_thread_adr_offset]; 
  30.  
  31.     local_refBuf[thread_adr_offset]   
  32. = p_ref[ref_lcu_adr_offset + global_thread_adr_offset ]; 
  33.     local_refBuf[thread_adr_offset + block_w] 
  34. = p_ref[ref_lcu_adr_offset + global_thread_adr_offset + block_w]; 
  35.     local_refBuf[thread_adr_offset + local_refBuf_stride * block_h]
  36. = p_ref[ref_lcu_adr_offset + global_thread_adr_offset + stride * block_h]; 
  37.     local_refBuf[thread_adr_offset + local_refBuf_stride * block_h + block_w]
  38. = p_ref[ref_lcu_adr_offset + global_thread_adr_offset + stride * block_h + block_w]; 

 

  1. barrier(CLK_LOCAL_MEM_FENCE);   
  2. {   
  3.   int i;   
  4.   int uiSum = 0;   
  5.   forint i = 0; i < block_h; i++ )   
  6. {   
  7. 计算sad 
  8. else 
  9. 比较最小SAD 保存bestcost  
  10.  
  11. if((local_y ==0) && (local_x == 0))   
  12.  {   
  13.   int best_sad  = mv_cost[local_y*2*SearchRange*3 + local_x*3 + 2];   
  14.   int best_mvx  = mv_cost[local_y*2*SearchRange*3 + local_x*3 + 0];   
  15.   int best_mvy  = mv_cost[local_y*2*SearchRange*3 + local_x*3 + 1];   
  16.   outputBuf[(lcu_y * get_num_groups(0) + lcu_x)*3 + 0] = best_mvx;   
  17.   outputBuf[(lcu_y * get_num_groups(0) + lcu_x)*3 + 1] = best_mvy;   
  18.   outputBuf[(lcu_y * get_num_groups(0) + lcu_x)*3 + 2] = best_sad;    
  19.   //printf("\nxxxxxxxxxxxxxxxxxxxxxxxxxxxxxx = %d, y =%d
  20. //, sad = %d",best_mvx, best_mvy, best_sad);   
  21.  }   

 

(feixiang_john)
本站文章除注明转载外,均为本站原创或编译欢迎任何形式的转载,但请务必注明出处,尊重他人劳动,同学习共成长。转载请注明:文章转载自:罗索实验室 [http://www1.rosoo.net/a/201506/17340.html]
本文出处:CSDN博客 作者:feixiang_john 原文
顶一下
(0)
0%
踩一下
(0)
0%
------分隔线----------------------------
发表评论
请自觉遵守互联网相关的政策法规,严禁发布色情、暴力、反动的言论。
评价:
表情:
用户名: 验证码:点击我更换图片
栏目列表
将本文分享到微信
织梦二维码生成器
推荐内容