在内核中设置固定大小的数组时，例如：
int my_array[100];
数组最终在哪个内存空间中？
特别是，我想知道这样的数组是否可以存储在 >= 2.0 设备上的寄存器文件或共享内存中，如果可以，要求是什么。
问问题
1407 次
2 回答
8
对于 Fermi（可能还有更早的架构），要将数组存储在寄存器文件中，必须满足以下条件：
- 数组仅用常量索引
- 有可用的寄存器
- 希望编译器也进行一些分析以确定对整体性能的影响
(1) 的原因是寄存器索引直接在 SASS 指令中编码。没有办法间接寻址寄存器。
限制（2）的寄存器数量的主要因素是：
- SASS 指令仅包含 6 位用于寄存器索引，这将内核中可以使用的寄存器数量限制为 64 个。实际数量为 63，因此保留一个用于某些用途。
- 一个 SM 有一个寄存器块，由同时运行的所有线程共享。
- 保存变量也需要寄存器，因此编译器必须平衡寄存器的使用以获得最佳的整体性能。
(1) 的潜在解决方法是循环展开。如果循环使用循环计数器作为数组的索引，则展开循环（使用#pragma unroll或手动）会导致数组索引变为常量，因为现在每个数组访问都有单独的 SASS 指令。
部分基于此 NVIDIA 演示文稿：本地内存和寄存器溢出。该文档还详细介绍了变量和数组的位置如何影响性能。
于 2012-06-09T02:39:36.310 回答
3
内核中的本地数组，正如您定义的那样，当没有足够的寄存器时，分配在寄存器和本地内存中。
如果要在共享内存中分配数组，则必须按如下方式指定：
__shared__ int my_array[100];
于 2012-06-08T16:19:12.940 回答