OpenCL代码的优化？

Question

我在OpenCL中有一段代码，它以4为步长添加数组元素，从给定点开始递减。

这是代码- rp是源数组，out是输出数组)：

__kernel void subFilter(__global unsigned char* rp,__global unsigned char *out,int istop,int bpp)
{
    int gid = get_global_id(0);//add the offset by bpp to access the next gid
    int i;
    unsigned char temp=0;

    if(gid>=bpp){
         i=gid;
        while(i>=0)
        {
            if((temp + rp[i])>255)
            {
                temp = temp - 256;
                temp=temp + rp[i];
            }
            else
            {
                temp=temp+rp[i];
            }
            i=i-bpp;
        }
        out[gid]=(temp & 0xff);  //masked
        //rp[gid]=66;
    }
    else if(gid<bpp)
    out[gid]=rp[gid];
}

现在，这种方法工作得很好，但比在CPU上运行相同的代码需要更多的时间。如果删除循环，可能会更快？

执行时，它的工作方式如下所示：

rp[0] = rp[0];        // same for rp[1],rp[2],rp[3]
rp[4] = rp[4]+rp[4-4] // steps of 4 till zero
// ...
rp[16]= rp[16]+rp[12]+rp[8]+rp[4]+rp[0]; // etc.

生成的rp[i]保存在代码中的temp中，最后保存在out[gid]中。

所以，要添加这个序列，需要循环，并且需要很多时间……如果有某种方法，可以缓存之前的和，或者可以完全删除循环，那就太好了。

我如何改进这段代码以消除循环？

Answer 1

首先，我建议您消除可能的uchar溢出：

__kernel void subFilter(
    __global unsigned char  *rp,
    __global unsigned char  *out,
    int                     istop,
    int                     bpp)
{
    int gid = get_global_id(0), i = gid;
    unsigned char temp = 0;

    if(gid>=bpp){
        i=gid;

        while(i>=0)
        {
            if(temp > 255 - rp[i])
            {
                temp -= 255 - rp[i];
            }
            else
            {
                temp += rp[i];
            }

            i -= bpp;
        }

        //masked
        out[gid]=(temp & 0xff);  
    }
    else if(gid<bpp){
        out[gid]=rp[gid];
    }
}

回到过滤算法。每个工作项都有其唯一的ID，它被用作循环计数器。一般来说，这是不好的，因为不同的指令集应用于相同的数据，这违反了SIMD的思想。这可能会导致线程分歧和性能下降。你的算法的本质是什么？请详细描述。可能存在支持GPU的实现。

Answer 2

消除循环中的条件代码可能会提高性能：

__kernel void subFilter(__global unsigned char* rp,__global unsigned char *out,int istop,int bpp)
{
    int gid = get_global_id(0);//add the offset by bpp to access the next gid
    int i;
    unsigned char temp=0;

    out[gid]=rp[gid];
    if(gid>=bpp){
        i=gid;
        while(i>=0)
        {
            temp+=rp[i];
            i-=bpp;
        }
        out[gid]=(temp & 0xff);  //masked
    }
}