如何在多个内核和多个GPU中共享全局内存？

Question

----------------a.c--------------------- 
variable *XX; 
func1(){ 
    for(...){ 
    for(i = 0; i < 4; i++) 
     cutStartThread(func2,args) 
    } 
} 
---------------b.cu------------------- 
func2(args){ 
    cudaSetDevice(i); 
    xx = cudaMalloc(); 
    mykernel<<<...>>>(xx); 
} 
-------------------------------------- 

最近，我想为我的程序使用多个GPU设备。我的节点上有四个Tesla C2075卡。我使用四个线程来管理四个GPU。更重要的是，每个线程中的内核都会多次启动。一个简单的伪代码如上。我有两个问题：

1. 可变XX是一个很长的字符串，只有在内核中读取。我想在mykernel的多次发布期间保留它。只有在mykernel第一次启动时才可以致电cudaMalloc并将指针传递给mykernel？或者我应该使用__device__限定符？

2. XX在四个线程中使用，所以我将其声明为文件a.c中的全局变量。 XX的多个cudaMalloc是否正确，还是应该使用数组，例如variable *xx[4]？

Answer 1

1. 对于由一个单一的设备上运行的内核使用情况，您可以拨打cudaMalloc一次以创建您的变量XX拿着字符串，然后路过cudaMalloc（即XX）创建的指针为准内核需要它。

   #define xx_length 20 
   char *XX; 
   cudaMalloc((void **)&XX, xx_length * sizeof(char)); 
   ... 
   kernel1<<<...>>>(XX, ...); 
   ... 
   kernel2<<<...>>>(XX, ...); 
   etc. 
   

2. 创建针对每个线程独立的XX变量，假设每个线程正被用于访问一个不同的设备。你如何做到这一点将取决于XX的范围。但是数组：在全局范围内

   char *XX[num_devices]; 
   

，应该没问题。

CUDA OpenMP sample作为如何使用多个线程来管理多个GPU的例子可能是有意义的。