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

关于cuda缩减示例：你为什么要麻烦转换它，它的工作原理基本上和我上面发布的opencl版本完全一样，只是每次迭代只减少一个硬编码大小（2*LocalWorkSize而不是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;
}

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

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

Answer 2

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