如何将 C++ 结构数组传递给 CUDA 设备？答案

【问题标题】：How can I pass a C++ array of structs to a CUDA device?如何将 C++ 结构数组传递给 CUDA 设备？
【发布时间】：2016-03-24 17:27:22
【问题描述】：

我花了 2 天时间试图弄清楚这一点，但一无所获。假设我有一个如下所示的结构：

struct Thing {
    bool is_solid;
    double matrix[9];
}

我想创建一个名为 things 的结构数组，然后在 GPU 上处理该数组。比如：

Thing *things;
int num_of_things = 100;
cudaMallocManaged((void **)&things, num_of_things * sizeof(Thing));

// Something missing here? Malloc individual structs? Everything I try doesn't work.

things[10].is_solid = true; // Segfaults

最好的做法是这样做而不是传递具有num_of_things 大数组的单个结构？在我看来，这可能会变得非常讨厌，尤其是当您已经拥有数组时（例如 matrix，它需要是 9 * num_of_things。

任何信息将不胜感激！

【问题讨论】：

您的代码对我来说很好用。 Here is my full test case（你应该提供这样的东西）。你的方法没有错；这将是您在托管场景中的典型做法。很可能您犯了以下几个错误之一： 1. 您没有为 cc3.0+ 设备进行编译。 2.你没有在cc3.0+设备上运行。 3.你的环境does not support managed memory usage.
无论如何，最好的建议是在遇到 cuda 代码问题时使用proper cuda error checking，我已经在我的测试用例中证明了这一点。尝试在我的测试用例中运行代码，看看你得到什么样的输出；我认为这对于正在发生的事情具有指导意义。
非常感谢！事实证明我做错了，因为编译器说没有足够的参数，所以我将第三个参数0 传递给cudaMallocManaged。我不认为它是相关的，所以我把它排除在外。它必须是 cudaMemAttachGlobal 或 cudaMemAttachHost，正如一位非常善良的人在 freenode/#cuda 上指出的那样。
您使用的是 CUDA 7.5 吗？尽管文档在这方面尚不清楚，但编译器将/应该接受仅带有 2 个参数的 cudaMallocManaged 调用（正如您在这个问题中所指出的那样，正如我在示例中所指出的那样）；在这种情况下，flags 变量将默认为cudaMemAttachGlobal。如果您说“编译器说的参数不够”，我只能假设这是 Visual Studio Intellisense？（实际上不是编译器）无论如何，您可以通过检查 cuda 包含目录中的 cuda_runtime.h 来说服自己相信 flags 参数的默认参数
谢谢，我检查了一下，它确实默认为cuda_runtime.h 中的那个，但仍然无法仅使用 2 个 args 进行编译。不知道为什么。我的设置是纯命令行、g++、CMake 的 FindCuda 和带有标志 -gencode arch=compute_30,code=sm_30 的 nvcc。 IRC 的那个家伙提到这可能是 C 的东西，但我认为不是这样，因为文件是 .cu 而其他的都是 c++。

标签： c++ pointers image-processing cuda parallel-processing

【解决方案1】：

在 cmets 中进行一些对话后，似乎 OP 发布的代码没有问题。我能够成功编译并运行围绕该代码构建的测试用例，OP 也是如此：

$ cat t1005.cu
#include <iostream>

struct Thing {
    bool is_solid;
    double matrix[9];
};

int main(){

  Thing *things;
  int num_of_things = 100;
  cudaError_t ret = cudaMallocManaged((void **)&things, num_of_things * sizeof(Thing));
  if (ret != cudaSuccess) {
    std::cout << cudaGetErrorString(ret) << std::endl;
    return 1;}
  else {
    things[10].is_solid = true;
    std::cout << "Success!" << std::endl;
    return 0;}
}
$ nvcc -arch=sm_30 -o t1005 t1005.cu
$ ./t1005
Success!
$

关于这个问题：

最好的做法是这样做而不是传递一个包含 num_of_things 大数组的单个结构？

是的，这是一种明智的做法，无论是否使用托管内存都可以使用。可以使用单个 cudaMemcpy 调用以简单的方式将不包含指向其他地方动态分配的数据的嵌入式指针的或多或少任何结构的数组传输到 GPU（例如，如果未使用托管内存。）

解决关于cudaMallocManaged的第三个（flags）参数的问题：

如果指定，则传递零是不正确的（尽管 OP 发布的代码没有提供任何证据。）您应该使用 the documented choices 之一。
如果没有指定，这仍然有效，并且提供了一个默认参数cudaMemAttachGlobal。这可以通过查看cuda_runtime.h 文件或简单地编译/运行上面的测试代码来确认。这一点似乎是文档中的一个疏忽，我已经在 NVIDIA 提交了一个内部问题来查看它。因此，未来文档可能会在这方面发生变化。

最后，proper cuda error checking 在您遇到 CUDA 代码问题时始终可用，使用此类代码可能会发现所犯的任何错误。 OP 在代码 cmets 中报告的 seg 错误几乎可以肯定是由于 cudaMallocManaged 调用失败（可能是因为错误地提供了零参数），因此有问题的指针 (things) 没有实际分配。该指针的后续使用将导致段错误。我的测试代码演示了如何避免该段错误，即使 cudaMallocManaged 调用由于某种原因失败，关键是正确的错误检查。

【讨论】：