OpenCL：减少示例，并保留内存对象/将 cuda 代码转换为 openCL

Question

我经历了几个例子，将一组元素减少为一个元素，但没有成功。有人在 NVIDIA 论坛上发布了这个。我已从浮点变量更改为整数。

__kernel void sum(__global const short *A,__global unsigned long  *C,uint size, __local unsigned long *L) {
            unsigned long sum=0;
            for(int i=get_local_id(0);i<size;i+=get_local_size(0))
                    sum+=A[i];
            L[get_local_id(0)]=sum;

            for(uint c=get_local_size(0)/2;c>0;c/=2)
            {
                    barrier(CLK_LOCAL_MEM_FENCE);
                    if(c>get_local_id(0))
                            L[get_local_id(0)]+=L[get_local_id(0)+c];

            }
            if(get_local_id(0)==0)
                    C[0]=L[0];
            barrier(CLK_LOCAL_MEM_FENCE);
}

这看起来对吗？第三个参数“大小”，应该是本地工作大小还是全局工作大小？

我的论点是这样设置的，

clSetKernelArg(ocReduce, 0, sizeof(cl_mem), (void*) &DevA);
clSetKernelArg(ocReduce, 1, sizeof(cl_mem), (void*) &DevC); 
clSetKernelArg(ocReduce, 2, sizeof(uint),   (void*) &size);  
clSetKernelArg(ocReduce, 3, LocalWorkSize * sizeof(unsigned long), NULL); 

第一个参数是输入，我试图保留在它之前启动的内核的输出。

clRetainMemObject(DevA);
clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocKernel, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL);
//the device memory object DevA now has the data to be reduced

clEnqueueNDRangeKernel(hCmdQueue[Plat-1][Dev-1], ocReduce, 1, NULL, &GlobalWorkSize, &LocalWorkSize, 0, NULL, NULL);
clEnqueueReadBuffer(hCmdQueue[Plat-1][Dev-1],DevRE, CL_TRUE, 0, sizeof(unsigned long)*512,(void*) RE , 0, NULL, NULL);

今天我打算尝试将以下 cuda 缩减示例转换为 openCL。

__global__ voidreduce1(int*g_idata, int*g_odata){
extern __shared__ intsdata[];

unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;
sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];
__syncthreads();


for(unsigned int s=blockDim.x/2; s>0; s>>=1) {
if (tid < s) {
sdata[tid] += sdata[tid + s];
}
__syncthreads();
}

// write result for this block to global mem
if(tid == 0) g_odata[blockIdx.x] = sdata[0];
}

有一个更优化的，（完全展开+每个线程多个元素）。

http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf

这可以使用 openCL 吗？

前几天灰熊给了我这个建议，

"...使用对 n 元素进行操作的缩减内核并将它们缩减为类似 n / 16 （或任何其他数字）的值。然后您迭代地调用该内核，直到您缩减到一个元素，这就是您的结果"

我也想试试这个，但我不知道从哪里开始，我想先做点什么。

Answer 1

只要只有一个工作组在进行缩减工作，您提供的第一个缩减代码就应该可以工作（所以get_global_size(0) == get_local_size(0)）。在这种情况下，内核的size 参数将是A 中的元素数量（与全局或本地工作大小没有真正的相关性）。虽然这是一个可行的解决方案，但在进行归约时让大部分 gpu 闲置似乎是天生的浪费，这正是我建议迭代调用归约内核的原因。只需对代码稍作修改即可实现这一点：

__kernel void sum(__global const short *A, __global unsigned long  *C, uint size, __local unsigned long *L) {
        unsigned long sum=0;
        for(int i=get_global_id(0); i < size; i += get_global_size(0))
                sum += A[i];
        L[get_local_id(0)]=sum;

        for(uint c=get_local_size(0)/2;c>0;c/=2)
        {
                barrier(CLK_LOCAL_MEM_FENCE);
                if(c>get_local_id(0))
                        L[get_local_id(0)]+=L[get_local_id(0)+c];

        }
        if(get_local_id(0)==0)
                C[get_group_id(0)]=L[0];
        barrier(CLK_LOCAL_MEM_FENCE);
}

使用小于size（例如4）的GlobalWorkSize 调用它会将A 中的输入减少4*LocalWorkSize 的因子，可以迭代（通过使用输出缓冲区作为下一次使用不同的输出缓冲区调用sum。实际上这并不完全正确，因为第二次（以及随后的所有）迭代需要A 的类型为global const unsigned long*，所以你实际上需要内核，但你明白了。

关于 cuda 缩减样本：你为什么要转换它，它的工作原理基本上与我在上面发布的 opencl 版本一样，除了每次迭代仅减少硬编码的大小（2*LocalWorkSize insted of size/GlobalWorkSize*LocalWorkSize）。

我个人使用几乎相同的方法来减少，尽管我将内核分成两部分并且仅在最后一次迭代中使用使用本地内存的路径：

__kernel void reduction_step(__global const unsigned long* A, __global unsigned long  * C, uint size) {
        unsigned long sum=0;
        for(int i=start; i < size; i += stride)
                sum += A[i];
        C[get_global_id(0)]= sum;
}

对于最后一步，使用了在工作组内部进行缩减的完整版本。当然，您需要reduction step 的第二个版本以使用global const short*，并且此代码是您的代码未经测试的改编（遗憾的是，我无法发布自己的版本）。这种方法的优点是内核完成大部分工作的复杂性要低得多，并且由于分支不同，wasted work 的数量也更少。这使它比其他变体快一点。但是，对于最新的编译器版本和最新的硬件，我都没有结果，因此该点可能不再正确，也可能不再正确（尽管我怀疑它可能是由于不同分支数量的减少）。

现在对于您链接的论文：当然可以在 opencl 中使用该论文中建议的优化，除了使用 opencl 不支持的模板，因此必须对块大小进行硬编码。当然，opencl 版本已经对每个内核进行了多次添加，如果您按照我上面提到的方法，将不会真正受益于通过本地内存展开减少，因为这仅在最后一步中完成，不应该采取对于足够大的输入，占整个计算时间的重要部分。此外，我发现展开的实现中缺乏同步有点麻烦。这只有效，因为进入该部分的所有线程都属于同一个经线。然而，在除当前 nvidia 卡（未来的 nvidia 卡、amd 卡和 cpu 之外）之外的任何硬件上执行时，这不是必需的（尽管我认为它应该适用于当前的 amd 卡和当前的 cpu 实现，但我不一定指望它）），所以我会远离它，除非我需要绝对的最后一点速度来减少（然后仍然提供通用版本并在我不识别硬件或类似的东西时切换到该版本）。

Answer 2

还原内核在我看来是正确的。在缩减中，大小应该是输入数组A 的元素个数。代码在sum 中累积每个线程的部分和，然后执行本地内存（共享内存）缩减并将结果存储到C。您将在每个本地工作组获得C 的部分金额。要么与一个工作组再次调用内核以获得最终答案，要么在主机上累积部分结果。