CUDA中的内核参数传递？

Question

我对CUDA内核的工作原理有一个新手疑问。

如果我有以下代码(使用来自这里的cuPrintf函数):

#include "cuPrintf.cu"

__global__ void testKernel(int param){
    cuPrintf("Param value: %d\n", param);
}

int main(void){

    // initialize cuPrintf
    cudaPrintfInit();

    int a = 456;    

    testKernel<<<4,1>>>(a);

    // display the device's greeting
    cudaPrintfDisplay();

    // clean up after cuPrintf
    cudaPrintfEnd();
}

执行的输出是：

Param value: 456
Param value: 456
Param value: 456
Param value: 456

我无法理解内核如何读取我传递的参数的正确值，它不是分配在主机内存中吗？GPU能够从主机内存中读取吗？

谢谢，

安德烈

- Andrea

---

函数变量由NVCC编译器隐式地进行了驱动程序内部传输和复制。 - Buddhika

Answer 1

根据CUDA C编程指南中的E.2.5.2章节，__global__函数参数通过以下方式传递到设备：

• 在计算能力为1.x的设备上，通过共享内存传递，并且限制为256字节。
• 在计算能力为2.x及更高版本的设备上，通过常量内存传递，并且限制为4 KB。

- phil

Answer 2

声明 void testKernel(int param) 表示传递的是值而不是引用。换句话说，栈中包含了 a 的值的副本，而不是指向 a 的指针。CUDA将栈复制到在GPU上运行的内核中。

- Jesse Hall

---

所以，如果我理解正确，在内核执行期间，“param”将位于GPU内存堆栈中。 因此，如果我多次读取它，我将不会访问任何“主机内存”，从而降低性能，对吗？ - Andrea

文档并没有明确说明内核参数存储在哪里。但是可以假设它们存储在快速内存中：寄存器、共享或常量中。正如kirbuchi在他的回答中所说，CUDA编程指南（3.0指南中的B.13节）指出：“执行配置的参数在实际函数参数之前进行评估，并且像函数参数一样，目前通过共享内存传递到设备”。 - Jesse Hall

太好了。再次感谢您的所有答案。 - Andrea

Answer 3

根据CUDA编程指南（附录B.16），参数通过共享内存传递到设备。执行配置的参数在实际函数参数之前进行评估，并且像函数参数一样，目前通过共享内存传递到设备。

- kirbuchi

---

该引用是指执行配置参数（4和1），而不是函数参数（a == 456），这正是OP所询问的。 - Steve Fallows

@Steve 你说得对，但它将它们与函数参数进行比较，它说函数参数也是通过共享内存传递的。应该强调那一部分 :) - kirbuchi

传递内核参数的方法因架构而异。计算能力为1.*的设备将值放入共享内存中。计算能力大于等于2.0的设备将值放入常量内存中。 - Greg Smith

这并不是很合理，因为您可以修改每个参数，比如说ar是一个int参数，因此您可以编写： at = threadIdx.x; 这并不合理，因为您修改了块中所有线程的唯一共享值。 - TripleS

Answer 4

当您调用内核时，参数将传递给内核；否则，您如何与GPU通信呢？这与在着色器中设置统一变量的想法相同。

- Mikola

Answer 5

在运行时API中，全局函数的参数会被隐式地从主机传输到设备并进行复制。

NVCC编译器生成的代码会将参数传输隐藏起来。您可以在CUDA编程指南中找到参数的大小和限制。

- Buddhika