如何将多个重复的参数传递给 CUDA 内核

Posted 2023-03-23

【中文标题】如何将多个重复的参数传递给 CUDA 内核

【英文标题】：How to pass multiple duplicated arguments to CUDA Kernel

【发布时间】：2013-01-29 05:23:20

【问题描述】：

我正在寻找一种优雅的方式来在 CUDA 内核中传递多个重复的参数，

众所周知，每个内核参数都位于每个 CUDA 线程的堆栈上，因此，内核传递给每个线程的参数之间可能存在重复，即位于每个堆栈上的内存。

为了尽量减少传递的重复参数的数量，我正在寻找一种优雅的方法。

为了解释我的担忧：假设我的代码如下所示：

   kernelFunction<<<gridSize,blockSize>>>(UINT imageWidth, UINT imageWidth, UINT imageStride, UINT numberOfElements,x,y,ect...)

UINT imageWidth、UINT imageWidth、UINT imageStride、UINT numberOfElements 参数位于每个线程库中，

我正在寻找一种技巧来发送更少的参数并从其他来源访问数据。

我正在考虑使用常量内存，但由于常量内存位于 global 上，所以我放弃了它。不用说内存位置应该很快。

【问题讨论】：

"众所周知，每个内核参数都位于每个 CUDA 线程的堆栈上" - 我们根本不知道，因为内核参数要么存储在块范围共享内存中（计算1.x）或上下文范围常量内存（计算 2.x、3.x）。这个问题的整个前提都是错误的。

【参考方案1】：

内核参数通过常量内存（或 sm_1x 中的共享内存）传入，因此没有您建议的复制。

参考programming guide:

__global__ 函数参数传递给设备：

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

当然，如果您随后修改代码中的一个变量，那么您正在修改本地副本（根据 C 标准），因此每个线程都将拥有自己的副本，或者在寄存器中，或者如果需要，在堆栈。

【讨论】：

是的，我明白了，我正在寻找避免重复的技巧

@TripleS 你能解释一下你认为哪里有重复吗？如果您将参数视为只读参数，它将在需要时从常量内存加载到寄存器中。如果您修改该值，它可能最终被推入堆栈，或者它可能驻留在寄存器中。如果您的参数在许多内核启动时保持不变，您可以考虑 (a) 模板化内核，或 (b) 将值复制到 constant 变量以减少每次启动参数传递的开销。

@TripleS：你正在尝试解决一个不存在的问题。