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

Question

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

我的问题类似于。这个问题和link中的问题有两个很大的区别：

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

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

#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 中？

Answer 1

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

我们不使用cudaMemcpyAPI。我们使用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]....）

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

Synchronizing Statically Allocated Struct Instances between CPU and GPU

memory

struct

cuda

static-allocation