当按参数值传递给内核函数时，参数在哪里复制？

Question

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

当我通过值传递参数，如下所示：

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

由于可变一个和b被传递到核函数添加如在函数调用栈复制的值，我猜到一些存储器空间将需要复制。

如果我是正确的，那些参数在GPU或Host的主内存中复制 的额外内存空间？

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

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

Answer 1

很简短的答案是所有的CUDA内核参数都通过值传递，这些参数被主机通过API复制到GPU上的专用内存参数缓冲区中。目前，此缓冲区存储在常量内存中，每次启动内核时参数限制为4kb - 请参阅here。

---

在更多的细节，所述PTX标准（技术上因为计算能力2.0硬件和CUDA ABI出现）定义的专用逻辑状态空间呼叫.param持有内核和设备参数的参数。请参阅here。从该文档引用：

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

委员会还指出，：

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

所以参数状态空间的精确位置是实现特定的。在CUDA硬件的第一次迭代中，它实际上映射到用于内核参数的共享内存和用于设备函数参数的寄存器。但是，自从计算2.0硬件和PTX 2.2标准以来，它在大多数情况下映射到内核的常量内存。该文件说，following on the matter：

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

静态大小的常量变量有一个可选的变量初始值设定项;没有显式初始化器的常量变量默认为 ，初始化为零。 由主机初始化 驱动程序分配的常量缓冲区，指向此类缓冲区的指针为 作为参数传递给内核。

[着重矿]

因此，尽管内核参数被存储在常数存储器中，这是不一样的恒定存储器，其通过在CUDA C或定义一个变量作为__constant__映射到访问的.const状态空间在Fortran或Python中是等效的。相反，它是由驱动程序管理的内部设备内存池，不能被程序员直接访问。