如何设计GPU芯片
Chapter 5 Gpu Memory System Design

第 5 章：GPU 内存系统设计

图形处理单元 (GPU) 已经发展成为高度并行、可编程的加速器,能够在广泛的应用程序上实现高性能和能源效率。内存系统是现代 GPU 架构的关键组件,因为它必须为大量并发线程提供快速的数据访问。在本章中,我们将探讨 GPU 内存系统设计的关键要素,包括 GPU 中使用的 DRAM 技术、内存控制器和仲裁、共享内存和缓存,以及高效内存利用的技术。

GPU 的 DRAM 技术

动态随机存取内存 (DRAM) 是现代计算系统(包括 GPU)主内存实现的主要技术。DRAM 具有高密度和相对较低成本的特点。然而,DRAM 也具有较高的访问延迟和较低的带宽,与片上存储器如缓存和寄存器文件相比。

GPU 通常采用经过优化的 DRAM 技术,以实现高带宽而非低延迟。GPU 中使用的一些常见 DRAM 技术包括:

  1. GDDR (Graphics Double Data Rate): GDDR 是一种专门为图形卡和游戏机设计的专用 DRAM 技术。它通过使用更宽的总线和更高的时钟速度,提供了比标准 DDR DRAM 更高的带宽。GDDR5 和 GDDR6 是最新的版本,提供高达 512 GB/s 和 768 GB/s 的带宽。

  2. HBM (High Bandwidth Memory): HBM 是一种高性能的 3D 堆叠 DRAM 技术,提供非常高的带宽和低功耗。HBM 将多个 DRAM 芯片堆叠在彼此之上,并通过硅通孔 (TSV) 进行连接,使数据传输速率远高于传统 DRAM。HBM2 可提供高达 1 TB/s 的带宽。

图 5.1 展示了传统 GDDR 内存和 3D 堆叠 HBM 之间的区别。

   GDDR 内存
