OpenCL 中的本地和全局内存栅栏有什么区别？

Question

我正在尝试使用 PyOpenCL 来减少总和，类似于示例：https://dournac.org/info/gpu_sum_reduction。我正在尝试对所有值为 1 的向量求和。结果在第一个元素中应该是 16384。然而，似乎只收集了一些要点。是否需要本地索引？是否有任何竞争条件（当我运行两次结果不一样时）？以下代码有什么问题？

import numpy as np
import pyopencl as cl

def readKernel(kernelFile):
    with open(kernelFile, 'r') as f:
        data=f.read()
    return data

a_np = np.random.rand(128*128).astype(np.float32)
a_np=a_np.reshape((128,128))
print(a_np.shape)

device = cl.get_platforms()[0].get_devices(cl.device_type.GPU)[0]
print(device)
ctx=cl.Context(devices=[device])
#ctx = cl.create_some_context() #ask which context to use 
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags

a_g = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=a_np)

prg = cl.Program(ctx,readKernel("kernel2.cl")).build()

prg.test(queue, a_np.shape, None, a_g)

cl.enqueue_copy(queue, a_np, a_g).wait()
np.savetxt("teste2.txt",a_np,fmt="%i")

内核是：

__kernel void test(__global float *count){
    int id = get_global_id(0)+get_global_id(1)*get_global_size(0);
    int nelements = get_global_size(0)*get_global_size(1);

    count[id] = 1;
    barrier(CLK_GLOBAL_MEM_FENCE); 

    for (int stride = nelements/2; stride>0; stride = stride/2){
        barrier(CLK_GLOBAL_MEM_FENCE); //wait everyone update
        if (id < stride){
            int s1 = count[id];
            int s2 = count[id+stride];
            count[id] = s1+s2;
        }
    }
    barrier(CLK_GLOBAL_MEM_FENCE); //wait everyone update
}

Answer 1

问题是您的内核实现为在一个工作组内进行缩减，并且隐含地安排了许多工作组。

根据 GPU，每个工作组的最大工作项数量不同。 Nvidia 为 1024，AMD 和 Intel 为 256（旧 GPU 中的 Intel 为 512）。

让我们假设在这个例子中，GPU 上每个工作组的最大工作项数是 256。在这种情况下，最大 2d 工作组大小可以是 16x16，因此如果您使用该大小的矩阵，您的内核将 return 正确的结果。使用原始大小 128x128 并且在调度内核时未指定局部大小，实现会为您计算出全局大小 128x128 和局部大小（很可能）16x16，这意味着正在调度 8 个工作组。在当前内核中，每个工作组都从不同的 id 开始计算，但索引减少到 0，因此您有竞争条件，因此每个运行.

的结果不同

您有 2 个选项可以解决此问题：

重写您的内核以计算一个工作组中的所有内容，并使用全局、局部大小安排它：(16x16)、(16,16) 或每个工作组设备的最大工作项目数
使用全局、局部大小：(128x128),(16x16)，每个工作组将计算其结果，然后在 cpu 端必须对每个工作组求和以获得最终结果。

对于 128x128，第一个选项将是首选，因为它应该执行得更快并且应该更容易实现。

OpenCL 中的本地和全局内存栅栏有什么区别？

Whats the difference between local and global memory fence in OpenCL?

python-2.7

pyopencl