如题，在cuda程序中，内核启动后，内核参数位于GPU的本地内存还是全局内存中？
例如在一个cuda程序的LLVM IR中:
__global__ 内核(int param1):
%0 = alloca int
存储参数 1，%0
那么，在这种情况下，%0 指向哪里？本地内存还是全局内存？
此外，我还看到有时内核参数会直接保存在寄存器中并使用，而不是将其存储在任何内存中。这个决定是如何做出的？
最佳答案
正如 Robert Corvella 在他的评论中指出的那样:参数存储在 GPU 的常量内存中。
但是，
执行 alloca 并将 param1 存储到分配的空间移动将参数从常量内存复制到本地内存。
alloca 指令被降低到 PTX 代码中的堆栈分配。
在 clang 中，这是在代码生成期间处理函数参数的规范方式。然而，在 GPU 上，这可能(因为 PTX 在降低到 SASS 期间进行了优化，只是说:可以)导致性能下降，因为本地内存通过所有缓存级别下降到全局内存，并且比常量内存慢得多。
在 LLVM 中，您有 mem2reg 优化器 channel 。此遍将堆栈上的所有内存分配提升到寄存器。
对于内核参数，您很可能需要这种优化。 alloca 和 store 指令从您的 IR 中消失，参数将直接使用，而不是不必要的副本。
关于cuda - 内核参数数据存放在哪里？，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/32815036/