CUDA cudaMemcpy 结构体数组

Question

我希望清理我的项目中CUDA内核的参数。

---

现在，一个内核需要 3 个 uint32_t 数组，这导致代码非常丑陋：（id 表示全局线程 id，valX 是任意值）

__global__ void some_kernel(uint32_t * arr1, uint32_t * arr2, uint32_t * arr3){arr1[id] = val1; arr2[id] = val2; arr3[id] = val3;}

我希望用结构体来围绕所有这些数组:

typedef struct S{uint_32_t arr1, uint_32_t arr2, uint_32_t arr3, uint32_t size} S;

其中 size 表示结构体内每个 arrX 的长度。

我想要的是这样的效果：

__global__ void some_kernel(S * s){s->arr1[id] = val1; s->arr2[id] = val2; s->arr3[id] = val3;}

---

如果要针对像这样的结构体进行相应的cudaMalloc和cudaMemcpy操作，该怎么做呢？在我还未意识到的情况下，是否存在任何性能上的缺陷呢？

提前谢谢！

- Daniel Jünger

---

是的，我得到了这个：test.cu(27)：错误：不存在从“S”到“const void *”的适当转换函数...就在memcopy处。 - Daniel Jünger

为什么不直接按值传递结构体？不需要使用cudaMalloc或cudaMemcpy。 - talonmies

创建一个包含指针的结构体，使用cudaMalloc为每个指针分配内存。通过值传递方式传递该结构体。如果你不理解这一点，那么我认为你需要重新学习C++中指针、引用和值的相关知识。CUDA在概念上有些复杂，但在尝试编写CUDA代码之前，你需要彻底理解C或C++。你的第一个C++程序也不应该是你的第一个CUDA程序。 - talonmies

错误显示它无法将 S 转换为 const void *，你需要传递指向 memcopy 的指针而不是实际变量。所以 S *src, dst; src = (S *)malloc(n * sizeof(S)); cudaMalloc((void **)&dst, n*sizeof(S)); cudaMemcpy(dst, src, n * sizeof(S), cudaMemcpyHostToDevice); - triple_r

我不建议在构造函数和析构函数中进行设备内存分配和释放，除非你真的非常小心地管理作用域。否则，这可能会导致一些非常难以诊断的运行时错误。 - talonmies

Answer 1

你至少有两个选择。一个绝佳的选择已经由talonmies提供了（链接），但我会向你介绍“学习艰辛之路”的方法。

首先，你需要定义你的结构体：

typedef struct S {
    uint32_t *arr1;
    uint32_t *arr2;
    uint32_t *arr3; 
    uint32_t size;
} S;

...和内核定义（带有一些全局变量，但您不需要遵循该模式）：

const int size = 10000;

__global__ void some_kernel(S *s)
{
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    if (id < size)
    {
        s->arr1[id] = 1; // val1
        s->arr2[id] = 2; // val2
        s->arr3[id] = 3; // val3
    }
}

请注意，if 语句可以防止数组越界。

接下来，我们将介绍一些函数来准备数据、执行内核并打印一些结果。第一部分是数据分配：

uint32_t *host_arr1, *host_arr2, *host_arr3;
uint32_t *dev_arr1, *dev_arr2, *dev_arr3;

// Allocate and fill host data
host_arr1 = new uint32_t[size]();
host_arr2 = new uint32_t[size]();
host_arr3 = new uint32_t[size]();

// Allocate device data   
cudaMalloc((void **) &dev_arr1, size * sizeof(*dev_arr1));
cudaMalloc((void **) &dev_arr2, size * sizeof(*dev_arr2));
cudaMalloc((void **) &dev_arr3, size * sizeof(*dev_arr3));

// Allocate helper struct on the device
S *dev_s;
cudaMalloc((void **) &dev_s, sizeof(*dev_s));

这没什么特别的，你只需要分配三个数组和一个结构体。更有趣的是如何处理将这样的数据复制到设备中：

// Copy data from host to device
cudaMemcpy(dev_arr1, host_arr1, size * sizeof(*dev_arr1), cudaMemcpyHostToDevice);
cudaMemcpy(dev_arr2, host_arr2, size * sizeof(*dev_arr2), cudaMemcpyHostToDevice);
cudaMemcpy(dev_arr3, host_arr3, size * sizeof(*dev_arr3), cudaMemcpyHostToDevice);

// NOTE: Binding pointers with dev_s
cudaMemcpy(&(dev_s->arr1), &dev_arr1, sizeof(dev_s->arr1), cudaMemcpyHostToDevice);
cudaMemcpy(&(dev_s->arr2), &dev_arr2, sizeof(dev_s->arr2), cudaMemcpyHostToDevice);
cudaMemcpy(&(dev_s->arr3), &dev_arr3, sizeof(dev_s->arr3), cudaMemcpyHostToDevice);

除了普通的数组拷贝之外，你注意到还必须将它们与结构体“绑定”起来。为此，您需要传递指针地址。因此，只有这些指针被复制。

接下来的内核调用，再次将数据复制回主机并打印结果：

// Call kernel
some_kernel<<<10000/256 + 1, 256>>>(dev_s); // block size need to be a multiply of 256

// Copy result to host:
cudaMemcpy(host_arr1, dev_arr1, size * sizeof(*host_arr1), cudaMemcpyDeviceToHost);
cudaMemcpy(host_arr2, dev_arr2, size * sizeof(*host_arr2), cudaMemcpyDeviceToHost);
cudaMemcpy(host_arr3, dev_arr3, size * sizeof(*host_arr3), cudaMemcpyDeviceToHost);

// Print some result
std::cout << host_arr1[size-1] << std::endl;
std::cout << host_arr2[size-1] << std::endl;
std::cout << host_arr3[size-1] << std::endl;

请记住，在任何严肃的代码中，您应该始终检查CUDA API调用是否存在错误。

- Grzegorz Szpetkowski

---

如果您首先在主机内存中构建设备结构，然后将其复制到dev_s中，您可以使用单个memcpy替换“//注意：将指针与dev_s绑定”部分中的三个内存复制。这将更简单和更快速。 - talonmies