我对 CUDA 编程很陌生，关于内存模型的一些内容对我来说还不是很清楚。比如，它是如何工作的？例如，如果我有一个简单的内核
__global__ void kernel(const int* a, int* b){
some computation where different threads in different blocks might
write at the same index of b
}
所以我想a 会在所谓的常量内存中。但是b 呢？既然不同block中的不同线程都会往里面写，那它是怎么工作的呢？我在某处读到，可以保证在同一块中的不同线程同时写入全局内存的情况下，至少会写入一个，但不能保证其他的。我是否需要担心这一点，例如，让块中的每个线程都写入共享内存，一旦它们都完成，让一个将它们全部写入全局内存？还是 CUDA 替我处理？
【问题讨论】：
所以我想
a会在所谓的常量内存中。
是的，a 指针 将在常量内存中，但不是因为它被标记为const（这是完全正交的）。 b 指针也在常量内存中。 所有内核参数在常量内存中传递（CC 1.x 除外）。 a 和 b 指向的内存理论上可以是任何东西（设备全局内存、主机固定内存、UVA 可寻址的任何东西，我相信）。它所在的位置由您（用户）选择。
我在某处读到，保证在同一块中的不同线程并发写入全局内存的情况下，至少会写入一个，但不能保证其他线程。
假设您的代码如下所示：
b[0] = 10; // Executed by all threads
那么是的，这是一个（良性）竞争条件，因为所有线程都将相同的值写入相同的位置。写入的结果已定义，但写入次数未指定，执行“最终”写入的线程也是如此。唯一的保证是至少发生一次写入。在实践中，我相信每个 warp 都会发出一次写入，如果您的块包含多个 warp（它们应该如此），这会浪费带宽。
另一方面，如果您的代码如下所示：
b[0] = threadIdx.x;
这是明显的未定义行为。
我是否需要担心这一点，例如让一个块中的每个线程都写入共享内存，一旦它们都完成，让一个将它们全部写入全局内存？
是的，这就是通常的做法。
【讨论】：
__synctreads(); if(threadId.x == 0) for(i=...) b[i] = s[i]，假设我有一个共享数组 s ，每个线程都写它必须写的东西。那么共享内存呢？写入是原子的吗？还是我也应该自己承担？感谢您的回答，如果我的问题是基本的，我很抱歉，我是 CUDA 的新手。