计算机组成与设计学习——GPU篇（五）：并行存储子系统

除GPU本身外，内存子系统是决定图形系统性能的最重要因素。图形处理工作负载要求极高的内存读写传输速率。像素写入和混合（读取-修改-写入）操作、深度缓冲区的读写以及纹理贴图读取，再加上命令与对象顶点及属性数据的读取，构成了内存访问流量的大部分。

现代GPU具有高度并行性，如图B.2.6（见本专栏第二篇文章）所示。例如，GeForce 8800能够在600MHz的频率下每时钟周期处理32个像素。通常情况下，每个像素需要进行一次4字节像素的颜色读取和写入，以及深度的读取和写入。为了生成一个像素的颜色，通常平均需要读取两到三个各为4字节的纹理元素。因此，在典型情况下，对于每一个时钟周期，需求量为28字节乘以32个像素，等于896字节。显然，这对内存系统的带宽需求是巨大的。

为了满足上述需求，GPU内存系统具有以下特征：

• 宽度大：这意味着GPU与内存设备之间有大量的引脚用于传输数据，并且内存阵列本身由许多DRAM芯片组成，以提供全总线的完整数据宽度。这种设计确保了在单个时钟周期内可以同时传输大量数据。
• 速度快：采用积极的信号处理技术来最大化每个引脚的数据速率（比特/秒）。这通常涉及高速、低电压信号传输以及先进的内存控制器设计，以实现极高的带宽性能。
• 高效利用时钟周期：GPU致力于在每一个可用的时钟周期中进行数据的读取或写入操作。为了达到这一目标，GPU并不特别追求最小化到内存系统的延迟，而是更侧重于提高吞吐量（利用率效率）。高吞吐量和低延迟本质上是相互冲突的，因此现代GPU架构往往牺牲部分延迟换取更高的带宽使用效率。
• 压缩技术：采用有损和无损两种压缩技术。有损压缩需要程序员了解其对数据可能产生的影响；而无损压缩则对应用程序透明，能够在不影响数据完整性的情况下，根据机会动态地减少所需传输的数据量。
• 缓存和工作合并结构：GPU采用缓存机制以及工作合并结构来减少芯片外部的通信流量。这些设计有助于确保数据移动所花费的时钟周期能够被尽可能充分地利用起来，通过将多个小规模或者相关的工作单元合并为一次较大规模的内存访问，从而减少内存访问次数，提升整体性能。

DRAM 考量

在设计GPU时，必须充分考虑DRAM的独特特性。DRAM芯片内部通常被组织为多个（通常是四到八个）独立的存储库，每个存储库包含一个2的幂次方数量的行（通常大约为16,384行），而每行又包含一个2的幂次方数量的位（通常为8192位）。DRAM对控制其运行的处理器施加了一系列时序要求。例如，激活一行可能需要几十个时钟周期，但一旦激活后，在接下来的四个时钟周期内就可以通过新的列地址随机访问该行内的所有数据。采用双倍数据速率（DDR）同步DRAM能够在接口时钟的上升沿和下降沿传输数据（参见第5章）。因此，对于工作频率为1GHz的DDR DRAM而言，每个数据引脚每秒可以传输2吉比特的数据。图形用DDR DRAM通常具有32根双向数据引脚，这意味着每个时钟周期可以从DRAM中读取或写入8字节数据。

GPU内部拥有大量内存访问请求生成器。逻辑图形管线的不同阶段都有各自的请求流：命令与顶点属性获取、着色器纹理加载与存取、像素深度与颜色读写等。在每个逻辑阶段，往往有多达多个独立单元来提供并行处理吞吐量，它们各自作为独立的内存请求者。从内存系统角度看，存在大量无关联的飞行中请求。这与DRAM所偏好的引用模式自然不匹配。

解决方案是让GPU的内存控制器维持指向不同DRAM存储库的独立流量堆，并等待针对特定DRAM行的足够多的待处理流量积累起来，然后一次性激活该行并将所有流量同时传输。需要注意的是，虽然累积待处理请求有利于提高DRAM行局部性从而高效利用数据总线，但这会导致请求者的平均延迟增加，因为请求会花时间等待其他请求完成。设计时必须确保没有特定请求等待过长的时间，否则某些处理单元可能会因等待数据而饥饿，最终导致相邻处理器闲置。

GPU内存子系统通常被组织为多个内存分区，每个分区包括一个完全独立的内存控制器以及一两个完全且独占式拥有的DRAM设备。为了实现最佳负载均衡并尽可能接近理论上的n个分区性能，地址被细粒度地均匀交错分布于所有内存分区中。分区交错步幅通常是一个几百字节的数据块大小。内存分区的数量旨在平衡处理器和其他内存请求者的数量，以达到最优化的设计。

