CUDA共享内存中的load/store操作是否具备原子性？

首先纠正一个核心认知偏差：共享内存的bank冲突串行化规则，不等价于默认给普通load/store提供原子性保障。

• 手册里提到的同bank访问串行化，是硬件解决bank冲突的调度逻辑：多个线程访问同一bank的不同地址时，硬件会拆分请求多轮顺序执行。但这个串行化是「访存请求批次」粒度的，从来没有承诺过「同地址的并发普通读写是原子的」，更不承诺跨线程的访存可见性、顺序性。

PTX指令的语义差异 #

ld.weak.shared.cta 和 ld.relaxed.shared.cta 效果完全不一致，原子标记不是仅给编译器做优化提示用的：

• ld.weak 是普通变量访存的默认生成指令：它仅保证单线程视角下的访存逻辑符合程序顺序，不承诺跨线程的读写完整性、可见性。编译器可以对这类指令做任意优化：比如把共享内存变量缓存到寄存器跳过实际访存、合并多次写、重排访存顺序；硬件层面也不保证同地址并发读写不会出现半字撕裂的情况。
• ld.relaxed 是原子操作对应的访存指令：首先硬件层面对齐基元类型的单次读写是原子的，不会出现读写撕裂；其次会强制禁止编译器对该访存做消除、非法重排，保证每次操作都真实访问对应内存；同时保证同scope下其他线程能看到完整的读写结果，只是不承诺跨操作的内存先后顺序。

共享内存普通int与block域atomic的差异 #

在变量自然对齐的前提下，仅从硬件单次读写的原子性来看，两者在部分架构上可能碰巧表现一致，但从编程模型语义上完全不等价，不能互换：

• 普通__shared__ int的访存不受原子规则约束，编译器可能做出各种导致跨线程访问异常的优化：比如循环中反复读同一个共享int时，编译器会直接把值放到寄存器，不再重新读取共享内存的最新值；也可能把多个写操作合并、重排顺序，最终其他线程读到的值完全不符合预期。
• __shared__ cuda::atomic<int, cuda::thread_scope_block>的relaxed序load/store，会严格遵守原子语义：每次读写都真实访问共享内存，禁止非法重排，硬件保证读写原子性，跨线程可见性符合对应内存序的承诺。

全局内存普通int与block域atomic的差异 #

两者完全不等价，不存在可互换性：

• 全局内存没有共享内存的bank串行调度机制，普通全局内存访存和原子访存走的是完全不同的硬件路径。普通访存不会经过全局内存的原子仲裁逻辑，哪怕是对齐的int类型，并发普通读写也可能出现读写撕裂。
• 编译器对普通全局内存变量的优化权限和共享内存普通变量一致，会做缓存、重排、消除访存等优化，根本不保证跨线程访问的正确性。

实操结论：只要是跨线程并发访问的共享变量，不管存在共享内存还是全局内存，不管访问范围是同warp、同block还是全设备，只要存在并发读写，就必须使用cuda::atomic对应类型，不要依赖对硬件隐式行为的假设。旧代码里用volatile+手动加内存栅栏的写法是CUDA未提供标准原子类型时期的妥协方案，极易出错，不建议在新代码中使用。

内容的提问来源于stack exchange，提问作者Pierre T.