看来我误解了本地内存，认为访问速度很快，并且在内核中分配一个大数组可以利用它，这将是一件好事。然而，经过一番谷歌搜索后，看起来本地内存实际上是全局内存的一部分，因此访问速度会很慢。所以现在我试图准确理解当我在内核中分配一个大数组时会发生什么。考虑到寄存器非常有限，数组肯定无法容纳在该空间中。剩余部分是否会溢出并写入本地分配的全局内存空间？它的各个部分是否根据需要移入和移出寄存器？如果我必须在内核中多次读写它，那么它是否以某种方式缓存，可以减轻它是全局内存的事实？线程实际可以使用多少内存来动态分配数组，这个限制只是全局内存量除以线程数吗？
如果这对你答案中的数字很重要，我正在使用 V100。谢谢！
本地指的是逻辑空间（至少，这是我要开始的地方）。当您像这样创建线程局部变量时，就会发生这种情况：
__global__ void k(int *d, ...){
int a = *d;
...
}
该变量a位于逻辑局部空间中。
这样的变量可能有一个或多个物理支持。此处的描述可能因所使用的 GPU 略有不同。要获得 GPU 到 GPU 的变化，请参阅调整指南以了解详细信息。
几乎可以肯定的是，除非编译器完全优化它，否则该变量将在其生命周期中的一部分时间物理上支持在 GPU 寄存器中。CUDA 程序员无法直接控制这个或这些选择。
根据您正在执行的其他操作，编译器可能会将此变量定位或“溢出”到本地内存。这可能对物理产生影响的第一个地方是 L1 缓存。它可能会也可能不会从 L1 缓存转换到 L2 缓存。如果它找到进入 L2 缓存的路径，它可能会也可能不会真正存储在“最终”物理支持中：GPU DRAM 内存。
大多数 CUDA 程序员会认为寄存器支持（可能还有 L1 或 L2 支持）是“快速”的。寄存器是最快的。CUDA 程序员可能会认为 GPU DRAM 内存的支持至少与寄存器相比“慢”。
同样，作为 CUDA 程序员，您无法直接控制这些决策。编译器（以及硬件，当涉及到缓存驱逐时）会为您做出决定。
您要做的决定是使用（或不使用）逻辑本地空间。使用（或不使用）本地存储及其上述表现肯定会对性能产生影响。
看起来本地内存实际上是全局内存的一部分
我会通过说局部空间和全局空间在逻辑上是不同的来区分这一点。然而，两者最终都可能得到 GPU DRAM 内存的支持。（在某些情况下，全局空间也可以由主机内存支持，甚至可以由驻留在不同 GPU 上的内存支持。）
剩余部分是否会溢出并写入本地分配的全局内存空间？
可能是的（为了正确性，将“GPU DRAM”替换为“全局内存”）。编译器可以告诉您它决定执行的溢出加载和存储 ( -Xptxas -v)。
它的各个部分是否根据需要移入和移出寄存器？
是的，由编译器调度，通过创建实际指令将寄存器内容移动到内存，反之亦然。
如果我必须在内核中多次读写它，那么它是否以某种方式缓存，可以减轻它是全局内存的事实？
是的，缓存的作用可能取决于特定的 GPU。总是涉及L2缓存，也可能涉及L1缓存。
线程实际可以使用多少内存来动态分配数组，这个限制只是全局内存量除以线程数吗？
是的，这是硬件限制，不，你的计算并不总是正确的。有一个硬件上限，可以用 检索deviceQuery，并且还有一个必须满足的计算。简而言之，每个线程所需的本地内存不能超过可用GPU内存除以可以活动的线程总数（SM 数量乘以每个 SM 的最大线程数）。GPU线程可能需要堆栈，并且堆栈也存在于逻辑本地空间中，并影响线程对本地内存的需求。
尽管您没有直接询问它，但还有一个关于如何获取本地数组的主题：
int a[16];
完全由编译器在寄存器中处理。如上所述，不具备这样做的充分条件；这取决于您正在做什么以及编译器可能选择做什么。编译器通常具有控制启发式方法来最大限度地提高代码的性能，因此编译器未在寄存器中定位数组这一事实并不一定是缺陷。即使您已经满足了必要的条件，也无法“强制”此编译器行为。
然而，编译器可能需要在寄存器中定位数组的必要条件。
（上一个链接提供了一个相当复杂的案例的比较示例，显示了编译器无法在编译时确定索引的一种实现，以及可以在编译时确定索引的另一种实现。我并不是建议您应该转置数组，或以这种方式转置。这只是用于理解/学习。）