device 函数中定义的普通变量放在哪里？

Question

在 CUDA 中，我知道如果变量被定义为 __ shared __ 将被放置在共享内存中，如果变量被定义为 __ 常量 __.Also 则将被放置在常量内存中，那些是使用 cudamalloc() 分配的内存放在 GPU 全局内存中。但那些没有像 __ shared __ 、 __ constant __ 和 register 这样的前缀的变量放在哪里？例如变量i如下：

__device__ void func(){
   int i=0;
   return;
}

Answer 1

局部变量要么放在硬件寄存器中，要么放在本地内存（实际上是全局内存）中。

但是，在您的示例中，变量 i 将被编译器删除，因为它未被使用。

Answer 2

GPU 有一堆 space 专用于直接存储在 GPU 计算单元（即流式多处理器）中的许多寄存器。 寄存器不会存储在内存中，除非发生一些寄存器溢出（通常是当您在给定内核中使用了太多寄存器时）。与所有内存字节不同，寄存器没有地址。 CPU 也会发生同样的事情，除了 CPU 寄存器的数量通常比 GPU 上的少得多。例如，Intel Skylake 内核有 180 个整数寄存器和 168 个向量寄存器，而指令集架构仅限于 16 integer/vector 个寄存器。请注意，在寄存器溢出的情况下，寄存器的值会临时存储在本地内存中（如果可能，通常存储在 L1 缓存中）。这是基本 (Nvidia Fermi) GPU 的整体内存层次结构：

有关更多信息，请考虑阅读：Local Memory and Register Spilling。

Answer 3

自动变量，即函数范围内没有记忆space规范的变量，放置在以下位置之一：

优化掉后：

1.1 Nowhere - 如果实际上不需要该变量。这实际上发生了很多，因为 CUDA 函数通常是内联的，一些变量成为调用函数中变量的副本。 Example（注意 bar() 的编译中 foo() 中的 x - 完全消失了）。

1.2 程序编译代码中的立即值 - 如果变量的值是常量，并且没有得到更新，它的值可能只是被“插入”到代码中。 Here's an example 两个变量取常量，用常量代替它们的总和。
未优化时：

2.1 寄存器 - 如果你的变量不能是 optimized-away，更好的选择是将它保存在硬件寄存器中（对称多处理器核心之一上）GPU。 Example（变量x和y放在寄存器%r1和%r2中）。最佳和最高性能的选项，编译器

2.2 'Local' 内存 - CUDA 线程的 'local memory' 是全局设备内存中的一个区域，（原则上）只能由它访问线程。

现在，很明显，本地内存的使用速度要慢得多。那编译器什么时候选择呢？

CUDA编程指南gives us答案：

当自动变量太大无法放入当前线程的寄存器文件时（每个线程通常获得 63 到 255 个 4 字节寄存器）。
当自动变量是一个数组，或者有一个数组成员，它用non-constant偏移索引.不幸的是，NVIDIA GPU 多处理器不支持寄存器索引。
当内核过度使用其可用的寄存器配额时已经充满了其他变量或被编译代码使用——即使变量本身非常小。这被称为 寄存器溢出。

device 函数中定义的普通变量放在哪里？

Where is a ordinary variable defined inside a device function placed?

c++

cuda

__device__ 函数中定义的普通变量放在哪里？

Where is a ordinary variable defined inside a __device__ function placed?

c++

cuda

device 函数中定义的普通变量放在哪里？

Where is a ordinary variable defined inside a device function placed?