无分支发散时，仍需bar.warp.sync保证内存顺序吗？

是的，即使Warp内所有线程无分支发散、同步执行存储指令，依然需要bar.warp.sync来保证共享内存通信的内存顺序和可见性，原因如下：

1. 指令同步发射≠内存操作同步完成 #

Warp内的SIMT架构保证所有线程同步发射指令，但这并不意味着存储操作会同时完成并对其他线程可见。硬件层面，共享内存的存储操作可能因流水线延迟、bank冲突处理等微架构细节，出现细微的完成时序差异。没有栅栏同步的话，后续的加载操作可能在其他线程的存储操作尚未提交到共享内存阵列时就执行，导致读取到旧值。

2. PTX内存模型的明确要求 #

根据PTX文档的说明：

bar.warp.sync可保证参与barrier的线程间的内存顺序。因此warp内线程若要通过内存通信，可先存储到内存，执行bar.warp.sync，再安全读取其他线程存储的值。

PTX的内存一致性模型中，Warp内线程的内存操作没有隐式的同步保证——哪怕指令无分支发散。bar.warp.sync的核心作用是提供Warp级内存栅栏：

• 确保所有在barrier前的存储操作，对所有参与barrier的线程全局可见；
• 阻止barrier后的加载操作被硬件重排到barrier之前执行。

3. 实际代码场景的验证 #

假设我们有如下无分支的Warp内共享内存通信代码：

// 所有线程同步执行存储指令
st.shared.f32 [shared_mem+tid*4], val;
// 若无bar.warp.sync，此处加载可能读取到旧值
ld.shared.f32 r0, [shared_mem+(tid+1)%32*4];

即使存储指令是同步发射的，没有bar.warp.sync的情况下，加载操作可能在其他线程的存储操作未完成时执行，导致读取数据不正确。只有插入bar.warp.sync后，才能确保所有存储操作完成并可见：

st.shared.f32 [shared_mem+tid*4], val;
bar.warp.sync 0x7fffffff; // 同步所有Warp内线程
ld.shared.f32 r0, [shared_mem+(tid+1)%32*4];

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