CUDA写入其他warp看不到的全局内存

Question

我试图向CUDA的新手解释全球记忆.我想出了以下虚拟内核,它阻止其他warp中的其他线程,直到选定的warp将全局变量设置为另一个值:

__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()在主机代码中放置它等待所有线程).这让我想知道其他warp是否没有收到is_done变量的变化.但是,据我所知,全局内存意味着可以在设备级别(即至少是网格中的所有块)中看到该值.

我的问题如下:CUDA是否有任何缓存/优化使得这种不一致的全局内存视图可以发生？有没有办法从驻留在全局内存中的变量访问"最新"值？

Answer 1

CUDA是否进行了任何缓存/优化,使得这种不一致的全局内存视图可以发生？有没有办法从驻留在全局内存中的变量访问"最新"值？

是的,有缓存行为.您可以使用volatile 限定符对其进行修改.

这是一个有效的例子:

$ cat t310.cu
#include <stdio.h>

#ifndef USE_VOLATILE
__device__ __managed__ int is_done = 0;
#else
__device__ volatile __managed__ int is_done = 0;
#endif

__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");
    }
}



int main(){

  with_sync<<<1,1024>>>();
  cudaDeviceSynchronize();
}
$ nvcc -o t310 t310.cu
$ ./t310
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
^C
$ nvcc -o t310 t310.cu -DUSE_VOLATILE
$ ./t310
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
I'm done!
$

(如果不清楚上面的第一次运行被Ctrl-C终止,由于挂起)

特斯拉P100 PCIE CUDA 10.0,CentOS 7