时间戳

首页 > 解决方案 > Numba cuda：为什么一维数组的总和不正确？

python - Numba cuda：为什么一维数组的总和不正确？

问题描述

我正在练习 numba 和 cuda 编程。我试图用 cuda 总结一系列的。总和不正确。我认为必须在最后正确同步和收集数据。

 @cuda.jit
def my_kernel(const_array, res_array):

    sbuf = cuda.shared.array(512, float32)

    # Thread id in a 1D block
    tx = cuda.threadIdx.x
    # Block id in a 1D grid
    ty = cuda.blockIdx.x
    # Block width, i.e. number of threads per block
    bw = cuda.blockDim.x
    # Compute flattened index inside the array
    pos = tx + ty * bw

    sbuf[tx] = 0

    if pos < const_array.shape[0]:

        sbuf[tx] = const_array[pos] # do the computation

    cuda.syncthreads()
    if cuda.threadIdx.x == 0:
        for i in range(bw):
            res_array[0] += sbuf[i] 


    return


data_size = 10000000
res = numpy.zeros(1, dtype=numpy.float64)
const_array = numpy.ones(data_size, dtype=numpy.int8)

threadsperblock = 512
blockspergrid = math.ceil(data_size / threadsperblock)

my_kernel[blockspergrid, threadsperblock](const_array, res)

print(res)

每次我运行这段代码时，它都会检索不同的值，例如 28160.0，但当然它必须是 10m。

并提示？

标签： pythoncudanumba

解决方案

问题似乎是您没有对整组块进行求和。您的向量维度为 10000000 和 512 个线程，这意味着您需要对所有块求和 19532 个块。这是在标准 CUDA 语言中通过启动多个内核（主要用于旧设备）或使用原子操作来实现的。具体来说，您的问题出在代码的这一部分：

if pos < const_array.shape[0]:
    sbuf[tx] = const_array[pos] # do the computation    cuda.syncthreads()
if cuda.threadIdx.x == 0:
    for i in range(bw):
        res_array[0] += sbuf[i]

在前两行中，您将数据从全局复制到数组 sbuf 的共享内存中。但是，不同块中的所有线程同时尝试将其本地数据添加到 res_array 处的全局内存地址中，这不是顺序的，不同的线程可能只是读取相同的数据两次并给您错误的结果。解决方法是先在共享内存中进行部分求和，然后再进行原子求和，避免异步读写操作

if cuda.threadIdx.x == 0:
    sum = 0
    for i in range(bw):
        sum  += sbuf[i] 
    cuda.atomic.add(res_array, 0, sum)

那应该可以解决您的问题。

问候。

推荐阅读