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

下面是一个简单的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函数并行执行，所使用的寄存器数量就会成倍上升。