(Py)OpenCL 同时修改所有线程的值答案

【问题标题】：(Py)OpenCL modify value from all threads simultanously(Py)OpenCL 同时修改所有线程的值
【发布时间】：2020-11-16 18:03:09
【问题描述】：

很简单，我有以下内核，它修改了C[0] 的值，其中C 是一个只有一个元素的数组。

__kernel void sigma(__global float *A, __global float *B, __global float *C) {
    int i = get_global_id(0);
    printf("Adding %.2f + %.2f", A[i], B[i]);
    C[0] += A[i] + B[i];
}

问题是，最后C[0] 具有最后完成的线程的值，特别是在这个例子中我得到以下内容

Adding 1.00 + 0.00
Adding 2.00 + 1.00
Adding 3.00 + 1.00
Adding 4.00 + 1.00
[5.]

最后C[0] 是4.00 + 1.00。我想要的是C[0] 是(1.00 + 0.00) + (2.00 + 1.00) + (3.00 + 1.00) + (4.00 + 1.00)。所以我希望将每个线程的A[i] 和B[i] 添加到C[0]。

另外我不只是在寻找补充，我希望它与任何功能或操作兼容。

这可能是多余的，但在主机代码中，我只是在做最低限度的工作以将数据传递给内核。这个问题是否与宿主代码有关？

import pyopencl as cl, numpy as np; 
ctx = cl.create_some_context(); queue = cl.CommandQueue(ctx); mf = cl.mem_flags
M = np.array([1, 2, 3, 4]).astype(np.float32) # A
V = np.array([0, 1, 1, 1]).astype(np.float32) # B
a = np.array([0]).astype(np.float32) # C
# Transfer data to GPU
A_GPU = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=M) 
B_GPU = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=V)
C_GPU = cl.Buffer(ctx, mf.WRITE_ONLY, a.nbytes)
c = np.zeros(shape=a.shape, dtype= np.float32) # array to copy the result
kernel.sigma(queue, M.shape, None, A_GPU, B_GPU, C_GPU)
cl.enqueue_copy(queue, c, C_GPU).wait()

【问题讨论】：

标签： python opencl pyopencl

【解决方案1】：

内核的编写方式效率低下，并且在许多情况下会导致不确定的结果。通常，您应该避免让多个线程（工作项）同时写入同一个地址。在您的情况下，C[0] += A[i] + B[i]; 是一个 read modify write 序列，从并行访问的角度来看，这会使事情变得更糟。

您需要使用的是并行缩减模式。此模式允许您有效且确定地将数组中的所有元素求和为单个值（或在您的情况下为 2 个数组）。它还支持除加法之外的许多其他关联运算符。

网上有很多资源可以根据自己的情况进行调整。

【讨论】：

【解决方案2】：

正如@Elad Maimoni 所说，在 PyOpenCl 中执行缩减的一种方法是使用版本 2.0 中提供的函数 work_group_reduce_add。

在python中可以这样实现

内核

kernel = cl.Program(ctx, """
__kernel void resum(__global float *A, __global float *B, __global float *a) {
    int i = get_global_id(0);
    a[0] = work_group_reduce_add(A[i] + B[i]);
}
""").build(options='-cl-std=CL2.0') # Build using cl 2.0

主机代码的其余部分

# Some example arrays
a = np.array([1, 2, 3, 4, 5, 6, 7, 8, 9])
b = np.array([9, 8, 7, 6, 5, 4, 3, 2, 1])
c = np.array([0])
d = np.array([0])

# Create GPU data
a = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a)
b = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b)
c = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=c)

kernel.resum(queue, (10, ), None, a, b, c)
cl.enqueue_copy(queue, d, c)
print(d)

这将输出[90]。

【讨论】：