2

我一直在经历一些例子，将一组元素减少为一个元素，但没有成功。有人在 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（或任何其他数字）。然后您迭代地调用该内核，直到您缩减到一个元素，这就是您的结果”

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

4

2 回答 2

8

只要只有一个工作组在减少工作，您提供的第一个减少代码就应该工作（所以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);
}

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

关于 cuda 缩减样本：你为什么要转换它，它的工作原理基本上与我上面发布的 opencl 版本完全一样，除了每次迭代仅减少硬编码的大小（2*LocalWorkSizeinsted 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 版本已经对每个内核进行了多次添加，如果您按照我上面提到的方法，将不会真正受益于通过本地内存展开减少，因为这仅在最后一步中完成，不应该采取足够大的输入的整个计算时间的重要部分。此外，我发现展开的实现中缺乏同步有点麻烦。这只有效，因为进入该部分的所有线程都属于同一个经线。然而，这不是

于 2012-01-14T20:27:49.513 回答
1

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

于 2012-01-14T20:26:43.730 回答