在 CUDA __constant__ 内存和多 GPU 上？答案

【问题标题】：On CUDA __constant__ memory and multi-GPUs?在 CUDA __constant__ 内存和多 GPU 上？
【发布时间】：2014-02-20 12:51:58
【问题描述】：

在共享内存编程模型中，任何全局变量对每个线程都是可见的。

在 CUDA 中，constant 内存的声明方式与共享内存系统中的全局变量类似，这让我有点担心：

考虑以下代码：

__constant__ int array[1024];

void hostFunction(int DeviceID, cudaStream_t streamIdx)
{
    cudaSetDevice(DeviceID);
    someKernel<<<100,1024,0, streamIdx>>>(...);
    //The function someKernel will use data stored in array[] on current device;
};

那么，array[] 的内容是否是每个 cuda 上下文/设备的本地内容，这样我们就可以安全地更新每个设备的“私有”array[]，而不必担心更改分配在其他 cuda 设备上的 array[] 的值？

顺便说一句：我搜索了该网站，有一些相关的问题，但是我无法从其中找到任何明确的答案。

【问题讨论】：

标签： c++ c arrays cuda

【解决方案1】：

那么，array[] 的内容是否对每个 cuda 上下文/设备都是本地的，这样我们就可以安全地更新每个设备的“私有”array[]，而不必担心更改分配在其他 cuda 设备上的 array[] 的值？

是的，单行代码

__constant__ int array[1024];

在您的程序访问的每个设备上创建一个分配。

然后您可以使用例如：

在每个设备上单独加载__constant__ 内存：

cudaSetDevice(0);
cudaMemcpyToSymbol(array, my_device_0_constant_data, 1024*sizeof(int));

并为您希望使用的每个设备重复上述操作。

关于__device__ variables，也有类似的说法。

这是一个完整的例子：

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

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__constant__ int my_const_data;

__device__ int my_dev_data;

__global__ void my_kernel(int my_dev){

  printf("device %d constant data is: %d\n", my_dev, my_const_data);
  printf("device %d __device__ data is: %d\n", my_dev, my_dev_data);
}

int main(){

  int num_dev = 0;
  cudaGetDeviceCount(&num_dev);
  cudaCheckErrors("get device count fail");
  if (num_dev == 0) {printf("no cuda devices found!\n"); return 1;}
  for (int i = 0; i < num_dev; i++){
    int cdata = i;
    int ddata = 10*i;
    cudaSetDevice(i);
    cudaMemcpyToSymbol(my_const_data, &cdata, sizeof(int));
    cudaMemcpyToSymbol(my_dev_data, &ddata, sizeof(int));
    cudaCheckErrors("memcpy to symbol fail");}
  for (int i = 0; i < num_dev; i++){
    cudaSetDevice(i);
    my_kernel<<<1,1>>>(i);
    cudaDeviceSynchronize();}
  cudaCheckErrors("kernel fail");
  return 0;
}

$ nvcc -arch=sm_20 -o t223 t223.cu
$ ./t223
device 0 constant data is: 0
device 0 __device__ data is: 0
device 1 constant data is: 1
device 1 __device__ data is: 10
device 2 constant data is: 2
device 2 __device__ data is: 20
device 3 constant data is: 3
device 3 __device__ data is: 30
$

【讨论】：