内核参数数据在哪里？答案

【问题标题】：Where does the kernel parameter data resides?内核参数数据在哪里？
【发布时间】：2015-09-28 02:48:41
【问题描述】：

如题，在cuda程序中，内核参数在内核启动后驻留在GPU的本地内存还是全局内存中？

例如，在 cuda 程序的 LLVM IR 中：

__global__ kernel(int param1):

%0 = alloca int

store param1, %0

那么，在这种情况下，%0 指向哪里？本地内存还是全局内存？

另外，我看到有时内核参数被保存并直接在寄存器中使用，而不是存储在任何内存中。这个决定是如何做出的？

【问题讨论】：

在 CUDA 7 或 7.5 支持的所有设备上，它位于 constant memory。如果您转储生成的机器代码 (SASS)，您将能够观察到这一点。
@RobertCrovella 谢谢！那么CUDA 6呢？如何查看它在 SASS 中的内存？是否为此在 SASS 中标记了特定的限定符？

标签： cuda llvm gpgpu llvm-ir

【解决方案1】：

正如 Robert Corvella 在他的评论中指出的那样：参数存储在 GPU 的持续内存中。

然而，执行alloca 并将 param1 存储到分配的空间移动将参数从常量内存复制到本地内存。 alloca 指令降低到 PTX 代码中的堆栈分配。在 clang 中，这是在代码生成期间处理函数参数的规范方法。但是，在 GPU 上，这可能（因为 PTX 在降低到 SASS 期间进行了优化，只是说：可以）导致性能损失，因为本地内存通过所有缓存级别一直到全局内存，并且比常量内存慢得多。

在 LLVM 中，您有 mem2reg 优化器通道。此过程将堆栈上的所有内存分配提升到寄存器。对于内核参数，您很可能需要这种优化。 alloca 和 store 指令从您的 IR 中消失，参数将直接使用，而不是不必要的副本。

【讨论】：