__syncthreads() 是否可以防止写后读内部扭曲共享内存的危害？

Question

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

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

因此，我们需要求助于一些显式同步。显然，块级 __syncthreads() 有效。这就是 does:

__syncthreads() is used to coordinate communication between the threads of the same block. When some threads within a block access the same addresses in shared or global memory, there are potential read-after-write, write-after-read, or write-after-write hazards for some of these memory accesses. These data hazards can be avoided by synchronizing threads in-between these accesses.

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

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

另一方面，__threadfence_block() does not seem to suffice。有什么"in-between"那两个等级的实力吗？

备注：

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

Answer 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() 不会弄乱自己带来的危险。）

__syncthreads() 是否可以防止写后读内部扭曲共享内存的危害？

Does anything short of __syncthreads() prevent read-after-write intra-warp shared memory hazards?

synchronization

cuda

memory-barriers

gpu-shared-memory