3

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

当我按值传递参数时，如下所示：

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

由于变量a和b被传递给内核函数add作为函数调用堆栈中的复制值，我猜需要一些内存空间来复制。

如果我是对的，那么这些参数是否复制到 GPU 或主机的主内存中的额外内存空间？

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

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

4

1 回答 1

2

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

更详细地说，PTX 标准（从技术上讲，自计算能力 2.0 硬件和 CUDA ABI 出现以来）定义了一个专用的逻辑状态空间调用.param，它保存内核和设备参数参数。见这里。引用该文档：

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

它进一步指出：

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

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

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

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

[强调我的]

因此，虽然内核参数存储在常量内存中，但这与映射到.const状态空间的常量内存不同，后者可以通过定义__constant__CUDA C 中的变量或 Fortran 或 Python 中的等效项来访问状态空间。相反，它是由驱动程序管理的内部设备内存池，程序员不能直接访问。

于 2017-11-27T15:44:21.000 回答