```HBM内存
  ____________                   ______________________  
 |            |                 |  ___________________  |
 |   DRAM     |                 | |                   | |
 |   Chips    |                 | |      DRAM Dies    | |
 |            |                 | |___________________| |
 |            |                 |           .          |
 |            |                 |           .          | 
 |            |                 |           .          |
 |____________|                 |  ___________________  |
      |                         | |                   | |
     PCB                        | |  Logic Die (GPU)  | |
                                | |___________________| |
                                |______________________|

图5.1：GDDR和HBM内存架构的比较。

DRAM技术的选择取决于GPU的具体需求,如功耗预算、尺寸因素和目标应用。用于游戏和专业图形的高端GPU通常使用GDDR6,因为它具有高带宽,而HBM2则更常见于数据中心和HPC GPU,因为它更注重能源效率。

## 内存控制器和仲裁

内存控制器负责管理GPU和外部DRAM之间的数据流。它们处理来自GPU核心的内存请求,调度DRAM命令,并优化内存访问模式,以最大化带宽利用率和最小化延迟。

GPU内存控制器通常采用多通道设计,以提供高带宽和对DRAM的并行访问。每个内存通道连接到一个或多个DRAM芯片,并拥有自己的命令和数据总线。内存控制器将内存请求分布在可用通道上,以最大化并行性并避免通道冲突。

图5.2显示了一个具有四个通道的简化GPU内存控制器。

GPU核心 | | | | | 内存内存控制器 | | 控制器 | |_____________| | | | | Ch0 Ch1 Ch2 Ch3 | | | | DRAM DRAM DRAM DRAM

图 5.2: 具有四个通道的 GPU 内存控制器。

内存仲裁是决定在有多个未完成请求时应该先服务哪些内存请求的过程。 GPU 采用各种仲裁策略来优化内存系统性能和公平性:

1. **先到先服务 (FCFS)**: 最简单的仲裁策略, 请求按到达顺序服务。 FCFS 是公平的, 但由于缺乏请求重新排序, 可能会导致性能欠佳。

2. **轮流 (RR)**: 请求按循环顺序服务, 确保所有请求方都有平等的优先级。 RR 提供公平性, 但可能无法针对局部性或紧急性进行优化。

3. **基于优先级的**: 根据各种标准(如请求类型(如读与写)、来源(如纹理与 L2 缓存)或请求的老化)为请求分配优先级。 优先级较高的请求先被服务。

4. **截止时间感知**: 根据请求的截止时间进行调度, 以确保及时完成。 这对实时图形应用程序尤其重要。

5. **局部性感知**: 内存控制器尝试将访问相邻内存位置的请求一起调度, 以最大化行缓冲区命中率并最小化 DRAM 预充电和激活开销。

先进的 GPU 内存控制器通常采用这些仲裁策略的组合, 以实现性能、公平性和实时要求之间的最佳平衡。

## 共享内存和缓存

GPU 采用分层内存系统, 包括软件管理和硬件管理的缓存, 以减少主存储器的延迟和带宽需求。

### 共享内存

共享内存是一个软件管理的内存空间, 在线程块(NVIDIA)或工作组(AMD)的线程之间共享。 它具有很低的访问延迟,下面是该 Markdown 文件的中文翻译。对于代码部分,只翻译注释,代码本身不翻译。

kgroup（OpenCL）作为一种用户控制的缓存,允许程序员明确地管理线程块内的数据移动和重用。

共享内存通常使用快速、多端口的 SRAM 银行来实现,以提供低延迟、高带宽的访问。每个银行每个周期可以服务一个内存请求,所以硬件必须仲裁同时访问同一个银行的并发访问,以避免冲突。

图 5.3 展示了 GPU 核心中共享内存的组织结构。

线程块

| _________________ | | | 线程 0 | | | || | | . | | . | | . | | _________________ | | | 线程 N-1 | | | || | |_______________| | | | | | 共享内存 | | ____________ | | | 银行 0 | | | |____| | | | 银行 1 | | | || | | . | | . | | . | | | 银行 M-1 | | | |__________| | ||

图 5.3: GPU 核心中共享内存的组织结构。

合理使用共享内存可以显著提高 GPU 内核的性能,通过减少对较慢的片外 DRAM 的访问。但是,它需要仔细的编程来确保高效的数据共享和避免银行冲突。

### 硬件管理的缓存

除了软件管理的共享内存之外,GPU 还使用硬件管理的缓存来自动利用数据局部性并减少对 DRAM 的访问。GPU 中最常见的硬件管理缓存类型有:

1. **L1 数据缓存**: 一个小的、每个核心专有的缓存,用于存储最近访问的全局内存数据。L1 缓存通常专属于每个 GPU 核心,用于减少全局内存访问的延迟。

2. **纹理缓存**: 一种专门设计用于优化对只读数据的访问的缓存。这是一个关于 GPU 纹理数据缓存的 Markdown 文件。

纹理缓存针对 2D 空间局部性进行了优化,并支持硬件加速的过滤和插值操作。

3. **常量缓存**:一个小的只读缓存,用于存储频繁访问的常量数据。常量缓存会被广播到一个 warp 中的所有线程,这使得共享这些数据的多个线程很高效。

4. **L2 缓存**:一个更大的共享缓存,位于 GPU 核心和主内存之间。L2 缓存存储了从 L1 缓存中逐出的数据,用于减少对 DRAM 的访问次数。

图 5.4 展示了一个典型的 GPU 内存层次结构,包括硬件管理的缓存。

GPU 核心 0 GPU 核心 1 GPU 核心 N-1

| | | | | | | L1 数据 | | L1 数据 | | L1 数据 | | 缓存 | | 缓存 | | 缓存 | || || || | | | | | | | 纹理 | | 纹理 | | 纹理 | | 缓存 | | 缓存 | | 缓存 | || || || | | | | | | | 常量 | | 常量 | | 常量 | | 缓存 | | 缓存 | | 缓存 | || || |______________| | | | |_________________|_________________| | | | | | L2 缓存 | |_____________| | | 主内存

图 5.4: 带有硬件管理缓存的 GPU 内存层次结构硬件管理的缓存有助于通过自动利用数据局部性和减少DRAM访问次数来提高GPU应用程序的性能。然而,它们也可能引入缓存一致性和一致性的挑战,特别是在CUDA和OpenCL等并行编程模型的背景下。

## 有效利用内存的技术

有效利用GPU内存系统对于实现高性能和能源效率非常关键。优化GPU应用程序内存使用的一些关键技术包括:

1. **合并访问(Coalescing)**: 安排线程蜂群中的内存访问到相邻的内存位置,允许硬件将它们合并成一个更宽的内存事务。合并访问可以最大化DRAM带宽的利用率,并减少内存事务的数量。

2. **数据布局优化**: 以内存中的方式组织数据结构,以最大化空间局部性并最小化缓存未命中。这包括诸如结构数组(SoA)布局的技术,它将相同类型的数据元素分组在一起,以及结构数组(AoS)布局,它将属于同一结构的数据元素保持在一起。

3. **缓存和预取**: 有效利用硬件管理的缓存,通过利用内存访问模式中的时间和空间局部性来实现。这可以通过数据分块(data tiling)等技术来实现,数据分块将数据划分为更小的块,使其适合缓存,以及软件预取(software prefetching)等技术,它可以在需要之前将数据显式加载到缓存中。

4. **内存访问调度**: 重新排序内存访问,以最大化行缓冲区命中率,并最小化DRAM预充电和激活开销。这可以通过内存控制器中的硬件机制或通过访问模式优化和数据布局转换等软件技术来实现。

5. **压缩**: 应用数据压缩技术,以减小在内存和GPU核心之间传输的数据量。这可以有效降低内存访问延迟,从而提高性能和能源效率。缓解带宽瓶颈和减少与数据移动相关的能源消耗。

6. **内存虚拟化**: 采用虚拟内存技术为 GPU 应用程序提供统一、连续的地址空间。这允许更灵活的内存管理,并支持按需页面交换等功能,有助于减少内存占用和提高系统利用率。

图 5.5 展示了这些技术在 GPU 内存系统中的应用。

GPU 核心 | | | | | 内存合并 | || | | | | | 数据布局优化| || | | | | | 缓存和预取 | || | | | | | 内存访问 | | 调度 | || | | | | | 压缩 | || | | | | | 内存 | | 虚拟化 | || | DRAM

图 5.5: GPU 内存系统中提高内存利用效率的技术。

1. **内存合并**: 将同一个 warp 中的线程访问相邻内存位置的操作合并为一个更宽的内存事务。这样可以最大化 DRAM 带宽的利用率,并减少内存事务的数量。

   示例:
   ```c
   // 未合并的访问模式
   int idx = threadIdx.x;
   float val = input[idx * stride];
   
   // 合并后的访问模式
   int idx = threadIdx.x;
   float val = input[idx];
  1. 数据布局优化: 合理组织内存中的数据结构,以最大化空间局部性,减少缓存缺失。这包括结构体数组 (SoA) 布局和数组结构体 (AoS) 布局等技术。Here is the Chinese translation for the provided markdown file, with the code comments translated:

