CUDA中的全局内存写入是否被视为原子操作？

Question

在CUDA中，全局内存写操作是否被视为原子操作？

考虑以下CUDA内核代码：

int idx = blockIdx.x*blockDim.x+threadIdx.x;
int gidx = idx%1000;
globalStorage[gidx] = somefunction(idx);

全局内存写入globalStorage是原子的吗？例如，没有竞态条件，使得并发的内核线程写入存储在globalStorage中相同变量的字节，这可能会搞乱结果（例如部分写入）？

请注意，我这里不是说原子操作，如加/减/按位等操作，而是直接全局写入。

编辑：重写示例代码以避免混淆。

- user0002128

---

请认真思考您的问题 - 写出一个完全有效的答案，然后发现问题稍后被重写而使得答案无效是不太好玩的。 - talonmies

你似乎不理解什么是原子性。 原子性是指执行多个操作（经典的：读取-修改-写入），这些操作按顺序执行，不能被其他干扰操作打断。因此，仅谈论单个操作（例如写入）并问它是否具有原子性是没有意义的。通常情况下，在全局内存中对同一位置进行多次写入或在全局内存中读取和写入同一位置的组合将肯定导致竞争条件和意外结果的可能性。 - Robert Crovella

如果您询问多个写入的顺序，CUDA中没有执行顺序的保证（即使在单个线程或warp基础上），因此对同一位置的多个写入将产生未定义的结果。其中一个写入将成功，即将在一段时间后到达该位置。但是除此之外的任何事情都是未定义的。我认为，您的问题仍然不清楚。 - Robert Crovella

Answer 1

CUDA中的内存访问并不是隐式原子操作。然而，如果在运行内核时idx对于每个线程都有唯一的值，那么您最初展示的代码本质上不会出现内存竞争问题。

因此，您最初的代码：

int idx = blockIdx.x*blockDim.x+threadIdx.x;
globalStorage[idx] = somefunction(idx);

如果内核启动使用1D网格且globalStorage大小合适，则会很安全，而您的第二个版本：

int idx = blockIdx.x*blockDim.x+threadIdx.x;
int gidx = idx%1000;
globalStorage[gidx] = somefunction(idx);

不会是因为多个线程可能会写入globalStorage中的相同条目。在这种情况下，没有原子保护或序列化机制可以产生可预测的结果。

- talonmies