OpenCL中的本地记忆围栏和全局内存围栏有什么区别?

Question

我试图用pyopencl进行减简和示例类似：https：//dournac.org/info/gpu/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，每个工作组有不同数量的最大工作项。对于1024年的NVIDIA，AMD和Intel 256（Intel在较旧的GPU中具有512）。

对于此示例，让我们假设您的GPU上的每个工作组的最大工作项为256。在这种情况下，最大2D Worgroup尺寸可以为16x16，因此，如果您使用矩阵的大小，则内核将返回正确的结果。使用原始尺寸的128x128，并且在安排内核时未指定本地大小，该实现会计算您和您将获得全局大小的128x128和本地尺寸（很可能）16x16，这意味着正在安排8个worgroups。在当前内核中，每个工作组正在从不同的id开始计算，但是索引缩小到0直到0，因此您的比赛条件有不同的结果。

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

1. 重写您的内核以计算一个工作组中的所有内容，并使用全局，本地大小进行安排：（16x16），（16,16）或每个工作组设备的最大工作项目
2. 使用全局，本地大小：（128x128），（16x16），每个工作组将计算其结果，然后在CPU方面必须总结每个工作组才能获得最终结果。

对于128x128，第一个选项将是首选的，因为它应该更快地执行，并且应该更简单地实现。