将结构作为参数传递给 CUDA 内核的行为

Posted 2023-02-18

【中文标题】将结构作为参数传递给 CUDA 内核的行为

【英文标题】：Behaviour of passing struct as a parameter to a CUDA kernel

【发布时间】：2021-06-10 05:10:15

【问题描述】：

我对 CUDA 编程比较陌生，所以我想澄清当我将结构传递到内核时的行为。我定义了以下struct 来在某种程度上模仿知道自己大小的 3D 数组的行为：

struct protoarray 
    size_t dim1;
    size_t dim2;
    size_t dim3;
    float* data;
;

我创建了两个protoarray类型的变量，在主机端和设备端通过malloc和cudaMalloc为数据动态分配空间，并更新dim1、dim2和dim3以反映我想要的数组大小这个结构来表示。我在this thread 中读到struct 应该通过副本传递。所以这就是我在内核中所做的事情

__global__ void kernel(curandState_t *state, protoarray arr_device)
    const size_t dim1 = arr_device.dim1;
    const size_t dim2 = arr_device.dim2;
    
    for(size_t j(0); j < dim2; j++)
        for(size_t i(0); i < dim1; i++)
            // Do something 
        
    

结构体是通过拷贝传递的，所以它的所有内容都被拷贝到每个块的共享内存中。这是我得到奇怪行为的地方，我希望你能帮助我。假设我在主机端设置了arr_device.dim1 = 2。在内核内部调试并在for 循环之一处设置断点时，检查arr_device.dim1 的值会产生类似16776576 的东西，没有大到足以导致溢出的地方，但这个值正确复制到dim1 中作为@ 987654336@，这意味着for 循环按照我的预期执行。作为一个附带问题，使用size_t 这是必不可少的unsigned long long int 不好的做法，因为GPU 是由32 位内核组成的？

一般来说，将struct 和class 作为参数传递到内核中有多安全，是否应该不惜一切代价避免这种不好的做法？我想将指向类的指针传递给内核是很困难的，因为它们包含指向动态分配内存的成员，并且如果我想通过值传递它们，它们应该非常轻量级。

【问题讨论】：

您确定size_t 在 CPU 和 GPU 上的大小相同吗？ （我相信，没有构造函数的 structs 通常是安全的。）

@KenY-N：CUDA 保证 sizof(size_t) 和 sizeof(void *) 在 GPU 和设备上是相同的无论使用什么平台。不同平台可能不一样

这让我很担心——“结构是通过副本传递的，所以它的所有内容都被复制到每个块的共享内存中”。这根本不是事情的运作方式。无论问题是什么，它都不是由您在此问题中显示的任何内容引起的。 minimal reproducible example，请

@KenY-N 全部更改为unsigned int，现在值匹配。您能否请我参考一些关于struct 与内核中的构造函数的问题的来源？ @talonmies ““结构是通过副本传递的，因此它的所有内容都被复制到每个块的共享内存中”这基本上是来自我链接的 NVIDIA 开发者论坛线程的引用。

@SlavaK.: 好的，在 2009 年确实如此。自 2010 年以来就不是这样了。所有内核参数都放在所有 CUDA 硬件上的专用常量内存库中，除了第一代 G80/ G90 零件

【参考方案1】：

这是部分答案，因为没有proper program to look into，很难/不可能猜出为什么您会在arr_device.dim1 中看到无效值。

结构体是通过拷贝传递的，所以它的所有内容都会拷贝到每个块的共享内存中。

不正确。内核参数存储在常量内存中，它是设备全局的，而不是特定于块的。它们不存储共享内存（特定于块）。

当线程运行时，它通常将参数从常量内存读取到寄存器中（同样，不是共享内存）。

一般来说，将 struct 和 class 作为参数传递给内核有多安全

我个人对此问题的经验法则是：如果结构/类...

可轻松复制；和 它的结构/类的所有成员都是为主机端和设备端定义的，或者至少 - 在设计时考虑了 GPU 的使用；

那么传递给内核应该是安全的。

将结构和类作为参数传递到内核 [ - ] 是 [它] 应该不惜一切代价避免的不好的做法？

否。但请记住，大多数 C++ 库只提供主机端代码；并没有考虑在 GPU 上使用。所以我会谨慎使用未经大量审查的非平凡类。

我认为将指向类的指针传递给内核是很困难的，因为它们包含指向动态分配内存的成员

是的，这可能会有问题。但是 - 如果您使用 cuda::memory::managed::allocate()、cuda::memory::managed::make_unique() 或 cudaMallocManaged() - 那么这应该“正常工作”，即相关内存页面将在访问时根据需要被提取到 GPU 或 CPU。见：

Unified Memory in CUDA for beginners Beyond GPU Memory Limits with Unified Memory on Pascal

如果我想按值传递[对象到内核]，它们应该非常轻量级。

是的，因为每个线程必须先从常量内存中读取每个参数，然后才能使用该参数。虽然恒定内存允许这种情况相对较快地发生，但您仍然希望最大限度地减少大量开销。

还请记住，您不能通过 (C++) 引用将任何内容传递给内核；这都是“按值” - 对象本身或指向它的指针。

【讨论】：

历史旁注：很早的 CUDA（在它公开之前）只允许一个内核参数。所以我写了很多内核，使用struct 作为所有内核参数的集合。该代码的很大一部分保留了多年，并且在我记得的任何系统上都没有任何问题。

@konovification：如果这回答了您的问题，请将其标记为已接受（按 V 标记）。如果您喜欢它（无论是否接受） - 支持它；这就是所有必要的感谢。