将结构作为参数传递给 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 中作为@ 987654336@，这意味着for 循环按照我的预期执行。作为一个附带问题，使用size_t 这是必不可少的unsigned long long int 不好的做法，因为GPU 是由32 位内核组成的？

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

Answer 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++) 引用将任何内容传递给内核；这都是“按值” - 对象本身或指向它的指针。