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