将结构数组从主机复制到设备 cuda

Question

假设我有一个结构如下：

typedef struct values{
int one, int two, int three
} values;

现在，假设我在主机上创建一个值数组并填充随机数据

values vals*;
__device__ values* d_vals;
int main(){
     vals = (values*)malloc(sizeof(values) * A_LARGE_NUMBER);
     PopulateWithDate(); //populates vals with random data
}

现在我希望能够将这些值复制到设备中，这样我就可以像这样在内核中访问它们：

__global__ void myKernel(){
     printf("%d", d_vals[0].one);//I don't really want to print, but whenever I try to access I get an error
}

无论我尝试什么，都会遇到非法内存访问错误。

这是我目前的尝试：

int main(){
     vals = (values*)malloc(sizeof(values) * A_LARGE_NUMBER);
     PopulateWithDate(); //populates vals with random data

     values* d_ptr;
     cudaGetSymbolAddress((void**)&d_ptr, d_vals);
     cudaMalloc((void**)&d_ptr, A_LARGE_NUMBER * sizeof(values));

     cudaMemcpyToSymbol(d_ptr, &vals, sizeof(values) * A_LARGE_NUMBER);
     cudaDeviceSynchronize();
     dim3    blocksPerGrid(2, 2);
     dim3    threadsPerBlock(16, 16);

    myKernel<< <blocksPerGrid, threadsPerBlock >> >();
}

Answer 1

就您目前所展示的内容而言，使用__device__ 指针变量只会造成不必要的复杂性。只需使用cudaMalloc 进行设备存储的普通动态分配，否则遵循类似于任何 CUDA 示例代码（例如 vectorAdd）的方法。这是一个例子：

$ cat t1315.cu
#include <stdio.h>
#define A_LARGE_NUMBER 10

struct values{
int one, two, three;
};

values *vals;

__global__ void myKernel(values *d_vals){
     printf("%d\n", d_vals[0].one);
}

void PopulateWithData(){
  for (int i = 0; i < A_LARGE_NUMBER; i++){
    vals[i].one = 1;
    vals[i].two = 2;
    vals[i].three = 3;
  }
}


int main(){
     vals = (values*)malloc(sizeof(values) * A_LARGE_NUMBER);
     PopulateWithData(); //populates vals with random data

     values* d_ptr;
     cudaMalloc((void**)&d_ptr, A_LARGE_NUMBER * sizeof(values));
     cudaMemcpy(d_ptr, vals, A_LARGE_NUMBER *sizeof(values),cudaMemcpyHostToDevice);
     dim3    blocksPerGrid(1,1);
     dim3    threadsPerBlock(1, 1);

    myKernel<< <blocksPerGrid, threadsPerBlock >> >(d_ptr);
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_35 -o t1315 t1315.cu
$ cuda-memcheck ./t1315
========= CUDA-MEMCHECK
1
========= ERROR SUMMARY: 0 errors
$

您在所展示的内容中存在各种其他基本（非 CUDA）编码错误，我不会尝试全部解决。

如果你真的想保留你的__device__ 指针变量，并用它来指向设备数据（结构数组），那么你仍然需要使用cudaMalloc，整个过程需要额外的步骤。您可以按照答案here 中的示例进行操作。

按照该示例，对上述代码进行了一组更改，以使其使用 __device__ 指针变量而不是作为内核参数传递的指针：

$ cat t1315.cu
#include <stdio.h>
#define A_LARGE_NUMBER 10

struct values{
int one, two, three;
};

values *vals;
__device__ values *d_vals;

__global__ void myKernel(){
     printf("%d\n", d_vals[0].one);
}

void PopulateWithData(){
  for (int i = 0; i < A_LARGE_NUMBER; i++){
    vals[i].one = 1;
    vals[i].two = 2;
    vals[i].three = 3;
  }
}


int main(){
     vals = (values*)malloc(sizeof(values) * A_LARGE_NUMBER);
     PopulateWithData(); //populates vals with random data

     values* d_ptr;
     cudaMalloc((void**)&d_ptr, A_LARGE_NUMBER * sizeof(values));
     cudaMemcpy(d_ptr, vals, A_LARGE_NUMBER *sizeof(values),cudaMemcpyHostToDevice);
     cudaMemcpyToSymbol(d_vals, &d_ptr, sizeof(values*));
     dim3    blocksPerGrid(1,1);
     dim3    threadsPerBlock(1, 1);

    myKernel<< <blocksPerGrid, threadsPerBlock >> >();
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_35 -o t1315 t1315.cu
$ cuda-memcheck ./t1315
========= CUDA-MEMCHECK
1
========= ERROR SUMMARY: 0 errors
$