3.1 内存层次模型：从寄存器到全局内存
2.2 内存层次模型（寄存器、共享内存、本地内存、全局内存、常量/纹理内存）
2.2 内存层次模型（寄存器、共享内存、本地内存、全局内存、常量/纹理内存）
在CUDA编程模型中，若将计算单元比作战场上的士兵，那么内存系统便是其赖以生存的后勤补给线。没有高效、合理的内存访问策略，再强大的并行计算能力也会沦为“巧妇难为无米之炊”的空谈。NVIDIA GPU的内存层次结构并非简单的堆叠，而是一套精心设计、层层嵌套、各司其职的资源调度体系。理解这一模型，不仅是优化性能的关键，更是掌握GPU计算哲学的核心。
一、为何需要如此复杂的内存层次？
现代GPU拥有成千上万的线程并发执行，若所有线程都直接访问同一块主存，带宽将成为不可逾越的瓶颈。正如一座城市若只有一条主干道，无论有多少车辆，通行效率终将受限。为此，GPU引入了多级缓存与专用存储区域——从最快但容量极小的寄存器，到速度较慢但容量巨大的全局内存，每一层都在速度、容量与共享性之间做出权衡。这种分层设计，本质上是对“局部性原理”（Locality Principle）的极致利用：时间局部性（重复访问相同数据）和空间局部性（访问相邻数据）被不同层级的内存分别捕获，从而在整体上提升数据吞吐效率。
二、寄存器：线程的私有高速缓存
寄存器是GPU中最快速的存储单元，每个CUDA核心在执行线程时，都会分配一组专用寄存器。这些寄存器直接集成在流式多处理器（Streaming Multiprocessor, SM）内部，访问延迟几乎为零。一个线程在其生命周期内所使用的局部变量（如循环计数器、中间计算结果等），只要未被显式声明为其他内存类型，通常会被编译器自动分配到寄存器中。
然而，寄存器资源并非无限。每个SM拥有的寄存器总量是固定的（例如，在Ampere架构中，每个SM最多可配置65536个32位寄存器）。当一个线程块（Block）启动时，编译器会根据内核函数的寄存器使用量决定该块能占用多少寄存器。若单个线程使用的寄存器过多，会导致每个SM能同时驻留的线程数减少，进而降低硬件的利用率——这种现象被称为“寄存器压力”（Register Pressure）。因此，过度使用局部变量或复杂表达式可能导致性能下降，尽管它们看似“免费”。
值得强调的是，寄存器是线程私有的，不同线程之间无法直接共享寄存器内容。这种隔离性保证了线程执行的独立性，但也意味着跨线程通信必须借助更高层级的共享内存或全局内存。
三、共享内存：线程块内的协作桥梁
如果说寄存器是士兵随身携带的弹药包，那么共享内存就是整个作战小队共用的弹药箱。共享内存（Shared Memory）位于SM内部，由同一个线程块中的所有线程共享，访问速度仅次于寄存器，且支持程序员显式控制。通过__shared__关键字声明的变量即驻留在共享内存中。
共享内存的核心价值在于数据重用与协作通信。典型应用场景包括矩阵乘法中的分块（tiling）、归约（reduction）操作、以及任何需要多次访问同一数据集的算法。例如，在矩阵乘法中，将输入矩阵的子块加载到共享内存后，多个线程可反复读取这些数据而无需反复访问较慢的全局内存，从而显著提升计算强度（arithmetic intensity）。
然而，共享内存的使用需谨慎处理bank conflict问题。NVIDIA GPU将共享内存划分为多个bank（通常为32个），每个bank在每个时钟周期只能服务一次访问。若多个线程同时访问同一bank的不同地址（非广播情况），则会发生冲突，导致访问串行化，性能骤降。例如，若32个线程同时访问shared_mem[threadIdx.x * 2]，由于地址间隔为2字节（假设为float类型），可能全部落在偶数bank上，引发严重冲突。解决方法包括添加填充（padding）以打破对齐模式，或重新设计数据布局。
图注：共享内存被划分为多个bank，线程块内所有线程均可访问，但需避免bank conflict。
四、本地内存：寄存器溢出的“后备仓库”
本地内存（Local Memory）是一个容易被误解的概念。它并非物理上独立的存储器，而是全局内存的一部分，仅在线程的寄存器不足以容纳其局部变量时，由编译器自动将部分变量“溢出”（spill）至此。因此，本地内存的访问速度与全局内存相当，远慢于寄存器或共享内存。
尽管名为“本地”，但它并不具备任何性能优势。实际上，大量使用本地内存通常是性能反模式的标志。开发者可通过nvcc编译器的--ptxas-options=-v选项查看每个内核的寄存器使用量和本地内存使用量。若本地内存使用量非零，应审视代码是否存在过大的数组声明（如float temp[1000];）或复杂的控制流导致寄存器分配失败。
值得注意的是，本地内存对程序员是透明的——你无法显式声明变量存储于本地内存。它的存在纯粹是编译器在资源受限下的妥协方案。
五、全局内存：海量数据的终极归宿
全局内存（Global Memory）是GPU中容量最大（可达数十GB）、速度最慢的内存层级，通常对应于设备的显存（DRAM）。所有线程均可读写全局内存，它是主机（CPU）与设备（GPU）之间数据交换的主要通道。
全局内存的性能高度依赖于合并访问（Coalesced Access）。理想情况下，当一个warp（32个线程）中的线程连续访问连续的内存地址时，硬件可将这些请求合并为少量的内存事务（memory transaction），极大提升带宽利用率。例如，若warp中线程i访问地址A + i \times 4（假设为float），则32次访问可合并为一次128字节的事务。
反之，若访问模式分散（如随机访问或跨步过大），则每个线程可能触发独立的内存请求，导致大量未使用的数据被加载（称为“wasted bandwidth”），性能急剧下降。因此，数据布局的设计（如使用Structure of Arrays而非Array of Structures）对全局内存效率至关重要。
此外，现代GPU架构（如Volta及以后）引入了L1缓存和统一的L2缓存，可在一定程度上缓解非合并访问的惩罚。但缓存命中率依赖于访问模式的可预测性，不能替代良好的内存访问设计。
六、常量内存与纹理内存：特殊用途的加速器
常量内存（Constant Memory）和纹理内存（Texture Memory）是两种具有特殊缓存机制的只读内存空间，适用于特定场景。
常量内存容量较小（通常64KB），但通过常量缓存（constant cache）进行广播。当一个warp中的所有线程访问相同的常量地址时，只需一次内存读取，数据被广播给所有线程，效率极高。然而，若warp内线程访问不同的常量地址，则会串行化访问，性能反而不如全局内存。因此，常量内存最适合存储所有线程共同使用的参数，如物理常数、配置参数等。
纹理内存最初为图形渲染设计，现已泛化为一种具有空间局部性优化的只读内存。它通过专门的纹理缓存（texture cache）提供高效的二维或三维空间邻近数据访问。即使访问模式不完全连续，只要在局部区域内，纹理缓存也能有效提升命中率。此外，纹理内存支持硬件插值（如双线性插值）和边界处理（如clamp、wrap），在图像处理、科学可视化等领域仍有独特价值。尽管在通用计算中其使用已逐渐被全局内存+L1/L2缓存取代，但在某些特定算法中仍具优势。
七、性能权衡与最新进展
不同内存层级的性能差异巨大。以NVIDIA A100 GPU为例，寄存器带宽可达数TB/s，共享内存约10–20 TB/s，而全局内存峰值带宽约为2 TB/s。这种数量级的差距决定了：算法设计必须围绕内存访问模式展开，而非单纯追求计算密度。
近年来，NVIDIA在内存层次上持续创新。Ampere架构引入了异步复制引擎（Async Copy Engines），允许在计算的同时从全局内存预取数据到共享内存，隐藏内存延迟。Hopper架构进一步推出线程块集群（Thread Block Clusters）和共享内存持久化（Persistent Shared Memory），允许多个线程块协作使用更大的共享内存池，突破传统block级别的限制。
更值得关注的是统一内存（Unified Memory）的发展。通过cudaMallocManaged分配的内存可在CPU和GPU间自动迁移，简化了编程模型。结合内存访问提示（如cudaMemAdvise）和预取（cudaMemPrefetchAsync），开发者可在保持简洁代码的同时逼近手动内存管理的性能。这标志着GPU内存模型正从“显式控制”向“智能协同”演进。
八、结语：内存即算法
在GPU的世界里，内存访问模式往往比算法本身的复杂度更能决定程序的成败。一个O(n)的算法若伴随大量非合并全局内存访问，其实际性能可能远逊于一个O(n log n)但具备完美局部性的实现。因此，CUDA程序员必须像建筑师一样思考：如何在寄存器、共享内存、全局内存之间搭建高效的数据流水线？如何让每一个字节的移动都服务于计算的洪流？
未来的GPU架构将继续模糊内存层级的界限——更大的共享内存、更快的片上网络、更智能的缓存策略，都在指向一个目标：让开发者更专注于问题本身，而非底层细节。然而，理解当前内存层次的本质，仍是驾驭这一强大工具的前提。毕竟，唯有知其所以然，方能游刃有余于并行之海。