【发布时间】：2019-04-02 13:39:25
【问题描述】：
我试图向刚接触 CUDA 的人解释全局内存。我想出了以下虚拟内核，它阻止其他线程中的其他线程，直到选定的线程将全局变量设置为另一个值：
__global__ void with_sync()
{
while (threadIdx.x / 32 != 0)
{
if (is_done != 0)
{
break;
}
}
if (threadIdx.x / 32 == 0)
{
is_done = 1;
printf("I'm done!\n");
}
}
变量is_done在函数外部声明为__device__ __managed__ int（如果我错了，请纠正我，这意味着变量将驻留在全局内存空间中。
但是，当我像这样执行这个内核（单个块中的 1024 个 1D 线程）时：
with_sync<<<1, 1024>>>();
cudaDeviceSynchronize();
I'm done 按预期打印出来。但是，CUDA 程序并没有终止（我在主机代码中放置了cudaDeviceSynchronize() 以便它等待所有线程）。这让我想知道其他经线是否没有收到 is_done 变量的变化。但是，我知道全局内存意味着可以在设备级别看到该值（即至少，网格中的所有块）。
我的问题如下：CUDA 是否进行了任何缓存/优化，从而导致这种不一致的全局内存视图可能发生？有没有办法从驻留在全局内存中的变量中访问“最新”值？
【问题讨论】：
-
@RobertCrovella 这没有回答为什么它打印出“我完成了”，这意味着它正在运行的任何经纱都执行了前面的指令，即将
is_done变量设置为 1 . -
@RobertCrovella 有趣。我在 Volta (Tesla V100) 和 Pascal (Tesla P4) GPU 上运行它，它们都以相同的方式运行（打印出 32 个“I'm done”实例并停止）。
-
是的，volta 问题与您之前的代码示例有关，其中的行为可能会有所不同。 volta 问题与您当前的代码示例无关。