碎片整理,它将属于同一结构的数据元素保持在一起。

示例:

// 数组结构(AoS)布局
struct Point {
    float x;
    float y;
    float z;
};
Point points[N];
 
// 结构数组(SoA)布局
struct Points {
    float x[N];
    float y[N];
    float z[N];
};
Points points;

  1. 缓存和预取:有效利用硬件管理的缓存,利用内存访问模式中的时间局部性和空间局部性。这可以通过数据切片技术(将数据划分为适合缓存的较小块)和软件预取(在需要之前将数据预加载到缓存)来实现。

    示例:

    // 数据切片
for (int i = 0; i < N; i += TILE_SIZE) {
    for (int j = 0; j < N; j += TILE_SIZE) {
        // 处理一个适合缓存的数据切片
        for (int ii = i; ii < i + TILE_SIZE; ii++) {
            for (int jj = j; jj < j + TILE_SIZE; jj++) {
                // 对 A[ii][jj] 执行计算
            }
        }
    }
}

  2. 内存访问调度: 重新安排内存访问以最大化行缓冲命中率,并最小化DRAM预充电和激活开销。这可以通过内存控制器中的硬件机制或者访问模式优化和数据布局转换等软件技术来实现。

  3. 压缩: 应用数据压缩技术,减小在GPU核心和内存之间传输的数据大小。这可以缓解带宽瓶颈,并降低数据移动所需的能耗。

    示例:

    • 差分编码: 存储连续值之间的差异,而不是实际值。
    • 游程编码: 用单个实例和计数来替换重复的值。
    • 霍夫曼编码: 为更频繁出现的值分配更短的位序列。

  4. 内存虚拟化:虚拟化**: 利用虚拟内存技术为 GPU 应用程序提供统一的连续地址空间。这允许更灵活的内存管理,并支持按需分页等功能,从而可以减少内存占用,提高系统利用率。

    示例:

    • CUDA 中的统一虚拟寻址 (UVA)：使 GPU 线程可以直接使用单个指针访问 CPU 内存,简化了异构系统中的内存管理。

