在 CPU 和 GPU 之间同步静态分配的结构实例答案

【问题标题】：Synchronizing Statically Allocated Struct Instances between CPU and GPU在 CPU 和 GPU 之间同步静态分配的结构实例
【发布时间】：2021-03-04 20:03:58
【问题描述】：

我有一个包含数组的结构，我想将内容从 CPU 内存中该结构的一个实例复制到 GPU 内存中的另一个实例。

我的问题类似于this one。这个问题与链接中的问题有两个很大的区别：

我没有使用结构数组。我只需要一个。
结构的所有实例都是静态分配的。

为了回答我自己的问题，我尝试修改答案中的代码如下：

#include <stdio.h>
#include <stdlib.h>

#define cudaCheckError() { \
    cudaError_t err = cudaGetLastError(); \
    if(err != cudaSuccess) { \
      printf("Cuda error: %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
      exit(1); \
    } \
  }

struct Test {
    char array[5];
};

__global__ void kernel(Test *dev_test) {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test[0].array[i]);
    }
}

__device__ Test dev_test; //dev_test is now global, statically allocated, and one instance of the struct

int main(void) {

    int size = 5;
    Test test; //test is now statically allocated and one instance of the struct

    char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
    memcpy(test.array, temp, size * sizeof(char));

    cudaCheckError();
    cudaMemcpy(&dev_test, &test, sizeof(Test), cudaMemcpyHostToDevice);
    cudaCheckError();
    kernel<<<1, 1>>>(&dev_test);
    cudaCheckError();
    cudaDeviceSynchronize();
    cudaCheckError();

    //  memory free
    return 0;
}

但是这段代码会抛出运行时错误：

Cuda error: HelloCUDA.cu:34: invalid argument

有没有办法将test 复制到dev_test 中？

【问题讨论】：

标签： memory struct cuda static-allocation

【解决方案1】：

使用静态分配的__device__ 变量时：

我们不使用cudaMemcpy API。我们使用cudaMemcpyToSymbol（或cudaMemcpyFromSymbol）API
我们不将__device__ 变量作为内核参数传递。它们在全球范围内。您只需在内核代码中使用它们即可。

以下代码解决了这些问题：

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

#define cudaCheckError() { \
    cudaError_t err = cudaGetLastError(); \
    if(err != cudaSuccess) { \
      printf("Cuda error: %s:%d: %s\n", __FILE__, __LINE__, cudaGetErrorString(err)); \
      exit(1); \
    } \
  }

struct Test {
    char array[5];
};

__device__ Test dev_test; //dev_test is now global, statically allocated, and one instance of the struct

__global__ void kernel() {
    for(int i=0; i < 5; i++) {
        printf("Kernel[0][i]: %c \n", dev_test.array[i]);
    }
}


int main(void) {

    int size = 5;
    Test test; //test is now statically allocated and one instance of the struct

    char temp[] = { 'a', 'b', 'c', 'd' , 'e' };
    memcpy(test.array, temp, size * sizeof(char));

    cudaCheckError();
    cudaMemcpyToSymbol(dev_test, &test, sizeof(Test));
    cudaCheckError();
    kernel<<<1, 1>>>();
    cudaCheckError();
    cudaDeviceSynchronize();
    cudaCheckError();

    //  memory free
    return 0;
}
$ nvcc -o t10 t10.cu
$ cuda-memcheck ./t10
========= CUDA-MEMCHECK
Kernel[0][i]: a
Kernel[0][i]: b
Kernel[0][i]: c
Kernel[0][i]: d
Kernel[0][i]: e
========= ERROR SUMMARY: 0 errors
$

（您在内核代码中的数组使用也没有意义。dev_test 不是数组，因此您无法对其进行索引：dev_test[0]....）

【讨论】：