1

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

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

所以，我们需要求助于一些显式的同步。显然，块级__syncthreads()工作。这是做什么的：

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

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

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

另一方面，类似的东西__threadfence_block() 似乎还不够。这两个强度级别之间有什么“中间”吗？

笔记：

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

1 回答 1

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()不会造成自己的危险。）

于 2017-04-20T13:32:44.857 回答