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

Question

我对 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

中，如

 2

，这意味着

for

循环按照我的预期执行。作为一个附带问题，使用

size_t

这是必不可少的

unsigned long long int

不好的做法，因为 GPU 是由 32 位核心组成的？

一般来说，将

struct

和

class

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

Answer 1

这是部分答案，因为如果没有适当的程序来研究，很难/不可能猜测为什么您会在

arr_device.dim1

中看到无效值。

结构体是通过复制传递的，因此它的所有内容都被复制到每个块的共享内存中。

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

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

一般来说，将结构和类作为参数传递到内核中有多安全

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

• 是可简单复制构造的；和
• 其结构/类的所有成员都是为主机端和设备端定义的，或者至少是为 GPU 使用而设计的；

那么它应该可以安全地传递给内核。

_{（现在它更像是一条经验法则，我实际上在 CUDA 的 现代 C++ 包装器 中，在用于启动内核的 最里面的代码中强制执行该规则的第一部分。）}

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

不。但请记住，大多数 C++ 库仅提供主机端代码；并且并不是为了在 GPU 上使用而编写的。因此，在没有经过大量审查的情况下，我会对使用不平凡的类持谨慎态度。

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

是的，这可能会有问题。但是 - 如果您使用 cuda::memory::managed::allocate()

、

cuda::memory::managed::make_unique()

 或 

cudaMallocManaged()

 - 那么这应该“正常工作”，即在访问时，相关内存页面将根据需要提取到 GPU 或 CPU。参见：

• 适合初学者的 CUDA 统一内存
• Pascal 上的统一内存超越 GPU 内存限制

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

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

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