全局内存读写的撕裂保证是什么？Pascal架构对齐32位整数读写疑问

核心结论 #

• 非自然对齐的32位整数：会出现访问撕裂
• 自然对齐的32位整数：在Pascal及现代GPU架构中，单个线程的加载/存储操作无撕裂，但多线程并发读写的结果仍未定义

详细说明 #

1. 撕裂的本质 #

内存访问撕裂指单个读写操作被拆分为多个底层内存事务，导致其他线程可能读取到半更新的数据，或写操作产生不完整的覆盖结果。GPU全局内存事务以缓存行（通常32/64/128字节）为单位调度，但对齐的特定大小数据访问会被合并为原子事务。

2. 自然对齐的32位整数 #

当32位整数地址是4的整数倍（自然对齐）时：

• 在Pascal及后续架构（Kepler及以上均支持）中，单个线程的ld.global.u32（加载32位无符号整数）和st.global.u32（存储32位无符号整数）操作会被硬件保证为原子完成——整个32位值的读写在单个内存事务中执行，不会被拆分。因此其他线程不会读到半更新的中间值，写操作也不会出现部分覆盖的情况。
• 注意：这仅保证单个操作无撕裂，不保证多线程并发读写的结果确定性。比如10个线程同时写同一个值，最终结果是任意一个线程的写入值，无法预测；读写混合时，读线程只会拿到写前或写后的完整值，不会拿到混合状态。

3. 非自然对齐的32位整数 #

未自然对齐的32位整数会跨越两个缓存行，单个线程的读写操作会被拆分为两个独立的内存事务。在事务执行间隙，其他线程的读写可能修改其中一个缓存行的数据，导致当前线程读取到新旧混合的不完整值，或写入仅完成一部分，必然出现撕裂。

4. Pascal架构的特殊说明 #

Pascal架构从Kepler继承了对齐32位内存操作的原子性保障，完全支持无撕裂的单个对齐32位数据读写。若需要多线程并发读写的结果可预测，必须使用原子操作（如atomicAdd、atomicExch）来保证操作互斥。

内容的提问来源于stack exchange，提问作者Box Box Box Box