A100/H100上全局存储与bar.sync交互：存储可见性时机及延迟问询

给定以下PTX代码：

st.global.v4.b32 ...
bar.sync

已知bar.sync会阻塞执行，直到st.global.v4.b32全局存储操作对块内所有线程可见。现明确两个核心问题：

1. 全局存储操作完成可见性的节点是哪一步？
2. 全局存储后接bar.sync的延迟是多少？

全局内存存储通常采用"fire and forget"模式：存储指令写入存储缓冲区后，线程即可继续执行，存储缓冲区异步将数据提交至全局内存。针对A100和H100 GPU，结论如下：

可见性节点结论 #

bar.sync允许线程继续执行的节点是数据写入L2缓存时，原因如下：

• 写入存储缓冲区：仅代表线程发起了存储请求，数据仍为线程私有，无法被块内其他线程读取，屏障不会放行。
• 写入L1缓存：L1是SM（流式多处理器）私有缓存，同一块内线程可能跨warp或SM（部分场景），L1级别的一致性无法覆盖整个线程块，因此不满足可见性要求。
• 写入L2缓存：L2是GPU全局共享的缓存，一旦数据写入L2，GPU的内存一致性模型会保证同一块内所有线程都能读取到最新存储结果，此时屏障解除阻塞。
• 写入DRAM：属于后台异步操作，屏障无需等待这一步——L2的一致性已满足块内线程可见性，数据刷入DRAM由硬件自动完成。

延迟说明 #

st.global.v4.b32后接bar.sync的延迟，约等于全局存储请求到达L2缓存的时间。在A100和H100上，该延迟通常为几十到上百个时钟周期，具体数值会因内存访问模式、缓存命中率等因素波动，远小于等待DRAM写入的延迟（通常数百至上千时钟周期）。

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