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

Question

我是 CUDA 编程的新手，所以我想阐明将结构传递到内核时的行为。我定义了以下 struct 以在某种程度上模仿知道其自身大小的 3D 数组的行为：

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

我创建了两个protoarray类型的变量，在主机端和设备端通过malloc和cudaMalloc动态分配space给数据，并更新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

这是部分答案，因为如果没有 proper program to look into，则 difficult/impossible 可以猜测为什么您会在 arr_device.dim1.

中看到无效值

The struct is passed by copy, so all its contents are copied into shared memory of each block.

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

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

Generally, how safe is it to pass struct and class into kernels as arguments

我个人对此事的经验法则是：如果 struct/class...

是平凡可复制的；和
struct/class 的所有成员都是为主机端和设备端定义的，或者至少 - 设计时考虑到了 GPU 的使用；

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

passing struct and class into kernels as arguments [ - ] is [it] bad practice that should be avoided at all cost?

否。但请记住，大多数 C++ 库仅提供主机端代码；并且不是为了在 GPU 上使用而编写的。所以我会小心使用非平凡的类而不经过大量审查。

I imagine that passing pointers to classes to kernels is difficult in case they contain members which point to dynamically allocated memory

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

and that they should be very lightweight if I want to pass [objects to kernels] by value.

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

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

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

Behaviour of passing struct as a parameter to a CUDA kernel

struct

cuda

class