是否可以在块级别进行原子写入？
例如，考虑以下内容：
__global__ kernel (int atomic)
{
atomic+=blockid.x; //should be atomic for each block
}
【问题讨论】：
__global__？）。要么不要放那个标签，要么明确声明这是一个 CUDA 问题。
您可以在 CUDA 中执行一些原子操作。请参阅 CUDA 编程指南中的附录 B.11 原子函数。即：
__global__ void kernel (int *result)
{
atomicAdd(result, blockIdx.x); //
}
你也可以交换变量的值
__global__ void kernel (int *result)
{
atomicExch(result, blockIdx.x); //
}
这两个示例都在全局内存中运行。
在共享内存上运行的原子函数和在 64 位字上运行的原子函数仅适用于计算能力为 1.2 及以上的设备。
问候。
【讨论】：
您可以在共享内存上执行原子操作，但不能像您在代码 sn-p 中尝试这样做的方式：您的内核的 int 参数是线程特定的变量；即使所有线程都获得了您在启动时提供的相同值，它们也不会将其存储在共享内存中 - 以原子方式对其进行操作是没有意义的。
如果您已将 int * 传递给某个缓冲区 - 那将是全局内存中的缓冲区。您可以对全局内存中的数据执行设备范围的原子操作，如@pQB 的answer 中所述。但是您询问了块级原子操作……这对全局数据意义不大。尽管如此，如果您的一个线程写入某个全局地址，它可以全部__threadfence_block() 停止，直到该写入的效果对块中的所有其他线程可见。
CUDA 也支持正确的块级原子，但在共享内存上。在this Parallel4All blog entry 或相关的section CUDA Programming Guide 中了解如何使用共享内存。
如果你有一些__shared__ int x，你确实可以对其执行块级原子操作，语法与全局原子操作相同：atomicAdd(&x, 7) 将自动将 123 添加到 x 的值。但是 - 请记住，块中的所有线程都将执行相同的操作，并且您当然不想一次尝试多达 1024 次原子写入。通常你会有类似的东西
__shared__ some_buffer[NumFoosPerBar];
// ...
if (check_condition()) {
int foo_index = get_thread_foo_index_for(threadIdx.x);
atomicAdd(&some_buffer[foo_index], 7);
}
可能有多个线程写入同一位置，但不一定。当您确实期望多次写入时 - 不要使用原子，而是对要写入的值执行某种缩减。
【讨论】：
虽然不清楚您对块/块级别的含义，但听起来您只需要一个原子添加。
这些可以在#include <asm/atomic.h> 的内核中找到，您的代码将是
__global__ kernel (int atomic)
{
atomic_add(blockid.x,&atomic);
}
atomic 必须是 atomic_t 类型并且 blockid.x 是 int。
【讨论】：