Caches

GPU的工作负载通常具有非常大的工作集，生成单个图形帧时所需的数据量可能在几百兆字节的范围内。与CPU不同，构建足够大以容纳接近图形应用程序全部工作集的片上缓存并不实际可行。尽管CPU可以假设非常高的缓存命中率（99.9%或更高），但GPU面临的命中率更接近90%，因此必须处理大量飞行中的未命中情况。在这种情况下，虽然CPU可以在等待罕见的缓存未命中的时候合理地设计为暂停状态，但GPU需要在命中和未命中交错的情况下继续运行。这种设计被称为流式缓存架构。

GPU缓存必须为其客户端提供极高的带宽。以纹理缓存为例，一个典型的纹理单元每时钟周期可能会对每个像素执行两次双线性插值，而GPU中可能存在多个独立运行的此类纹理单元。每次双线性插值需要四个独立的纹理元素（texel），每个texel可能是一个64位值，常见的是由四个16位分量组成。因此，总带宽计算为2 × 4 × 4 × 64 = 每时钟周期2048比特。每个独立的64位纹理元素都分别寻址，这意味着缓存每时钟周期需要处理32个唯一地址。这就自然要求SRAM阵列采用多库或多端口排列方式，以便能够应对如此高密度和多样化的访问需求。这样的设计允许并行读取多个存储单元，从而实现高速数据传输，并满足GPU流水线中各个阶段对于内存带宽的极高需求。

MMU

现代GPU具备将虚拟地址转换为物理地址的能力。以NVIDIA GeForce 8800系列为例，其所有处理单元在生成内存地址时，都在一个40位的虚拟地址空间中进行操作。在计算过程中，加载（load）和存储（store）线程指令使用32位的字节地址，这些地址通过添加一个40位的偏移量扩展成40位的虚拟地址。

GPU内部有一个内存管理单元（Memory Management Unit, MMU），负责执行从虚拟地址到物理地址的转换工作。当处理器或渲染引擎中的层级翻译旁路缓冲器（Translation Lookaside Buffers, TLB）出现未命中时，硬件会从本地内存中读取页表信息来响应这些请求。GPU的页表条目不仅包含物理页的比特信息，还指定了每个页面所采用的压缩算法。

此外，GPU使用的页面大小范围从4KB至128KB不等，这允许根据数据的访问模式和特性选择不同的页面尺寸，以优化内存利用率和带宽效率。这种灵活的分页机制以及对虚拟内存的支持，使得GPU能够更高效地管理和利用其有限的显存资源，并支持复杂的图形渲染和并行计算任务。

存储空间

在B.3节中介绍到，CUDA（Compute Unified Device Architecture）向程序员提供了不同的内存空间，以便根据性能最优化的方式来存储数据值。接下来的讨论将以NVIDIA Tesla架构的GPU为例进行说明。

全局内存

全局内存存储在外部DRAM中，它不属于任何一个物理流式多处理器（Streaming Multiprocessor, SM），因为它的设计目的是为了不同网格中的不同并发线程阵列（Cooperative Thread Arrays, CTA）之间进行通信。实际上，引用全局内存同一位置的多个CTA可能并不会同时在GPU上执行；按照CUDA的设计，程序员无法知道CTA执行的相对顺序。

由于地址空间均匀分布于所有内存分区中，因此必须确保从任意一个流式多处理器到任意一个DRAM分区都有读写路径。

不同线程（甚至不同处理器）对全局内存的访问并不保证具有顺序一致性。线程程序看到的是放松的内存排序模型。在一个线程内部，对同一地址的内存读写操作顺序是保留的，但对于不同地址的访问顺序则可能不被保留。由不同线程请求的内存读写操作是无序的。

在同一个CTA内，可以使用屏障同步指令bar.sync来确保CTA内所有线程间的严格内存排序。membar 线程指令提供了一个内存屏障/栅栏操作，该操作会提交先前的内存访问并使其对其他线程可见，在继续执行前确保数据一致性。

线程还可以利用B.4节中（见本专栏第四节文章的图B.4.3）描述的原子内存操作来协调它们共享内存上的工作，以确保在多线程环境下数据访问的一致性和正确性。

共享内存

每个CTA（并发线程阵列）的共享内存仅对属于该CTA的线程可见，并且从创建CTA到其终止期间，共享内存才会占用存储空间。因此，共享内存可以驻留在芯片上。这种设计有多个优势。

