当将参数按值传递到内核函数时，参数被复制到哪里？

Question

我是一个CUDA编程的初学者，有一个问题。

当我按值传递参数时，就像这样：

__global__ void add(int a, int b, int *c) {
    // some operations
}

自变量a和b作为副本值在函数调用堆栈中传递给内核函数add，所以我猜测需要复制一些内存空间。

如果我是正确的，那么这些参数被复制到GPU还是主机的主内存中？

我关心这个问题的原因是我应该将一个大的结构体传递给内核函数。

我也考虑过传递结构体的指针，但这种方式似乎需要为结构体和每个成员变量调用cudamalloc。

- Hongmin Yang

---

将传递给内核调用的按值参数放置在__constant__内存中，这是一种特殊类型的设备内存。通常，在内核代码的开头，需要的参数将从__constant__内存复制到寄存器中。有关此内容，请参阅编程指南此处。 - Robert Crovella

Answer 1

非常简短的回答是，所有传递给CUDA内核的参数都是按值传递的，并且这些参数通过API由主机复制到GPU上的专用内存参数缓冲区中。目前，此缓冲区存储在常量内存中，每个内核启动的参数限制为4kb - 请参见此处。

---

更详细地说，PTX标准（从计算能力2.0硬件和CUDA ABI出现开始）定义了一个专用的逻辑状态空间调用.param，其中包含内核和设备参数参数。请参见这里。引用该文档中的内容：

每个内核函数定义都包括一个可选的参数列表。这些参数是在.param状态空间中声明的可寻址只读变量。通过使用ld.param指令，可以通过这些参数变量访问从主机传递到内核的值。内核参数变量在网格内的所有CTA之间共享。

它进一步指出：

注意：参数空间的位置是实现特定的。例如，在某些实现中，内核参数驻留在全局内存中。在这种情况下，参数空间和全局空间之间不提供访问保护。类似地，函数参数根据应用程序二进制接口（ABI）的函数调用约定映射到参数传递寄存器和/或堆栈位置。

因此，参数状态空间的精确位置取决于具体的实现。在CUDA硬件的第一次迭代中，它实际上映射到内核参数的共享内存和设备函数参数的寄存器。然而，自从compute 2.0硬件和PTX 2.2标准以来，在大多数情况下，它映射到内核的常量内存。文档中提到了以下相关内容：

常量（.const）状态空间是由主机初始化的只读内存。常量内存可以使用ld.const指令访问。常量内存大小受限，目前限制为64 KB，可用于保存静态大小的常量变量。还有额外的640 KB常量内存，组织为十个独立的64 KB区域。驱动程序可以在这些区域中分配和初始化常量缓冲区，并将指向缓冲区的指针作为内核函数参数传递。由于这十个区域不是连续的，因此驱动程序必须确保常量缓冲区被分配，以使每个缓冲区完全适合于一个64 KB区域，并且不跨越区域边界。

静态大小的常量变量具有可选的变量初始化程序；没有显式初始化程序的常量变量默认为零。由驱动程序分配的常量缓冲区由主机初始化，并将指向这样的缓冲区的指针作为参数传递给内核。

因此，虽然内核参数存储在常量内存中，但这并不是与在CUDA C中将变量定义为__constant__或Fortran或Python中的等效操作所访问的.const状态空间相同的常量内存。相反，它是由驱动程序管理的设备内存的内部池，程序员无法直接访问它。

- talonmies