在 OpenCL 中使用本地内存进行并行缩减

Question

我在 OpenCL 中实现了一个 reduce 内核来汇总大小为 N 的 input 向量中的所有条目。为了更容易测试，我用 1.0f 初始化 input 向量。所以结果应该是N。但事实并非如此！

这是我的 reduce-内核：

kernel void reduce(global float* input, global float* output, const unsigned int N, local float* cache)
{
    const uint local_id = get_local_id(0);
    const uint global_id = get_global_id(0);
    const uint local_size = get_local_size(0);

    cache[local_id] = (global_id < N) ? input[global_id] : 0.0f;
    barrier(CLK_LOCAL_MEM_FENCE);

    for (unsigned int s = local_size >> 1; s > 0; s >>= 1) {
        if (local_id < s) {
            cache[local_id] += cache[local_id + s];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if (local_id == 0) output[local_size] = cache[0];
}

这里是 OpenCL 的设置：

 const uint N = 8196;

 cl_float a[N];
 cl_float b[N];

 for (uint i=0; i<N; i++) {
      a[i] = 1.0f;
      b[i] = 0.0f;
 }

 cl::Buffer inputBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float)*N);
 cl::Buffer resultBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_float)*N);

 queue.enqueueWriteBuffer(inputBuffer, CL_TRUE, 0, sizeof(cl_float)*N, a);
 queue.enqueueWriteBuffer(resultBuffer, CL_TRUE, 0, sizeof(cl_float)*N, b);

 cl::Kernel addVectorKernel = cl::Kernel(program, "reduce");

 size_t localSize = addVectorKernel.getWorkGroupInfo<CL_KERNEL_WORK_GROUP_SIZE>(device); // e.g. => 512

 size_t globalSize = roundUp(localSize, N); // rounds up to a multiple of localSize

 addVectorKernel.setArg(0, inputBuffer);
 addVectorKernel.setArg(1, resultBuffer);
 addVectorKernel.setArg(2, N);
 addVectorKernel.setArg(3, (sizeof(cl_float) * localSize), NULL);


 queue.enqueueNDRangeKernel(
      addVectorKernel,
      cl::NullRange,    
      cl::NDRange(globalSize), 
      cl::NDRange(localSize)     
 );
 queue.finish(); // wait for ending

 queue.enqueueReadBuffer(resultBuffer, CL_TRUE, 0, sizeof(cl_float)*N, b); // e.g. => 1024

结果取决于工作组大小。我究竟做错了什么？是内核本身还是 OpenCL 的设置？

Answer 1

在循环的归约中，你需要这样：

for(unsigned int s = localSize >> 1; s > 0; s >>= 1)

您在初始化 s 时多移动了一位。

修复后，让我们看看您的内核在做什么。主机代码以 8192 的 globalSize 和 512 的 localSize 执行它，这导致 16 个工作组。在内核中，您首先将索引 2*global_id 处的两个连续内存位置的数据相加。对于 ID 为 15 的工作组，工作项 0，将位于索引 15*512*2 = 15,360 和 15,361，这超出了输入数组的边界。我很惊讶你没有崩溃。同时，这也解释了为什么你的价值是你期望的两倍。

要修复它，您可以这样做：

cache[localID] = input[globalID];

或者指定一个全局大小，它是当前大小的一半。

Answer 2

将总和写回全局内存时，您应该使用组的 ID。

if (local_id == 0) output[local_size] = cache[0];

该行将重复写入输出[512]。您需要每个工作组写入输出中的专用位置。

kernel void reduce(global float* input, global float* output, const unsigned int N, local float* cache)
{
    const uint local_id = get_local_id(0);
    const uint global_id = get_global_id(0);
    const uint group_id = get_group_id(0);
    const uint local_size = get_local_size(0);

    cache[local_id] = (global_id < N) ? input[global_id] : 0.0f;
    barrier(CLK_LOCAL_MEM_FENCE);

    for (unsigned int s = local_size >> 1; s > 0; s >>= 1) {
        if (local_id < s) {
            cache[local_id] += cache[local_id + s];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if (local_id == 0) output[group_id] = cache[0];
}

然后您需要对主机输出的值求和。请注意，宿主代码中的 'b' 不需要保存 N 个元素。每个工作组将只使用一个元素。

//replace (globalSize/localSize) with the pre-calculated/known number of work groups
for (i=1; i<(globalSize/localSize); i++) {
    b[0] += b[i];
}

现在 b[0] 是您的总计。

在 OpenCL 中使用本地内存进行并行缩减

Parallel reduction using local memory in OpenCL

opencl