首先，由于共享内存访问不需要与有限的片外带宽竞争，从而避免了全局内存引用时所需的带宽瓶颈。其次，在芯片上构建非常高的带宽内存结构以支持每个流式多处理器的读写需求是可行的。实际上，共享内存与流式多处理器紧密耦合在一起。

每个流式多处理器包含8个物理线程处理器。在一个共享内存时钟周期内，每个线程处理器可以处理两个线程的指令，这意味着在每个时钟周期内必须处理16个线程的共享内存请求。由于每个线程可以生成自己的地址，而且这些地址通常是唯一的，所以共享内存采用16个独立寻址的SRAM（静态随机存取存储器）库构建而成。

对于常见的访问模式，16个库通常足以保持吞吐量，但也可能存在极端情况；例如，所有16个线程可能恰好访问同一SRAM库的不同地址。为了确保任何线程通道都能将请求路由到任意一个SRAM库，需要实现一个16x16互连网络。这样就能保证即使在最不利的情况下，也能灵活地处理来自不同线程对共享内存的访问请求，维持较高的并行效率和内存访问性能。

本地内存

线程本地内存是仅对单个线程可见的私有内存。每个线程的本地内存从架构上来说比该线程的寄存器文件要大，并且程序可以计算出指向本地内存的地址。为了支持大规模的本地内存分配（请记住，总分配量是每个线程的分配量乘以活跃线程的数量），本地内存是在外部DRAM中分配的。

尽管全局内存和每个线程的本地内存位于片外，但它们非常适合在芯片上进行缓存。

常量内存

常量内存对于运行在流式多处理器（SM）上的程序来说是只读的（可以通过向GPU发送命令来写入）。它存储在外部DRAM中并在SM中缓存。由于在一个SIMT线程束中的大多数或全部线程通常会从常量内存中的同一地址读取数据，因此每个时钟周期内一个地址查找就足够了。常量缓存设计用于将标量值广播到线程束中的各个线程。

纹理内存

纹理内存用于存储大型只读数据数组。计算用的纹理具有与3D图形使用的纹理相同的属性和功能。虽然纹理通常是二维图像（像素值的2D数组），但也提供了一维（线性）和三维（体积）纹理。

计算程序通过tex指令引用纹理。操作数包括标识符以命名纹理，以及基于纹理维度的一个、两个或三个坐标。浮点坐标包含一个分数部分，用于指定采样位置，通常位于纹理元素（texel）之间的位置。非整数坐标会导致调用最接近的四个值（对于2D纹理）的双线性加权插值，然后将结果返回给程序。

纹理读取会在一个专门优化了数千个并发线程执行纹理读取的流式缓存层次结构中被缓存。某些程序利用纹理读取作为缓存全局内存的一种方式。

表面（surfaces）

表面 是一个通用术语，用于表示一维、二维或三维像素值数组及其相关的格式。定义了多种格式，例如一个像素可能被定义为四个8位RGBA整数分量，或者四个16位浮点分量。程序内核无需了解表面的具体类型。tex指令会根据表面格式将结果值转换为浮点数。

Load/Store 访问

带有整数字节寻址的加载/存储指令使得使用C和C++等传统语言编写和编译程序成为可能。CUDA程序使用加载/存储指令来访问内存。

为了提高内存带宽并减少开销，当地址位于同一块且满足对齐条件时，局部和全局加载/存储指令会将来自相同线程束的各个并行线程请求合并成单个内存块请求。将单个小内存请求合并成大的块请求，相比单独请求可以显著提升性能。大量的线程数，加上对外部DRAM中实现的本地和全局内存支持大量未完成的负载请求，有助于覆盖从加载到使用的延迟。

ROP（光栅操作处理器, Raster Operation Processors）

如图B.2.6所示（见本专栏第二篇文章），NVIDIA Tesla架构GPU包含了可扩展的流式处理器阵列（SPA），该阵列负责执行GPU的所有可编程计算任务；同时还有一个可扩展的内存系统，该系统包括了外部DRAM控制器和固定功能光栅操作处理器ROPs。ROP直接在内存上进行颜色和深度帧缓冲区的操作。

每个ROP单元都与特定的内存分区相匹配。这些ROP分区通过互连网络接收来自流式多处理器（SM）的数据。每一个ROP单元负责执行深度和模板测试及更新，并进行颜色混合运算。ROP和内存控制器协同工作，实现无损的颜色和深度压缩（最高可达8:1的比例），以此来减少对外部带宽的需求。此外，ROP单元还能在内存上执行原子操作，这对于同步多线程访问共享数据至关重要。