__syncthreads() 是否可以防止 read-after-write intra-warp 共享内存危害？答案

【问题标题】：Does anything short of __syncthreads() prevent read-after-write intra-warp shared memory hazards?__syncthreads() 是否可以防止 read-after-write intra-warp 共享内存危害？
【发布时间】：2017-09-17 01:54:16
【问题描述】：

我有一个warp，它将一些数据写入共享内存-没有覆盖，并且在从共享内存读取之后不久。虽然我的块中可能还有其他扭曲，但它们不会触及共享内存的任何部分或写入我感兴趣的扭曲读取的任何地方。

现在，我记得尽管 warp 以锁步方式执行，但我们不能保证共享内存写入之后的共享内存读取将返回应该由 warp 先前写入的相应值。（理论上这可能是由于指令重新排序或 - 正如@RobertCrovella 指出的那样 - 编译器优化了共享内存访问）

因此，我们需要借助一些显式同步。显然，块级__syncthreads() 工作。这就是does：

__syncthreads() 用于协调同一块的线程之间的通信。当块内的某些线程访问共享或全局内存中的相同地址时，对于其中一些内存访问，可能存在先读后写、先读后写或先写后写危险。通过在这些访问之间同步线程可以避免这些数据危害。

这对我的需求来说太强大了：

它也适用于全局内存，而不仅仅是共享内存。
它执行inter-warp同步；我只需要intra-warp。
它可以防止所有类型的危害R-after-W、W-after-R、W-after-W；我只需要R-after-W。
它也适用于多个线程执行写入共享内存中的同一位置的情况；在我的情况下，所有共享内存写入都是不相交的。

另一方面，类似__threadfence_block() does not seem to suffice。这两个强度级别之间有什么“中间”吗？

注意事项：

相关问题：CUDA __syncthreads() usage within a warp。
如果您要建议我使用改组，那么，是的，有时这是可能的 - 但如果您希望对数据进行数组访问，即动态决定您要访问的共享数据的哪个元素，则不行读。这可能会溢出到本地内存中，这对我来说似乎很可怕。
我在想volatile 可能对我有用，但我不确定使用它是否能达到我想要的效果。
如果您有一个假设计算机能力至少为 XX.YY 的答案，那就足够了。

【问题讨论】：

对于经线级别的使用，根据您第一段中的描述，volatile 可能是有意义的。没有一个清晰的例子是不可能确定的。根据您的描述，我看不出这与 classical parallel reduction tutorial（幻灯片 22）中的用法有何不同。根本问题不是指令重新排序，而是编译器将共享值优化到寄存器中
@RobertCrovella：这样的编译器优化不会确保我不需要__syncthreads() 开始吗？实际上 - 当它知道其他线程从编译时不确定的共享内存位置读取时，它怎么可能进行这种优化？
第一个问题没看懂。关于第二个问题，编译器对其他线程在做什么一无所知。 volatile 是您与编译器沟通的方式，即其他线程可能正在触及我正在查看的这个东西。这是对 GPU 编译器的一个常见误解——它以某种方式感知多线程，并且它的行为将根据可能在线程间发生的推断发生变化。它不是那样工作的。编程模型视图中只有一个线程，编译器也是。
@RobertCrovella：但编译器肯定在 warp 级别工作，即它决定（并因此知道）warp 中的所有线程将做什么。或者这仍然是一个误解？
我认为这是一种误解。

标签： cuda synchronization memory-barriers gpu-shared-memory

【解决方案1】：

如果我对@RobertCrovella 的理解正确，这段代码应该是安全的：

/* ... */
volatile MyType* ptr = get_some_shared_mem();
ptr[lane::index()] = foo();
auto other_lane_index = bar(); // returns a value within 0..31
auto other_lane_value = ptr[other_lane_index];
/* ... */

因为使用了volatile。（并且假设 bar() 不会造成自己的危险。）

【讨论】：