關於CUDA __constant__內存和多GPU？

Question

在共享內存編程模型中，每個線程都可以看到任何全局變量。在simliar方式就像在共享內存系統全局變量，這讓我有點宣佈

在CUDA，不斷記憶擔心：

考慮下面的代碼：

__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[]的值？

順便說一句：我搜索的網站，有一些相關的問題，但我找不到任何明確的答案。

Answer 1

然後，數組的內容[]本地爲每個CUDA上下文/設備，這樣我們就可以安全地更新每個設備的「私人」 []數組，而不必擔心改變數組的值[]分配其他cuda設備？

是的，代碼

__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 
$