opencl:提高kernel有效并发约束(Effective concurrency constraint)的一点心得

版权声明:本文为博主原创文章,转载请注明源地址。 https://blog.csdn.net/10km/article/details/51352893

下面是一个简单的kernel函数,从integ_count_mat矩阵中用vload8函数取出A1,A2,A3,A4四个向量执行A4+A1-A2-A3,结果存入density_mat,代码中只用到了一个向量类型的变量sum

__kernel void object_density_filter( 
          matrix_info_cl im_info
        , const __global  ushort *integ_count_mat
        , matrix_info_cl om_info
        , __global ushort8* density_mat
        ,int  face_dist_size
        ,ushort sum_threshold
         ){
    ushort8 sum;
    int start_y=(int)get_global_id(1),x_v=get_global_id(0);
    density_mat+= start_y * om_info.row_stride>>3;
    int index_a1=start_y*im_info.row_stride;                    //A1
    int index_a2=index_a1+face_dist_size;                       //A2
    int index_a3=(start_y+face_dist_size)*im_info.row_stride;   //A3
    int index_a4=index_a3+face_dist_size;                       //A4

    // compute faces sum in each window which size speciialed win_size : A4+A1-A2-A3
    sum=( vload8(x_v,integ_count_mat +index_a4) + vload8(x_v,integ_count_mat+index_a1)  
        - vload8(x_v,integ_count_mat +index_a2) - vload8(x_v,integ_count_mat+index_a3) );
    // fake object filter by density_const_param.sum_threshold
    density_mat[x_v]=sum>=sum_threshold?sum:(ushort8)(0);
}

但是使用CodeXL进行静态代码分析显示,这个kernel居然用到41个VGPRS(向量寄存器)!因此导致有效并发约束(Effective concurrency constraint(Max waves per SIMD))只能为5,怎么修改代码都无法提高。
这里写图片描述

我代码中明明只有一个向量类型的变量啊,这多出来的40个VGPRS用到哪里了?排除自己代码中使用向量寄存器的可能后,我怀疑到了vload8这个函数,vload8是opencl built-in函数,怎么实现的我们是不知道的。极有可能是它用了大量VGPRS,于我尝试把这行代码

sum=( vload8(x_v,integ_count_mat +index_a4) + vload8(x_v,integ_count_mat+index_a1)  
    - vload8(x_v,integ_count_mat +index_a2) - vload8(x_v,integ_count_mat+index_a3) );

改为

sum-= sum,//将sum赋值为0,使用sum=0这种方式初始化,一样会增加VGPRS的使用数量
sum+= vload8(x_v,integ_count_mat +index_a4),
sum+= vload8(x_v,integ_count_mat +index_a1),
sum-= vload8(x_v,integ_count_mat +index_a2),
sum-= vload8(x_v,integ_count_mat +index_a3);    

上面的公式等价于A4+A1-A2-A3,只是利用,号运算符强制让4个vload8函数串行执行,这样就防止了编译器对A4+A1-A2-A3这个公式进行并行优化。
于是立即就有了效果,
这里写图片描述
使用的VGPRS降到了6个。并发限制提高到了满格10.
总结:
在表达式中使用built-in函数时,要注意表达式的书写方式,过长的表达式,要拆分开,尽可能用,操作符进行改造,以防止编译器进行并行做优化,因为一表达式的built-in函数并行执行,所使用的寄存器数量就会成倍上升。

阅读更多
想对作者说点什么?

博主推荐

换一批

没有更多推荐了,返回首页