【发布时间】：2012-06-12 17:23:36
【问题描述】：
在内核中设置固定大小的数组时，如：
int my_array[100];
数组在哪个内存空间结束？
特别是，我想知道这样的数组是否可以存储在 >= 2.0 设备上的寄存器文件或共享内存中，如果可以，要求是什么。
【问题讨论】：
-
不是数组的声明方式，而是它的访问方式决定了内存的存储位置。
【发布时间】：2012-06-12 17:23:36
【问题描述】：
对于 Fermi（可能还有更早的架构），要将数组存储在寄存器文件中，必须满足以下条件：
- 数组只用常量索引
- 有可用的寄存器
- 希望编译器也进行一些分析以确定对整体性能的影响
(1) 的原因是寄存器索引直接在 SASS 指令中编码。没有办法间接寻址寄存器。
限制(2)的寄存器数量的主要因素有：
- SASS 指令仅包含 6 位用于寄存器索引，这将可在内核中使用的寄存器数量限制为 64。实际数量为 63，因此保留一个用于某些用途。
- 一个 SM 有一个寄存器块，由同时运行的所有线程共享。
- 还需要寄存器来保存变量，因此编译器必须平衡寄存器的使用以获得最佳的整体性能。
(1) 的潜在解决方法是循环展开。如果循环使用循环计数器作为数组的索引，则展开循环（使用#pragma unroll 或手动）会使数组索引变为常量，因为现在每个数组访问都有单独的 SASS 指令。
部分基于此 NVIDIA 演示文稿：Local Memory and Register Spilling。该文档还详细介绍了变量和数组的位置如何影响性能。
【讨论】：
-
对于 Kepler 和 Maxwell 微架构来说，寄存器不能被间接寻址仍然是正确的吗？
内核中的本地数组，正如您定义的那样，分配在寄存器中，当没有足够的寄存器时分配在本地内存中。
如果你想在共享内存中分配数组，你必须指定它如下：
__shared__ int my_array[100];
【讨论】：
-
添加
__shared__限定符不仅会更改存储，还会将数组的范围从线程本地更改为在块中的所有线程之间共享。 -
所以编译器更喜欢将数组存储在寄存器文件中，但是如果寄存器成为占用的限制因素，数组会被推送到本地内存吗？
-
@RogerDahl 据我所知，是的，编译器会尝试使用寄存器，然后使用本地内存。