内存一致性模型CUDA 4.0和全局内存？

Question

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

我阅读了CUDA编程指南,但我仍然不清楚CUDA如何处理与全局内存相关的内存一致性.(这与内存层次结构不同)基本上,我正在运行试图破坏顺序一致性的测试.我使用的算法是Peterson的内核函数内两个线程之间互斥的算法:

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;

这非常简单.每个线程通过将其标志设置为1并通过将转向设置为另一个线程来获得关键部分来请求关键部分.在评估时while(),如果其他线程没有设置其标志,则请求线程可以安全地进入临界区.现在,这种方法的一个微妙问题是,如果编译器重新排序写入,那么写入将turn在写入之前执行flag.如果发生这种情况,两个线程将同时在CS中结束.这很容易用普通的Pthreads来证明,因为大多数处理器都没有实现顺序一致性.但GPU呢？

这两个线程都将处于相同的warp中.并且他们将以锁步模式执行他们的语句.但是当它们到达turn变量时,它们正在写入相同的变量,因此内部执行变为序列化(无论顺序是什么).现在,此时,获胜的线程是否继续进入while条件,还是等待另一个线程完成其写入,以便两者可以同时进行评估while()？路径将再次分叉while(),因为只有其中一个会赢,而另一个等待.

运行代码后,我得到它一直打破SC.我读取的值总是1,这意味着两个线程每次都以某种方式进入CS.这怎么可能(GPU按顺序执行指令)？(注意:我已经编译了它-O0,因此没有编译器优化,因此没有使用volatile).

Answer 1

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