5

更新：while()编译器优化了下面的条件，因此两个线程都只是跳过条件并进入 CS，即使带有-O0标志。有谁知道为什么编译器会这样做？顺便说一句，声明全局变量volatile会导致程序由于某种奇怪的原因挂起......

我阅读了CUDA 编程指南，但我仍然不清楚 CUDA 如何处理与全局内存相关的内存一致性。（这与内存层次结构不同）基本上，我正在运行测试试图打破顺序一致性。我正在使用的算法是彼得森的算法，用于内核函数内部两个线程之间的互斥：

flag[threadIdx.x] = 1; // both these are global
turn = 1-threadIdx.x;

while(flag[1-threadIdx.x] == 1 && turn == (1- threadIdx.x));
shared_gloabl_variable_x ++;

flag[threadIdx.x] = 0;

这是相当简单的。每个线程通过将其标志设置为一个来请求临界区，并通过轮到另一个线程来表现良好。在评估 时while()，如果其他线程没有设置其标志，则请求线程可以安全地进入临界区。现在这种方法的一个微妙问题是，如果编译器重新排序写入，以便在写入到turn之前执行写入flag。如果发生这种情况，两个线程将同时在 CS 中结束。这很容易用普通的 Pthreads 证明，因为大多数处理器不实现顺序一致性。但是 GPU呢？

这两个线程将在同一个经线中。他们将以锁步模式执行他们的语句。但是当他们到达turn变量时，他们正在写入同一个变量，因此内部扭曲执行变得序列化（不管顺序是什么）。现在在这一点上，获胜的线程是继续进入 while 条件，还是等待另一个线程完成其写入，以便两者可以同时评估while()？路径将再次在 处分叉while()，因为只有其中一个会赢，而另一个等待。

运行代码后，我让它不断地破坏 SC。我读取的值是 ALWAYS 1，这意味着两个线程每次都以某种方式进入 CS。这怎么可能（GPU 按顺序执行指令）？（注意：我用 编译它-O0，所以没有编译器优化，因此没有使用volatile）。

4

2 回答 2

3

编辑：由于您只有两个线程并且1-threadIdx.x可以工作，因此您必须使用线程 ID 0 和 1。线程 0 和 1 将始终是所有当前 NVIDIA GPU 上同一扭曲的一部分。Warp 以 SIMD 方式执行指令，具有针对不同条件的线程执行掩码。您的 while 循环是一个发散的条件。

  • 当turn和flags 不是 volatile时，编译器可能会重新排序指令，您会看到进入 CS 的两个线程的行为
  • Whenturn和flags are volatile，您会看到挂起。原因是其中一个线程将成功写入转弯，因此turn将是 0 或 1。假设turn==0：如果硬件选择执行线程 0 的发散分支部分，那么一切正常。但是如果它选择执行线程 1 的发散分支的一部分，那么它将在 while 循环上旋转，并且线程 0 永远不会轮到它，因此会挂起。

您可以通过确保您的两个线程处于不同的经线中来避免挂起，但我认为经线必须同时驻留在 SM 上，以便可以从两者发出指令并取得进展。（可能适用于不同 SM 上的并发扭曲，因为这是全局内存；但这可能需要 __threadfence() 而不仅仅是 __threadfence_block()。）

一般来说，这是一个很好的例子，说明为什么这样的代码在 GPU 上是不安全的，不应该使用。我意识到这只是一个调查实验。一般来说，CUDA GPU 没有——正如你提到的大多数处理器没有——实现顺序一致性。

原始答案

  1. 变量turn和flag必须是volatile，否则flag不会重复加载 ，条件turn == 1-threadIdx.X不会被重新评估，而是将被视为true。
  2. store to和 store to__threadfence_block()之间应该有一个正确的顺序。flagturn
  3. 在共享变量增量之前应该有一个__threadfence_block()（也应该声明volatile）。您可能还需要__syncthreads()或至少__threadfence_block()在增量之后确保它对其他线程可见。

我有一种预感，即使在进行了这些修复之后，您仍然可能会遇到麻烦。让我们知道怎么回事。

顺便说一句，你在这一行有一个语法错误，所以很明显这不是你真正的代码：

while(flag[1-threadIdx.x] == 1 and turn==[1- threadIdx.x]);
于 2012-04-23T04:26:44.393 回答
2

在没有额外的内存屏障（例如 __threadfence()）的情况下，全局内存的顺序一致性仅在给定线程内强制执行。

于 2012-04-22T17:58:47.743 回答