多芯片模块 GPU

随着 GPU 性能和功耗需求的不断增加,传统的单芯片设计可能无法满足需求。多芯片模块 (MCM) 设计,即将多个 GPU 芯片集成到单个封装中,已成为解决这一问题的一种有前景的方案。

MCM GPU 设计提供了以下几个优点:

  1. 更高的内存带宽：通过集成多个内存堆栈或芯片,MCM GPU 可以提供比单芯片设计显著更高的内存带宽。

  2. 更好的可扩展性：MCM 设计允许集成更多计算单元和内存控制器,使 GPU 能够扩展到更高的性能水平。

  3. 更高的良品率和成本效益：MCM 设计中较小的单个芯片可以实现更高的制造良品率,并且相比大型单片芯片更具成本效益。

然而,MCM GPU 设计也带来了新的挑战,如:

  1. 芯片间通信：芯片之间高效的通信至关重要。需要高带宽、低延迟的互联,以最小化芯片间数据传输的开销。

  2. 供电和热量管理：MCM 设计需要谨慎的供电和热量管理策略,以确保最佳性能和可靠性。

  3. 软件支持：MCM GPU 可能需要对编程模型和运行时系统进行修改,以充分利用多芯片架构的优势。

该领域的研究正在这篇文章探讨了MCM GPU的设计和优化,包括内存系统架构、互联设计和资源管理等。

例如,Arunkumar等人[2017]提出了一种MCM GPU设计,使用高带宽、低延迟的互联将多个GPU芯片连接起来。作者还提出了一种内存系统架构,利用MCM设计增加的带宽和容量来提高性能和能源效率。

另一个例子是Milic等人[2018]的工作,提出了一种针对MCM GPU的资源管理方案,旨在提高资源利用率并减少芯片间通信开销。该方案结合硬件和软件技术来监控应用程序的资源使用和通信模式,并做出动态资源分配决策。

结论

内存系统是现代GPU架构的关键组件,其设计和优化可以显著影响整个系统的性能和效率。随着并行工作负载的需求不断增长,研究人员正在探索各种技术来提高GPU内存系统的性能、可扩展性和适应性。

这个领域的一些关键研究方向包括内存访问调度和互联设计、缓存效率、内存请求优先级和缓存旁路、利用inter-warp异质性、协调的缓存旁路、自适应缓存管理、缓存优先级、虚拟内存页面放置、数据放置,以及多芯片模块设计。

通过探索这些和其他技术,研究人员旨在开发能够满足并行工作负载日益增长需求的GPU内存系统,同时保持高性能和能源效率。随着GPU不断发展并在机器学习、科学计算和数据分析等领域找到新的应用,内存系统的设计和优化将变得越来越重要。茎干仍将是研究和创新的重要领域。

Chapter 4 Gpu Memory System DesignChapter 6 Gpu Performance Metrics and Analysis