Numba cuda：使用共享内存添加数字会导致覆盖

【问题标题】：Numba cuda: Using shared memory to add numbers results in overwritingNumba cuda：使用共享内存添加数字会导致覆盖
【发布时间】：2019-11-29 12:55:44
【问题描述】：

我一直在尝试使用共享内存添加数字，因此如下所示：

线程0：共享内存变量sharedMemT[0]加1

线程1：给共享内存变量sharedMemT[0]加1

同步线程并将 sharedMemT[0] 存储到 output[0]

但结果是……1？？

@cuda.jit()
def add(output):
    sharedMemT = cuda.shared.array(shape=(1), dtype=int32)
    sharedMemT[0] = 0
    cuda.syncthreads()

    sharedMemT[0] += 1
    cuda.syncthreads()
    output[0] = sharedMemT[0]

out = np.array([0])
add[1, 2](out)
print(out) # results in [1]

【问题讨论】：

标签： python cuda shared-memory numba

【解决方案1】：

恭喜，你参加了一场记忆竞赛。线程 0 和 1 同时运行，所以结果是不确定的，无论是对共享内存变量的操作，还是写回全局内存。

为了使其正常工作，您需要使用原子内存操作序列化对共享内存变量的访问，然后只有一个线程写回全局内存：

$ cat atomic.py

import numpy as np
from numba import cuda, int32

@cuda.jit()
def add(output):
    sharedMemT = cuda.shared.array(shape=(1), dtype=int32)
    pos = cuda.grid(1)
    if pos == 0:
        sharedMemT[0] = 0

    cuda.syncthreads()

    cuda.atomic.add(sharedMemT, 0, 1)
    cuda.syncthreads()

    if pos == 0:
        output[0] = sharedMemT[0]

out = np.array([0])
add[1, 2](out)
print(out)

$ python atomic.py
[2]

【讨论】：