更新: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
)。