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

【问题标题】：How to share global memory within multiple kernels and multiple GPUs?如何在多个内核和多个 GPU 内共享全局内存？
【发布时间】：2013-03-22 02:28:28
【问题描述】：

----------------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。更重要的是，每个线程中的内核都会被启动多次。一个简单的伪代码如上。我有两个问题：

变量XX 是一个很长的字符串，在内核中是只读的。我想在mykernel 的多次启动期间保留它。仅在首次启动mykernel 时调用cudaMalloc 并将指针传递给mykernel 是否可以？或者我应该使用__device__ 限定符？
XX在四个线程中使用，所以我在文件a.c中将其声明为全局变量。多个cudaMalloc 和XX 是否正确，或者我应该使用variable *xx[4] 之类的数组？

【问题讨论】：

不要在自己的代码中使用cuStartThread()。

标签： cuda qualifiers

【解决方案1】：

对于在单个设备上运行的内核，您可以调用一次cudaMalloc 来创建保存字符串的变量XX，然后将cudaMalloc 创建的指针（即XX）传递给任何需要它的内核。
```
#define xx_length 20
char *XX;
cudaMalloc((void **)&XX, xx_length * sizeof(char));
...
kernel1<<<...>>>(XX, ...);
...
kernel2<<<...>>>(XX, ...);
etc.
```
为每个线程创建一个单独的 XX 变量，假设每个线程都用于访问不同的设备。你如何做到这一点将取决于 XX 的范围。但是数组：
```
char *XX[num_devices]; 
```

在全局范围内，应该没问题。

CUDA OpenMP sample 作为如何使用多线程管理多个 GPU 的示例可能很有趣。

【讨论】：