GPU-架构基础笔记-全-

GPU 架构基础笔记（全）

001：引言 🚀

在本节课中，我们将介绍本系列将要学习的书籍《通用GPU架构》，并深入探讨GPU的非图形计算方面。这是计算机架构合成讲座的一部分，其中包含许多优秀内容。

概述

我们将探讨GPU架构的基础知识。许多人熟悉GPU，但关于其底层架构的优秀参考资料并不多。首先，我们需要思考GPU兴起的背景，以及它们在机器学习等工作负载中备受关注的原因。同时，我们也会探讨GPU是否会取代CPU的问题。

背景：并行时代的到来

上一节我们提到了GPU兴起的背景，本节中我们来看看推动这一转变的具体技术因素。

2005年左右，一个重要的变化发生了：登纳德缩放定律开始失效。登纳德缩放定律是我们依赖以获得代际性能提升的关键，它关系到我们能否在电压减半的同时使频率翻倍。但到了2005年左右，频率提升变得异常困难，我们遇到了所谓的“功耗墙”。这个定律可以追溯到1974年。

当晶体管难以继续缩小、时钟频率难以提升时，我们该怎么办？一个重要的方法是利用并行性。CPU主要利用指令级并行性，即通过重排序缓冲区等技术找出指令间的依赖关系，发掘固有的并行性。但并行性存在于不同层面，例如数据级并行性和任务级并行性。

因此，我们进入了并行时代。大约在2005年，多核处理器开始兴起，并成为当今的主流趋势，每代处理器的核心数量都在增加。

能效与向量硬件

上一节我们讨论了并行性的重要性，本节中我们来看看如何通过硬件设计来提升能效。

2010年左右，如果我们无法提升时钟频率，如何满足新硬件对能效的需求？很大程度上依赖于转向向量硬件，即进行并行计算。即使是CPU也集成了向量硬件，例如MMX扩展和SSE硬件，它们专门用于协助并行计算。

GPU是这种并行计算的绝佳来源，它们是大规模并行处理器。然而，提升性能不仅仅是发掘固有的并行性，另一个关键点是最小化数据移动。如果我们有大量需要处理的数据，而缓存和内存容量有限，我们就需要确保数据在正确的时间出现在正确的位置。

架构师的挑战：专用性与灵活性

计算机架构师面临的一个关键挑战是：如何在利用专用硬件带来的效率增益与支持广泛程序所需的灵活性之间取得平衡。我们不仅需要设计出非常快的硬件，还不能在出现需要加速的新应用时就将其淘汰，即使新应用与旧应用略有不同。

但这并不意味着专用硬件会消失。一个典型的例子是谷歌的张量处理单元，它是谷歌的机器学习加速器，在矩阵乘法等任务上表现出色。然而，与TPU等专用硬件相比，GPU有一个有趣的特点：GPU是图灵完备的，或者说拥有图灵完备的编程模型。这意味着GPU可以运行任意应用程序，而TPU等则限制更多。因此，GPU更加灵活。当新一代应用出现时，它不需要新的硬件，因为GPU足够通用，可以运行与之前模型显著不同的新模型。TPU则远不如GPU灵活。

这里存在一个权衡。因为GPU更灵活，所以在某些特定任务上，其性能肯定不如TPU。但对于能够充分利用GPU硬件的软件，GPU的效率可以比CPU高出一个数量级。不过，这并不意味着所有任务在GPU上都会表现更好，我们稍后会详细讨论这一点。

GPU硬件基础

那么，GPU会完全取代CPU吗？答案很可能是否定的。从根本上说，GPU和CPU做的是两件不同的事，它们为不同的目标进行了优化。

一个常见的类比是：CPU像跑车，速度非常快，但可能只能载两个人；GPU像巴士，速度相对较慢，但可以载20或30人。它们从根本上用于不同的任务。

如果我们有高度串行、依赖关系复杂、控制流复杂的代码，CPU可能是最佳选择。而对于那些依赖关系少、可以并行运行的大量任务，GPU通常是更好的选择。例如，线性代数运算在GPU上表现就非常出色。

因此，GPU更多地被视为协处理器，而非CPU的替代品。通常的模式是，CPU启动GPU上的工作，然后管理数据传输。GPU通常有自己的内存，你需要显式地将数据复制到GPU的本地内存，或者使用某种统一内存管理接口。GPU的内存与CPU的内存物理上是分开的。

CPU与GPU的分工

另一个GPU与CPU分离的重要原因是，有些任务从根本上更适合在CPU上执行。其中之一是与I/O设备交互。例如，从文件读取数据到GPU。从CPU的角度看，这很合理，因为我们通常在单线程中工作，可以打开文件并读取其内容。而在GPU上，我们通常从大量线程集合的角度思考问题。让GPU上的1000个线程都去打开同一个文件，这种想法是奇怪且不合理的。不过，已经有大量研究致力于为GPU提供文件系统支持，以及寻找更合理的方式让GPU直接进行I/O操作，例如将所有线程的请求合并为一个打开文件的请求。

典型的CPU-GPU系统设计

既然CPU不会消失，那么典型的CPU-GPU系统设计是怎样的呢？主要有两种类型。

第一种可能是最常见或最熟悉的：CPU及其关联的内存，以及GPU及其图形内存。GPU的本地内存通过某种总线与CPU连接。在很多情况下，GPU通过PCI Express总线连接到CPU。

另一种方式是CPU和GPU集成在一起，例如AMD的APU模型。这里提到的GPU并非指英特尔芯片上的集成显卡，而是特指APU中集成的更强大的GPU。特别是在共享同一物理内存方面，这变得非常有趣，尤其是在内存管理方面，如果你希望CPU和GPU可能访问相同的数据。

通常，我们称左边的为独立GPU，右边的为集成GPU。右边的设计典型如AMD的Bristol Ridge APU或移动GPU。高通的移动GPU称为Adreno GPU。

内存类型的差异

在考虑内存时，并非所有内存都是相同的。CPU的DRAM通常针对低延迟访问进行优化，而GPU的DRAM则针对高吞吐量进行优化。因此，在规格表中，你会看到CPU使用DDR4内存，而最新的GPU则使用GDDR内存，后者专门为图形和高带宽设计。对于低功耗移动设备，还有针对低功耗优化的DRAM。

执行流程

那么，在这些模型中的实际执行流程是怎样的呢？目前，一切从CPU开始。就像任何普通的C应用程序一样，你编译一个应用程序并在CPU上执行它。

通常，CPU上的部分负责调用一些GPU库来在GPU上分配内存。很多时候，你需要在CPU上分配内存，然后将其复制到GPU上分配的内存中。因此，这里存在数据复制。这对于APU架构或集成在同一芯片上、共享内存的情况来说，是一个有趣的研究领域，因为可能不需要这种复制。

CPU部分的应用协调数据的移动。很多时候，你需要显式地进行内存复制。但也有大量研究致力于实现数据的自动传输和GPU与CPU之间的实际分页，这最初是在英伟达的Pascal架构中引入的。

这是通过虚拟内存处理的，通常也称为统一内存，即你看到一个统一的地址空间。在集成在同一芯片上的系统中，有趣之处在于，你不需要程序员控制从CPU内存到GPU内存的复制。然而，在集成系统上直接共享内存会带来一些更微妙的问题，这通常与缓存一致性有关。当涉及GPU上许多线程同时访问许多数据时，缓存一致性会变得相当混乱。

GPU工作启动与内核

当然，在完成内存管理之后，我们必须告诉GPU执行某些操作。这通常通过驱动程序完成。GPU必须知道要做什么，因此你必须指定要运行的代码。我们通常将其称为内核。不要与Linux内核混淆，你可以将其想象为：你有一个主要应用程序，其中有一个执行大量操作的大外层循环，而在内部，你有一个计算密集的内层循环，你希望对其进行优化。这个程序的核心部分就是你希望加速并利用其并行性的部分，因此我们称之为内核。它只是程序中你希望利用其并行性的一小部分。

除了指定运行什么代码，你还需要指定其他事项，例如线程数量。与Linux系统上可能生成10或12个线程不同，在GPU上，我们通常谈论的是启动数千到数万个线程。当然，你还需要指定在GPU上分配的数据位于何处。

驱动程序必须进行大量的转换和管理工作，将其翻译并放置到GPU可访问的位置，以便GPU知道需要检查哪个位置以获取数据。然后，驱动程序还必须告诉GPU有新的工作需要完成，即CPU分配给它的任务。

现代GPU核心架构

现在我们已经了解了整体系统视图，让我们更深入地探讨现代GPU本身。我们将主要用两个不同的名称来指代它们：在英伟达，这些核心被称为流式多处理器；在AMD，它们被称为计算单元。这两个名称指的是非常相似，甚至几乎相同的东西。

GPU有不同的执行模型，即SIMT执行模型，它代表单指令多线程。我们将有一条指令，例如加法指令，但32个线程将在各自的数据上执行加法操作，或者按照它们对数组的索引方式执行。

GPU上的每个核心通常可以运行大约1000个线程，一般是1024个线程。在单个核心上执行的线程可以通过暂存器内存进行通信，并使用快速屏障操作进行同步。在核心内部，通常有一种称为暂存器内存的东西，在英伟达GPU中称为共享内存。你可以将其视为用户管理的缓存。如果你确切知道需要什么数据，并且希望获得最佳局部性，你可以直接将其加载进来，而不必担心它被其他请求逐出。

当然，它还有其他组件，例如指令和数据缓存，它们主要充当带宽过滤器。因为有如此多的线程同时运行，你不能持续地冲击内存系统。如果压力过大，任何系统都无法良好工作。因此，通过拥有缓存，它们或多或少起到了过滤器的作用。

在核心上运行的大量线程用于隐藏访问内存的延迟。与CPU相比，这是GPU的一个基本部分。因为GPU有如此多的线程和如此多的工作可供选择，即使需要访问内存获取数据，由于我们有成千上万的线程，我们实际上可以隐藏部分延迟，虽然不是全部，但可以做得相当不错。

GPU的布局与高计算吞吐量

让我们看看GPU的一般布局。通常，我们会有一个流式多处理器核心集群，几个核心通过某种互连网络分组在一起。它们会被分配给一个特定的内存分区，该分区对应物理内存的特定部分。由于这里有互连，我们仍然可以访问所有内存，但这些分区将物理内存分割开来。

那么，我们如何维持高计算吞吐量呢？为了维持高计算吞吐量，还必须与高内存带宽相平衡。理解我们不仅可以并行处理大量任务，还要如何与GPU架构的其他层级（尤其是内存系统）相平衡，这一点非常重要。

在GPU中，这种并行性是通过包含多个内存通道提供的。通过拥有多个内存通道，我们获得了更高的内存带宽，这不仅是因为我们使用了GDDR而不是DDR。正如我们所说，这与最后一级缓存的一部分相关联，在当今的GPU中，这就是L2缓存。在英伟达GPU中，会有L1缓存、L2缓存，然后是内存通道和后备内存。

这并不是说没有其他可能的组织方式，例如英特尔的Xeon Phi，它直接与GPU竞争超级计算市场，它将最后一级缓存与核心分布在一起。

GPU与CPU性能对比

当谈论CPU时，我们当然是在与最先进的技术进行公平比较，即高度并行工作负载上的超标量乱序执行CPU。GPU往往胜出的原因是，GPU的很大一部分面积用于算术逻辑单元，而用于控制逻辑的面积较少。需要记住的是，GPU没有分支预测等功能，也没有乱序执行。与CPU相比，GPU要简单得多。但它们以拥有极端数量的ALU来换取这一点，远超传统CPU。

但最重要的是要记住，存在一个平衡点，有些区域多核CPU表现更好，而有些区域大规模多线程的GPU表现更好。从这张图中我们可以清楚地看到，这里有一个多核区域。当我们开始增加线程数量时，会看到这里有一个下降。这个下降点代表了我们需要理解的权衡：CPU拥有非常高的频率、出色的串行处理能力、推测执行、分支预测和更高的时钟频率；而如果我们有大规模多线程数据，我们真正关心的是计算密度，我们只需要ALU，像分支预测这样的功能对我们帮助不大。但这里存在一个区域，即这个低谷，在这个区域中，CPU上利用指令级并行性的传统方法无法给我们带来更多性能，而我们拥有的线程数量又超过了CPU能有效支持的范围，但还没有达到大规模多线程的程度。如果将其放在GPU上，我们无法很好地利用可用带宽。因此，我们必须意识到这个低谷的存在，并决定如果我们处于这个低谷，是否应该通过某种方式优化代码，使其更适合CPU，或者它确实是并行的，我们只需要更努力地将其推向右侧，启动数百万个线程，真正利用其并行性。这些都是我们需要思考的问题。

在CPU上，我们会增加线程到一定程度，但缓存无法再支持这么多线程。每个线程最终都会破坏其他线程的缓存。因此，我们看到了性能下降。所以，根据平台的不同，并行性并不总是好事。同样，由于没有足够的线程，我们无法很好地隐藏片外延迟。

能效考量

随着登纳德缩放定律的终结，能效当然是当今市场的一大驱动力，有许多研究致力于设计更节能的方案。我们当然也不希望访问消耗大量能量的非常大的内存结构。因此，在谈论GPU和数千个线程时，我们必须考虑到这一点。在实际建模中，许多研究人员使用一个名为GPGPU-Sim的模拟器，它包含一个名为GPUWatch的能耗模型。

我们可以看到，在进行操作时，例如加法，只需要极少的皮焦耳能量；甚至从32KB SRAM访问32位数据也只需要大约5皮焦耳。但是，当我们从DRAM访问32字节数据时，突然需要640皮焦耳，大约是加法操作所需能量的64000倍。这确实会累积起来。因此，这强调了我们对功耗的担忧，尤其是对数据移动的担忧，特别是需要访问DRAM时。

GPU的历史背景

当然，我们应该谈谈GPU的背景，GPU并非凭空出现，它们有历史。早在20世纪60年代，计算机和计算机图形学刚刚兴起时，GPU就开始起步。早期GPU的许多驱动因素当然是动画渲染。

早期的显卡始于IBM的单色显示适配器，它只支持文本。然后我们逐渐发展到2D和3D加速。在谈论早期3D技术时，记住英伟达的GeForce 256总是很重要的。然而，最终我们开始通过顶点着色器和像素着色器为GPU提供一些可编程性，这更多是图形方面的内容。

当时的人们非常聪明，他们说：“我有这个并行硬件，为什么不尝试用它做些不同的事情呢？如果我不太关心图形呢？”于是，人们找到了非常有趣的方法，将线性代数映射到着色器上，将矩阵数据映射到纹理中并应用着色器，基本上是利用图形处理器进行通用计算。我们谈论的不是通用GPU的时代，而是严格的图形处理器时代。人们找到了一些巧妙的方法开始在其上进行线性代数等运算。

这真正推动了市场。人们如此广泛地进行这种操作，以至于像英伟达这样的公司表示：“如果你们想要这个，我们非常乐意提供。”例如，最终在GeForce 8系列中，出现了一些创新，例如能够从着色器写入任意内存地址以及暂存器内存。我们开始释放内存系统。然后在Fermi架构中启用了读写数据的缓存，现在我们开始接近更现代的架构。

还有其他创新，例如AMD的Fusion架构，将CPU和GPU集成在同一芯片上，并实现动态并行性，例如GPU自身启动工作。通常我们讨论的是CPU启动GPU上的工作，但如果GPU在某些情况下可以启动自己的工作呢？当然，最近引入并在Turing架构中进一步优化的功能包括张量核心，专门针对特定的非图形相关加速。

本系列后续内容

那么，在本系列中我们还将涵盖哪些内容呢？如果我们打算使用GPU，我们必须知道GPU运行什么。我有另一个专门关于CUDA编程的系列，但我们也将在本系列中详细讨论编程模型，不仅针对英伟达GPU，也针对AMD GPU，我们会发现它们非常相似，即使它们的语言差异很大。

当然，然后我们必须深入探讨真正酷的部分，即探索架构。具体来说，首先是GPU核心及其如何执行数千个线程。然后我们需要了解如何构建这种架构以提供高吞吐量和灵活的编程模型。我们还将介绍一些与GPU架构相关的最新研究，特别是针对核心本身的研究。

之后，在第4章中，我们将讨论GPU内存系统，主要是核心内的缓存，例如GPU核心内的L1缓存，以及一些更内部的组织结构，内存分区。当然，就像讨论核心一样，我们将总结一些与GPU架构相关的最新研究。

最后，我们将讨论一些当代的研究问题，以及我在GPU架构方面的工作，以及当前GPU架构领域人们正在关注的一些重大未解决问题。

本节课中我们一起学习了GPU架构的引言部分，包括其历史背景、基本概念、与CPU的对比、系统设计以及本系列后续内容的概览。

002：编程模型（第一部分） 🚀

在本节课中，我们将深入探讨GPU的编程模型。理解编程模型是学习GPU硬件架构的基础，因为编程方式直接依赖于底层硬件，而硬件设计也受到编程方式的影响。我们将从基础概念开始，逐步了解如何编写GPU程序。

概述

上一节我们从宏观视角了解了GPU的基本概念和历史。本节我们将聚焦于GPU的编程模型，学习如何通过API（如CUDA和OpenCL）来利用GPU的大规模并行计算能力。

执行模型基础

GPU采用宽SIMD（单指令多数据）硬件来挖掘应用程序中的数据级并行性。这意味着我们可以同时对不同的数据片段进行操作，因为它们之间没有依赖关系。

然而，我们实际上是通过CUDA（NVIDIA）或OpenCL（AMD）等API来编程这些硬件的。它们提供了一种类似MIMD（多指令多数据）的编程模型，允许程序员向GPU启动大量标量线程。

每个标量线程都可以遵循独立的执行路径。但在运行时，这些线程并非完全独立调度。它们被分组为线程束（NVIDIA称Warp，含32线程）或波前（AMD称Wavefront，含64线程），这些组内的线程以锁步方式执行相同操作。这种执行模型称为SIMT（单指令多线程）。

在SIMT模型下，我们仍可以拥有标量线程的视角，但需要通过活动掩码来指定线程束中哪些线程在某一时刻是活跃的。因此，编程时应考虑整个线程束在同时做什么，尽量避免线程束分化，即同一线程束内的线程执行不同操作。

GPU程序执行流程

一切始于CPU。CPU负责启动GPU代码的执行。

对于独立GPU（拥有独立显存），典型流程如下：

  1. 分配GPU内存：使用类似 cudaMalloc 的函数在GPU上分配内存。
  2. 数据传输：由于内存物理分离，需要使用 cudaMemcpy 等函数将数据从CPU内存复制到GPU内存。
  3. 启动内核：告知GPU要执行哪个内核函数。

对于集成GPU（与CPU共享物理内存，常见于移动设备或AMD APU），由于内存共享，无需单独分配和复制内存，只需直接启动内核。

内核与线程组织

内核通常由成千上万个线程组成。每个线程执行相同的程序，但可以有不同的控制流（例如，通过if-else语句）。

程序员通过定义网格和线程块来组织这些线程。

  • 线程块：线程的集合。有维度限制，会被调度到GPU的流多处理器上执行。
  • 网格：所有线程块的集合。

为了充分利用GPU，我们通常需要启动大量线程。线程块在幕后会被分解为线程束或波前，但程序员看到的是线程块和网格。

以下是如何计算需要启动的线程块数量的一个例子。假设每个线程块有256个线程，要处理 n 个元素：

int numBlocks = (n + 255) / 256;

这里 (n + 255) / 256 确保了即使 n 不是256的整数倍，我们也能启动足够多的线程块来覆盖所有元素，多余的线程可以通过条件判断使其不工作。

内核启动使用三重尖括号语法指定网格大小（线程块数量）和块大小（每个线程块的线程数）：

saxpy<<<numBlocks, 256>>>(n, 2.0, d_x, d_y);

SAXPY示例：从串行到并行

SAXPY是一个基础线性代数运算：Y = a * X + Y。我们通过它来对比CPU串行和GPU并行编程。

CPU串行版本核心是一个循环：

for (int i = 0; i < n; i++) {
    y[i] = a * x[i] + y[i];
}

GPU并行版本的核心是内核函数。它没有循环，每个线程根据其全局索引处理一个数据元素：

__global__ void saxpy(int n, float a, float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a * x[i] + y[i];
    }
}
  • __global__ 声明这是一个由CPU调用、在GPU上执行的函数（内核）。
  • blockIdx.x：当前线程块在网格x维度的索引。
  • blockDim.x：每个线程块在x维度的线程数（即块大小）。
  • threadIdx.x：当前线程在线程块x维度的索引。
  • int i：计算出的当前线程的全局索引，用于对应要处理的数据元素。
  • if (i < n)：防止为凑整而启动的多余线程访问越界内存。

性能优化概念

GPU编程中，程序员可以进行一些手动优化以提升性能。

共享内存/局部数据存储：GPU拥有速度极快的片上内存，NVIDIA称为共享内存，AMD称为局部数据存储。这是一种程序员可管理的暂存器内存。如果程序员能预知哪些数据会被频繁重用，可以手动将这些数据加载到共享内存中，从而获得极低的访问延迟（约几个时钟周期）。这类似于一个用户管理的L1缓存。

线程块内同步：在同一个线程块（或称为协作线程数组CTA）内，线程可以通过 __syncthreads() 这样的屏障操作进行同步。所有线程必须到达这个同步点，才能继续执行。这在协同使用共享内存时非常必要。

动态并行：NVIDIA的CUDA动态并行特性允许GPU内核在运行时启动新的内核，而不仅仅依赖于CPU来启动工作，这增加了编程的灵活性。

总结

本节课我们一起学习了GPU编程模型的基础知识。我们了解到，一个典型的GPU应用始于CPU进行内存分配和数据传输。然后，我们将问题（如SAXPY）映射到一个由线程块组成的网格上。在内核中，每个线程通过内置变量计算出自己在网格中的位置，并据此处理对应的数据。此外，我们还探讨了通过共享内存进行手动数据管理以优化性能，以及线程块内同步和动态并行等高级概念。

理解这些编程模型是后续深入学习不同GPU硬件架构细节的关键。在接下来的课程中，我们将探讨GPU的指令集架构及其在不同代际间的演变。

注：本教程内容基于公开的GPU编程知识整理，示例代码为概念演示。更深入的实践建议参考《Programming Massively Parallel Processors》等专业书籍或相关课程。

003：编程模型（第二部分）🚀

在本节课中，我们将深入探讨GPU的编程模型。我们将学习高级编程接口（如CUDA）如何被翻译成虚拟指令集（如PTX），并最终转换为机器指令集（如SASS）。同时，我们也会简要了解常用于AMD GPU的OpenCL。通过对比NVIDIA和AMD的不同实现方式，我们可以更好地理解GPU底层的工作原理。

从高级语言到虚拟指令集

上一节我们介绍了使用CUDA进行GPU编程的高级接口。本节中，我们来看看这些高级代码如何被翻译成更低层次的指令。

GPU编程的指令虚拟化始于早期的图形处理，例如OpenGL及其着色语言GLSL，以及微软的高级着色语言HLSL。对于NVIDIA GPU，其虚拟指令集架构被称为并行线程执行，简称PTX。

PTX是一个完全文档化的虚拟指令集，其设计类似于ARM或MIPS的RISC架构，因此相对易于理解。以下是一个简单的SAXPY（单精度a乘X加Y）CUDA内核代码示例：

__global__ void saxpy(int n, float a, float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a * x[i] + y[i];
    }
}

这段代码会被编译成PTX指令。PTX代码虽然更冗长，但结构清晰。它首先加载参数，然后通过特殊寄存器（如%ctaid、%ntid、%tid）计算线程在网格中的位置，接着使用setp指令设置谓词以执行条件判断（对应CUDA中的if语句），最后执行核心的加载、乘加和存储操作。

从虚拟指令集到机器指令集

PTX并非最终在硬件上执行的指令。在运行时，PTX会被进一步转换为流处理器汇编，即SASS。

这个过程由GPU驱动程序或独立的汇编器（如CUDA工具包中的ptxas）完成。与PTX不同，SASS并未完全公开其文档，这意味着开发者无法直接编写或完全理解所有SASS指令的精确语义。

尽管如此，学术界通过逆向工程对SASS进行了大量研究。例如，早期的“De CUDA”项目和后来的“MaxAs”项目都致力于理解并汇编SASS代码。NVIDIA也提供了cuobjdump工具，可以用于从二进制文件中提取SASS或PTX代码。

不同GPU架构（如Fermi、Maxwell、Pascal）的SASS指令集会存在显著差异。例如，从Kepler架构开始，NVIDIA在指令流中引入了控制指令，用于编码依赖关系、停顿计数等信息，以减少硬件中记分牌逻辑的开销。在Maxwell和Pascal架构中，大约每三条常规指令就伴随一条这样的控制指令。

此外，指令中还会包含操作数重用缓存等优化信息，旨在降低能耗或提升性能。

AMD的GCN架构与OpenCL

与NVIDIA不同，AMD为其GCN架构公开了硬件级别的指令集架构，这为学术研究和模拟器开发提供了便利。

AMD GCN架构的一个关键特点是其分离的标量和向量指令流水线。

  • 向量指令 以V为前缀，在向量单元上执行，为波前中的每个线程计算不同的32位值。
  • 标量指令 以S为前缀，在标量单元上执行，为波前中的所有线程计算一个共享的32位值。

这种设计非常高效。例如，在控制流处理中，像if (a > b)这样的条件判断对所有线程可能是一致的，只需由标量单元计算一次即可，避免了向量单元中所有线程重复计算相同值的资源浪费。

在GCN汇编中，EXEC是一个特殊的寄存器，用于谓词化执行，即通过掩码来控制波前中哪些线程在特定时刻是活跃的，以处理线程间的控制流分歧。

AMD的指令集手册还揭示了其硬件管理长延迟操作的方式。例如，通过S_WAITCNT指令，程序员可以让波前等待特定类型（如向量内存访问）的未完成操作数量降至阈值以下，从而实现数据依赖解析。

总结

本节课我们一起学习了GPU编程模型的底层转换过程。

  1. 我们首先看到高级CUDA代码如何被编译成PTX虚拟指令集。
  2. 接着，我们探讨了PTX如何在运行时被转换为未完全公开的SASS机器指令集，并了解了不同NVIDIA架构间指令集的差异以及其中嵌入的硬件优化信息。
  3. 最后，我们对比了AMD的GCN架构，了解了其分离的标量/向量流水线设计，以及其公开的指令集如何帮助理解硬件行为。

理解从高级编程模型到底层指令集的映射，是深入优化GPU程序性能和理解GPU架构细节的重要基础。下一节，我们将进入更核心的领域，探讨GPU的流多处理器核心、指令执行以及寄存器数据流。

004：SIMT核心 - 第一部分 🧠

在本节课中，我们将开始探讨GPU的实际微架构，从SIMT核心入手。SIMT（单指令多线程）是GPU并行编程模型的核心硬件实现。我们将首先聚焦于SIMT核心本身，特别是其前端调度机制，而内存系统部分将留到下一章讨论。

概述

GPU架构源于图形处理需求，需要处理巨大的纹理贴图数据，这些数据通常无法完全放入片上缓存。为了获得高性能，GPU架构设计必须能够维持高内存带宽，因此采用了GDDR甚至HBM等高带宽内存。同时，GPU需要并发执行成千上万个线程，这对缓存设计提出了挑战，因为大量线程可能访问分散的数据，难以有效利用局部性。然而，缓存仍然有效，例如在图像模糊等图形工作负载中，像素处理具有空间局部性，缓存可以过滤流量，减轻内存系统下层的拥塞。

本章的核心是理解单个SIMT核心的工作原理，它通常分为一个SIMT前端和一个SIMD（单指令多数据）后端。其调度主要涉及三个循环：指令获取循环、指令发射循环和寄存器访问调度循环。

SIMT核心组织架构

下图展示了SIMT核心的基本组织架构，清晰地划分了SIMT前端和SIMD数据通路。

  • SIMT前端：包含指令缓存（I-Cache）、解码逻辑、指令缓冲区（I-Buffer）、记分牌（Scoreboard）、发射逻辑以及SIMT堆栈。它负责获取指令、解码、管理依赖关系，并确定每个周期哪些线程是活跃的。
  • SIMD数据通路：包含操作数收集器（Operand Collector）、ALU（算术逻辑单元）和内存单元（用于加载/存储操作）。操作数收集器负责调度对寄存器文件的读取，以避免端口争用，然后将数据送入执行单元。

为了由浅入深地理解，我们将采用逐步逼近的描述方式，从最简单的单循环模型开始。

单循环近似模型

首先，我们从一个简化的视角出发，假设只有一个调度循环在工作。这类似于CUDA编程手册中向程序员抽象出的视图，但作为架构师，我们需要关心其下的硬件细节。

在这个模型中，调度的基本单位是线程束（Warp，NVIDIA）或波前（Wavefront，AMD），即一组（例如32或64个）始终一起执行同一条指令的线程。

以下是每个周期内发生的基本步骤：

  1. 选择线程束：硬件从线程块包含的多个线程束中选择一个进行调度。
  2. 获取指令：使用被选中线程束的程序计数器（PC）访问指令内存，获取下一条要执行的指令。
  3. 解码指令：解码指令，确定其操作类型以及需要使用的源寄存器。
  4. 获取操作数：从寄存器文件中读取源操作数的值。这一步会用到操作数收集器，我们稍后会详细讨论。
  5. 确定活跃掩码：并行于寄存器读取，硬件需要确定当前指令的SIMT执行掩码。这个掩码定义了线程束中哪些线程在当前控制流路径下是活跃的、应该执行该指令。
  6. 执行指令：在SIMD执行单元上执行指令。这些功能单元（如ALU、FPU、Tensor Core、Load/Store单元）通常是宽向量单元，宽度与线程束中的线程数匹配（例如32宽）。为了提高时钟频率和能效，这些单元也可以被设计得更窄并采用流水线化。

接下来，我们将深入探讨这个模型中最关键的部分：SIMT执行掩码和分支处理。

SIMT执行掩码与分支处理

SIMT模型允许程序员将每个线程视为独立的标量线程，这意味着不同线程可能进入不同的控制流路径（例如，有的线程执行if分支，有的执行else分支）。硬件如何用“单指令”来处理这种“多线程”的分支行为呢？答案是SIMT执行掩码和SIMT堆栈的组合。

现代GPU并非单纯依靠谓词（Predication）来实现，而是结合了传统谓词和SIMT堆栈来高效处理两个关键问题：嵌套控制流和完全跳过某些线程不执行的控制流路径。

为了说明这一点，让我们看一个具体的例子。考虑以下CUDA代码片段及其对应的PTX虚拟指令集表示：

// CUDA C 示例
do {
    // 块 A
    t1 = ...; t2 = ...; t3 = load(...); t4 = 0;
    if (t3 != t4) { // 外层 if-else
        // 块 B
        t5 = load(...);
        if (...) { // 内层 if-else
            // 块 C
            x += ...;
        } else {
            // 块 D
            y += ...;
        }
        // 块 E
        i++;
    } else {
        // 块 F
        z += 3;
    }
    // 块 G
} while (i < n);

在CPU上，单个线程会顺序执行并做出分支决策。但在GPU上，一个包含4个线程的线程束会以锁步方式开始执行块A中的所有指令。当遇到外层 if (t3 != t4) 时，问题就出现了：可能只有部分线程的条件为真，需要进入块B，而其他线程需要进入块F。

硬件通过SIMT执行掩码来标记哪些线程是活跃的。例如，进入外层if时，掩码可能变为 1110（假设第4个线程条件为假）。只有掩码为‘1’的线程会执行块B中的指令。当这些活跃线程再遇到内层if时，掩码可能会进一步分裂，例如变为 1100。

重新汇聚点是控制流分裂后，所有线程能够再次保证以锁步方式一起执行的位置。通常，这是引起分支的指令的直接后支配节点。在上例中，块E是内层if-else的重新汇聚点，块G是外层if-else的重新汇聚点。

SIMT堆栈工作原理

SIMT堆栈是管理这种掩码分裂与重新汇聚的核心硬件结构。堆栈中的每个条目通常包含：一个重新汇聚PC（即所有线程应重新汇合后继续执行的地址）和一个活跃掩码（表示哪些线程属于该条目的控制流路径）。

以下是处理流程：

  1. 初始状态：所有线程活跃，执行块A。遇到外层分支时，创建两个堆栈条目，分别对应路径B（掩码1110，重新汇聚于G）和路径F（掩码0001，重新汇聚于G）。通常，硬件会先执行活跃线程数多的路径（多数路径优先）。
  2. 执行与嵌套：假设先执行路径B（PC指向B，掩码1110）。执行到内层分支时，再次分裂。当前栈顶条目（对应路径B）的重新汇聚点G被保留，然后为新的内层分支创建两个新条目，例如路径C（掩码1000，重新汇聚于E）和路径D（掩码0110，重新汇聚于E）。
  3. 重新汇聚：硬件执行当前栈顶条目（例如路径C）。当执行到其重新汇聚点E时，弹出该栈顶条目。然后，下一个栈顶条目（路径D）变为当前执行路径。执行完路径D并到达重新汇聚点E后，再次弹出。此时，栈顶又回到了外层路径B的条目，但其PC现在指向了重新汇聚点E之后（即块E内的指令）。继续执行直到到达该条目的重新汇聚点G。
  4. 执行另一路径：到达G后，弹出路径B的条目。现在栈顶是路径F的条目，开始执行路径F（掩码0001）。执行完后到达重新汇聚点G，弹出该条目。此时堆栈为空，所有线程再次完全汇聚，继续执行块G之后的代码（如while循环判断）。

通过这种堆栈机制，GPU以串行化的方式执行了不同的控制流路径，同时在程序员视角下，每个线程仍然保持了独立的控制流逻辑。

总结

本节课我们一起学习了GPU SIMT核心架构的第一部分。我们首先概述了SIMT核心的整体组织，将其分为负责取指、解码和线程调度的前端，以及负责数据计算和存取的后端。然后，我们通过一个简化的单循环模型，介绍了线程束调度的基本步骤。

本章的重点是理解了GPU如何处理线程级分支。我们探讨了SIMT执行掩码的关键作用，它标识了线程束中哪些线程在当前指令下是活跃的。更重要的是，我们深入分析了SIMT堆栈的工作原理，它是硬件管理控制流分裂与重新汇聚的核心机制，通过维护重新汇聚点和活跃掩码，使得以锁步方式执行的硬件能够高效地模拟出独立线程的复杂分支行为。

在下一章中，我们将继续探讨SIMT核心的其他调度循环，并深入了解内存系统。同时，我们也会简要提及现代架构中关于SIMT死锁和无需堆栈的SIMT架构等更高级的话题。

005：SIMT核心 - 第二部分 🧠

在本节课中，我们将继续探讨GPU的核心架构。我们将从上节课结束的地方开始，即由基于栈的架构引发的“SIMT死锁”问题。根据NVIDIA发布的信息、专利文件以及当代研究，大约从2017年的Volta GPU架构开始，NVIDIA不再使用基于SIMT栈的实现来处理线程束（Warp）的分歧和同步。今天，我们将深入探讨什么是SIMT死锁、它如何源于架构，以及如何避免它。

SIMT死锁问题示例 🔄

上一节我们介绍了SIMT栈的基本工作原理，本节中我们来看看一个具体问题。让我们从一个简单的示例开始，展示使用SIMT栈的架构中存在的问题。

我们考虑一个包含32个线程的单一线程束。它初始化一个共享变量 mutex 为0，表示锁是空闲的。然后，每个线程执行一个原子比较并交换操作 atomicCAS。

int mutex = 0;
// 每个线程尝试获取锁
while(atomicCAS(&mutex, 0, 1) != 0) {
    // 等待
}
// 临界区
atomicExch(&mutex, 0); // 释放锁

atomicCAS 是一个编译器内置函数，会被翻译成PTX指令 atom.global.cas。逻辑上，它会读取第一个参数（mutex）的值，与第二个参数（0）比较。如果相等（即锁空闲），则将第三个参数（1）的值交换进去，表示锁已被占用。

由于这是原子操作，当同一个线程束内的多个线程同时对同一地址执行 atomicCAS 时，这些操作会被序列化。这意味着，在同一时刻，只有一个线程能成功获取锁（看到 mutex 为0并交换为1），其余31个线程将看到 mutex 为1，从而在 while 循环中等待。

从SIMT栈架构的角度看，这里发生了线程分歧：一个线程成功跳出 while 循环，而其他31个线程仍在循环内。这个 while 循环的出口也是线程的重新汇聚点。在标准的SIMT栈架构中，所有线程最终都会在此处重新汇聚。

问题在于，这个重新汇聚点位于释放锁的操作（atomicExch）之前。当那个成功获取锁的线程执行到汇聚点时，它的栈条目会被弹出，调度器会恢复执行仍在 while 循环中的其他31个线程。这就形成了一个循环依赖：

  • 跳出循环的线程在等待其他线程到达汇聚点才能继续执行（包括后续释放锁）。
  • 其他线程在等待锁被释放才能跳出循环。
    结果就是没有任何线程能继续执行，整个程序陷入停滞，这就是所谓的 SIMT死锁。

现代解决方案：屏障与掩码 🛡️

基于NVIDIA在2015年的一项专利以及相关研究，现代GPU（如Volta及之后架构）采用了一种非栈式的、基于屏障的方法来处理线程分歧与重汇聚。

以下是实现此机制的核心硬件字段，它们存储在寄存器中，供线程束调度器使用：

  • 屏障参与掩码：一个32位的掩码，用于跟踪线程束中哪些线程将参与某个特定的重汇聚屏障。
  • 屏障状态：用于跟踪在某个屏障处，哪些线程已经到达。
  • 线程状态：记录每个线程的当前状态（例如，准备执行、在某个汇聚屏障处阻塞、让出等）。

控制流与屏障初始化 🗺️

为了支持嵌套的控制流（如多重 if-else），线程束可能需要多个屏障参与掩码。软件（编译器）负责管理这些掩码，因为硬件管理任意深度的嵌套可能过于复杂。

屏障通过特殊的 bar.arrive 指令进行初始化。当线程束执行此指令时，所有活跃线程的对应位会在指定的汇聚屏障掩码中被设置。

// 示例PTX指令，初始化一个屏障
bar.arrive.sync 0; // 所有当前活跃线程将参与屏障0

随后，在代码的重汇聚点，会使用 bar.sync（或类似 wait）指令。

bar.sync 0; // 线程在此等待，直到参与屏障0的所有线程都到达

当线程执行 wait 指令时，它会被添加到对应屏障的状态寄存器中，并将其线程状态改为“阻塞”，直到屏障处的所有参与线程都到达。

“让出”状态是一个更抽象的概念，根据专利描述，它可能用于在线程可能陷入SIMT死锁的情况下，允许线程束内的其他线程越过汇聚屏障继续执行，从而打破循环依赖。

执行模型对比：栈式 vs 屏障式 ⏱️

让我们对比两种模型下的执行时序。

SIMT栈模型：
执行路径是严格有序的。线程束先完全执行分歧后的一个分支（如 if 块），到达汇聚点后弹出栈，再回头执行另一个分支（如 else 块），最后所有线程在最终汇聚点同步，一起执行后续代码。线程在整个过程中是“锁步”前进的。

屏障式模型（如Volta）：
执行顺序更加灵活。调度器可以在不同分支（称为“线程束分裂”）之间“乒乓”切换。例如，可以先执行 if 块的一部分指令，然后切换到 else 块执行一部分，再切回来。
这种灵活性的代价是，失去了线程的自动锁步保证。一个线程可能先执行到汇聚点后的代码（如 Z），但会被阻塞在那里（sink warp），等待其他分支的线程到达。只有当所有相关线程都到达屏障后，它们才会一起继续执行。因此，程序员有时需要使用 __syncwarp() 等显式同步指令来确保线程间的协调。

线程束调度策略 ⚙️

最后，我们简要探讨与线程束执行密切相关的调度策略。GPU通过大量线程束的切换来隐藏内存访问等长延迟操作的延迟。

一个简单的假设是轮询调度：每个周期从就绪的线程束中选择一个发出指令，依次循环。如果 线程束数量 × 线程束指令发出间隔 > 内存延迟，那么核心的执行单元就能始终保持忙碌。

然而，这存在权衡：

  • 增加线程束数量：提高并行度，更好地隐藏延迟。
  • 代价：需要存储更多线程状态（尤其是寄存器文件），在芯片面积固定的情况下，可能减少核心数量。现代GPU的寄存器文件已经非常庞大（约256KB）。

在实践中，内存延迟取决于应用程序的局部性。轮询调度在图形负载中效果很好，因为线程通常访问相同的数据（如纹理），缓存命中率高。但对于访问模式不规则、数据局部性差的负载（如某些图计算），轮询调度可能导致缓存抖动，反而降低性能。

研究表明，在这种情况下，限制同时活跃的线程束数量（即“节流”）可能更有效。这种称为“缓存感知的线程束调度”的策略，通过减少缓存冲突，让少数线程束的数据能常驻缓存，从而提高整体性能。这体现了调度策略需要根据负载特性进行优化。

总结 📚

本节课我们一起学习了以下内容：

  1. SIMT死锁：深入分析了基于SIMT栈的架构中，由于重汇聚点与同步操作顺序不当导致的循环依赖和程序停滞问题。
  2. 现代重汇聚机制：介绍了Volta及之后架构采用的、基于屏障和掩码的非栈式方法。我们了解了屏障参与掩码、屏障状态、线程状态等核心概念，以及 bar.arrive 和 bar.sync 指令的作用。
  3. 执行模型对比：对比了栈式模型的严格锁步执行与屏障式模型的灵活、非锁步执行，指出了后者需要显式同步的情况。
  4. 线程束调度：探讨了GPU通过大量线程束切换隐藏延迟的基本原理，分析了轮询调度的优缺点，并介绍了针对不同数据局部性负载的优化调度策略（如缓存感知调度）。

下一节，我们将基于更复杂的“双循环近似”模型，继续深入GPU核心架构的细节。

006：双循环近似 🚀

在本节课中，我们将要学习GPU核心架构的“双循环近似”模型。上一节我们介绍了单循环近似，它基于程序员的视角，通过多个线程束的切换来隐藏长延迟指令。本节中，我们将通过引入指令级依赖跟踪，来改进我们对底层GPU架构的理解，从而减少隐藏延迟所需的线程束数量，并允许从同一线程束中发射多条指令。

单循环近似的局限性

单循环近似模型的核心是线程束调度器根据程序计数器选择下一个要发射指令的线程束。然而，这个模型存在一个关键问题：它没有跟踪指令间的数据依赖关系。

这意味着，在单循环近似中，我们必须等待一个线程束的当前指令执行完毕，才能从该线程束发射下一条指令。因为调度器不知道下一条指令是否依赖于前一条指令的结果。

这导致了一个问题：为了隐藏长延迟指令（例如访存操作），我们必须在单个核心上维持大量活跃的线程束，以便在某个线程束等待时，调度器可以切换到其他线程束。虽然GPU本身是海量多线程的，但为每个线程束保存大量状态（如寄存器、程序计数器等）会消耗巨大的芯片面积和功耗。

双循环近似的核心思想

双循环近似旨在解决上述问题。其核心思想是：如果我们能跟踪指令间的依赖关系，就可以用更少的线程束来隐藏延迟，并允许从同一线程束中发射多条指令。我们只需要确保在发射指令前，其所有依赖都已满足。

为了实现这一点，GPU引入了指令缓冲区和独立的指令调度器。

  • 指令缓冲区：从指令缓存获取的指令会先存放在这里。
  • 指令调度器：负责检查指令缓冲区中的指令，判断哪些指令的依赖已解决，从而可以发射到流水线的后续阶段。

指令内存通常由一级指令缓存和更高级别的统一缓存支持。指令缓冲区有助于隐藏指令缓存未命中的延迟，可以看作是缓存层次结构中的又一层。

在缓存命中或未命中填充后，新的指令信息会被放入指令缓冲区。指令缓冲区的组织形式多样，但最直接的方法是为每个线程束预留一个或多个指令的存储空间。

依赖检测：记分牌

在CPU中，检测数据依赖主要有两种经典方法：记分牌和保留站。保留站用于消除名称依赖，但需要复杂的关联逻辑，在面积和能耗上代价高昂。这对于追求极致计算密度、设计倾向于简单化的GPU来说是个问题。

记分牌则支持有序或乱序执行。虽然乱序执行的记分牌可能很复杂，但用于单线程有序CPU的记分牌可以设计得非常简单。

在简单的有序记分牌中，每个寄存器由一个比特位表示。当一条将要写入该寄存器的指令发射时，对应的比特位被置位。后续指令在发射前，会检查其源操作数寄存器对应的记分牌比特位。如果比特位被置位（表示存在未完成的写操作），则指令必须等待，直到该比特位被清除（写操作完成）。这可以防止写后读和写后写冒险。结合有序发射，只要寄存器的读取是按序进行的，这种简单的记分牌也能防止读后写冒险。

由于其设计简单、面积和能耗低，这种方案对GPU很有吸引力。因此，GPU实现了有序记分牌。

GPU实现记分牌的挑战

然而，将CPU的简单记分牌直接移植到GPU上会面临挑战，根源在于GPU的海量多线程特性。

以下是两个主要挑战：

  1. 巨大的状态开销：由于海量多线程，每个线程束（32线程）可能拥有多达120个寄存器。如果一个核心支持多达64个线程束，那么实现记分牌就需要跟踪 64 warps * 120 registers = 7680 个比特位（约8Kb）。这本身就是一个不小的开销。
  2. 读端口爆炸：在简单的记分牌设计中，如果指令遇到依赖，它需要不断查询记分牌以检查依赖是否解除。在单线程设计中，这复杂度尚可。但在多线程有序处理器中，来自多个线程的指令可能都在等待更早的指令完成。如果所有这些指令都必须每个周期探测记分牌，将需要大量的读端口。例如，64个线程束，每个指令最多4个操作数，如果所有线程束每个周期都探测记分牌，将需要 64 warps * 4 operands = 256 个读端口。这个数字非常惊人，实现成本极高。

一种简单的规避方法是限制可参与调度的线程束数量。但这又与我们利用海量多线程来隐藏延迟的初衷相悖，形成了一个矛盾。

解决方案：基于每线程束的简化记分牌

2008年提出的一种设计可以同时解决状态开销和读端口问题。该设计的核心思想是：

不再为每个线程束的每个寄存器保留一个比特位，而是为每个线程姆维护一个小的条目表（研究表明约3-4个条目）。每个条目记录一条已发射但尚未执行完成的指令将要写入的寄存器标识符。

以下是该方案的工作流程：

  1. 指令入缓冲时检查依赖：当指令从指令缓存取出并放入指令缓冲区时，会将该线程姆记分牌中的条目与指令的源寄存器和目的寄存器进行比较。这会产生一个短的比特向量（例如3或4位），每一位对应记分牌中的一个条目。如果指令的某个操作数与记分牌中的某个条目匹配（即存在依赖），则对应的比特位被置位。这个比特向量会与指令一起编码并存放在指令缓冲区中。
  2. 调度条件：一条指令只有在它的整个依赖比特向量全为0（即所有依赖都已解除）时，才有资格被调度器考虑发射。这可以通过一个简单的或非门电路高效实现：只有当所有输入位都为0时，输出才为1，表示指令就绪。
  3. 依赖解除与更新：当一条执行完成的指令准备将结果写回寄存器文件时，它会清除在记分牌中为它分配的条目。同时，它还需要清除指令缓冲区中所有属于同一线程束、且依赖比特向量中对应位被置位的指令的相应依赖位。
  4. 条目耗尽处理：如果某个线程姆的记分牌条目已全部用完，那么要么暂停所有线程姆的指令获取，要么丢弃当前无法分配条目的指令，后续需要时重新获取。

双循环架构总结

在双循环近似架构中，两个“循环”或调度层次如下：

  • 第一循环（线程束调度/指令获取）：选择那些在指令缓冲区中有空间的线程姆，根据其程序计数器访问指令缓存，以获取下一条指令并放入指令缓冲区。
  • 第二循环（指令调度）：检查指令缓冲区中的所有指令，找出那些依赖已解除、有资格执行的指令。这些指令可以来自同一个线程姆。调度器从中选择指令发射到执行单元。

这种设计的关键优势在于，它解决了单循环近似中一个线程姆一次只能发射一条指令，从而需要大量线程姆的问题。现在，我们可以减少所需的线程姆数量，并能够从同一个线程姆中发射多条指令，只要它们之间的依赖得到满足。

课程总结

本节课中我们一起学习了GPU核心架构的“双循环近似”模型。我们首先回顾了单循环近似的局限性，即无法跟踪指令依赖，导致需要大量线程姆来隐藏延迟。接着，我们引入了通过指令缓冲区和记分牌实现依赖跟踪的双循环模型。我们探讨了将CPU简单记分牌移植到GPU时，因海量多线程特性而面临的巨大状态开销和读端口挑战。最后，我们详细介绍了一种基于每线程姆简化条目表的记分牌设计方案，它有效地解决了这些问题，允许更高效的指令级并行和线程束管理。

在下一节中，我们将继续探讨“三循环近似”，并深入了解当存在多条就绪指令时，GPU如何高效地访问寄存器文件等有趣主题。

007：SIMT核心 - 第四部分

在本节课中，我们将学习GPU调度模型中的“三循环近似”的第三个循环，即寄存器文件访问调度。我们将探讨为何简单的寄存器组设计会导致严重的组冲突，并介绍一种更高效的解决方案——操作数收集器。

概述

在前两节中，我们介绍了GPU调度模型的前两个循环：第一个循环从线程束中选择指令，第二个循环从指令缓冲区中选择已就绪的指令。本节我们将进入第三个循环，探讨如何调度对寄存器文件的访问，以解决因大量线程束导致的寄存器文件端口需求和组冲突问题。

简单的组化寄存器文件及其问题

为了隐藏长内存延迟，GPU需要维持大量活跃的线程束。这导致了巨大的寄存器文件需求，例如在Kepler、Maxwell和Pascal架构中，寄存器文件大小约为256KB。

一种简单实现寄存器文件的方法是：为每个周期、每条发射指令的每个操作数提供一个端口。但这会导致端口数量过多，不切实际。

因此，我们采用组化设计，使用多个单端口存储体来模拟大量端口。然而，简单的组化设计会引发组冲突问题。

以下是简单的组化寄存器文件微架构示意图：

指令解码 -> 寄存器号 -> 仲裁器 -> 寄存器组 -> 流水线寄存器 -> 执行单元

在简单的组化布局中，我们使用取模运算来映射寄存器到组。例如，对于4个组，映射规则为：寄存器号 % 4。

  • Warp 0: R0 -> Bank 0, R1 -> Bank 1, R2 -> Bank 2, R3 -> Bank 3, R4 -> Bank 0, ...
  • Warp 1: R0 -> Bank 0, R1 -> Bank 1, R2 -> Bank 2, R3 -> Bank 3, R4 -> Bank 0, ...

这种设计的局限性在于，它容易在同一线程束内和不同线程束间产生组冲突，导致寄存器读取停滞，利用率低下。

操作数收集器的引入

为了解决上述问题，现代GPU设计很可能采用了“操作数收集器”的概念。

我们用收集器单元取代了简单的流水线寄存器。当指令进入寄存器读取阶段时，会被分配一个收集器单元。多个指令可以同时在不同的收集器单元中收集它们的源操作数。

这种设计的关键优势在于：

  • 提高吞吐量：即使某些指令因组冲突而延迟，其他指令的操作数读取可以继续进行，从而提高了整体吞吐量。GPU是面向吞吐量的机器，短暂的延迟是可接受的。
  • 增加组级并行性：仲裁器现在可以从多个指令的众多源操作数中进行选择，更有可能找到可以并行访问不同组的操作数，从而更充分地利用所有存储体。
  • 容忍组冲突：操作数收集器通过调度读取操作来容忍偶尔发生的组冲突，这是第三个调度循环的核心。

优化的寄存器布局：交错组化

除了引入收集器，我们还可以优化寄存器到组的映射方式，以进一步减少冲突。

简单的取模布局（寄存器号 % 4）在线程束间容易产生冲突。当多个线程束执行相同或相近的代码时，它们会访问逻辑上相同的寄存器号，这些寄存器会被映射到同一个物理组，导致严重的线程束间组冲突。

解决方案是采用“交错组化”布局。其核心思想是为不同线程束的寄存器映射引入一个偏移量。

  • Warp 0: R0 -> Bank 0, R1 -> Bank 1, R2 -> Bank 2, R3 -> Bank 3 (标准取模)
  • Warp 1: R0 -> Bank 1, R1 -> Bank 2, R2 -> Bank 3, R3 -> Bank 0 (偏移+1)
  • Warp 2: R0 -> Bank 2, R1 -> Bank 3, R2 -> Bank 0, R3 -> Bank 1 (偏移+2)

这样，即使所有线程束都在访问R0，它们也会被分散到不同的组（Bank 0, 1, 2），从而实现了线程束间的组级并行，显著缓解了冲突。

潜在问题与解决方案：写后读冒险

使用操作数收集器时，一个潜在问题是写后读（Read After Write, RAW）冒险。由于指令的操作数读取可能因组冲突而延迟，一条较晚的指令可能比一条较早的指令更早收集齐操作数并准备执行。如果允许这种情况发生，会破坏程序的正确性。

研究提出了几种解决方案来控制这种冒险：

  1. 提交时释放（Release on Commit）：每个线程束最多只能有一条指令处于执行状态。这保证了指令顺序，但会严重降低性能（在某些情况下几乎减半）。
  2. 读取时释放（Release on Read）：每个线程束最多只能有一条指令在操作数收集器中收集操作数。这释放了收集器资源，性能影响较小（在所研究负载中小于10%）。
  3. 布隆过滤器（Bloom Filter）：使用一个小型的布隆过滤器来跟踪未完成的寄存器读取。这是一种更精细的跟踪机制，性能开销最小（小于几个百分点）。

此外，实践中的解决方案可能更复杂。例如，NVIDIA的Maxwell架构引入了读依赖屏障，这很可能通过特殊的控制指令（在SASS代码中看到的元数据指令）来管理依赖关系，避免特定指令的写后读冒险。

总结

本节课我们一起学习了GPU SIMT核心中调度模型的第三个循环——寄存器文件访问调度。

我们首先分析了简单组化寄存器文件设计的问题，即严重的组冲突导致资源利用率低下。接着，我们引入了操作数收集器的概念，它通过允许来自多个指令的操作数读取重叠，提高了吞吐量和组级并行性。然后，我们介绍了交错组化寄存器布局，通过为不同线程束引入偏移映射，有效减少了线程束间的组冲突。最后，我们探讨了使用操作数收集器时可能出现的写后读冒险问题，并概述了几种控制该冒险的解决方案。

通过这三个调度循环（线程束调度、指令就绪调度、寄存器访问调度），GPU能够高效地管理成千上万个线程，最大化硬件利用率，从而实现其强大的并行计算能力。

008：指令重放 🚀

在本节课中，我们将学习GPU如何处理流水线中的结构冒险。我们将探讨为何传统的“停顿”方法在GPU中并不理想，并介绍GPU采用的“指令重放”策略。

上一节我们讨论了寄存器读取调度，本节中我们来看看当指令在流水线中遇到资源冲突时，GPU如何应对。

结构冒险的来源

GPU流水线中存在多种导致结构冒险的原因。例如，在寄存器读取阶段，可能会耗尽运算单元或收集器单元。此外，内存系统也是结构冒险的常见来源。虽然我们尚未深入探讨内存系统，但可以理解，一个由线程束执行的内存指令可能需要分解为多个独立操作，每个操作在特定周期都可能占用流水线的一部分资源。

传统CPU的解决方案：停顿

在标准CPU流水线中，处理结构冒险的常见方法是停顿后续指令。具体做法是，当一条指令遇到资源冲突时，暂停所有更年轻的指令，直到冲突条件解除，该指令能够继续执行为止。

然而，这种方法在GPU中存在两个主要问题：

  1. 关键路径影响：GPU拥有大型寄存器文件和众多支持完整图形流水线的阶段。分发停顿信号可能会影响关键路径，进而可能影响图形处理等具有服务质量期限的任务的帧率。
  2. 吞吐量损失：停顿一个线程束的指令，可能导致其他不依赖该冲突资源的线程束指令也被迫停顿，从而造成不必要的吞吐量下降。

GPU的解决方案：指令重放

为了避免停顿带来的流水线阻塞、电路面积增加或时序开销，GPU采用了“指令重放”策略。这个概念在CPU中同样存在，通常用作推测执行的恢复机制。例如，CPU可能推测性地调度一条依赖于具有可变延迟的加载指令，如果加载未命中缓存，就需要重放该依赖指令。

但在GPU中，我们通常没有推测执行。GPU依赖大规模多线程来隐藏延迟。因此，指令重放在GPU中的主要目的是避免资源冲突导致的效率低下。

以下是GPU实现指令重放的一种可能方式：
GPU可以将指令保留在指令缓冲区中，直到确认该指令已完全执行完毕，或者其所有独立部分都已执行完成。这样，当指令因结构冒险而无法继续时，无需停顿整个流水线，只需在资源可用时，从缓冲区中重新取出该指令并再次尝试发射即可。

本节课中我们一起学习了GPU处理结构冒险的策略。我们了解到，与CPU采用停顿机制不同，GPU更倾向于使用指令重放来避免流水线阻塞和吞吐量损失。其核心思想是将遇到资源冲突的指令暂时保存在缓冲区中，待资源可用时再重新发射，从而维持高吞吐量。

下一节，我们将探讨与SIMT核心相关的研究方向与现状。

009：Warp Compaction

在本节课中，我们将学习GPU架构中处理分支发散问题的一个重要研究方向：Warp Compaction。我们将探讨其核心思想、面临的问题以及学术界和工业界提出的多种解决方案。

概述

上一节我们介绍了GPU处理分支发散的基础硬件机制。本节中，我们来看看如何通过Warp Compaction技术来优化分支发散带来的性能问题。其核心思想是重新组织线程，以提高SIMD执行单元的利用率。

分支发散的问题

理想情况下，一个Warp中的所有线程应执行完全相同的控制流路径。但由于线程的自主性，它们可能走向不同的分支，这就产生了分支发散。现代GPU包含特殊硬件（如SIMT栈）来处理分支发散，但基线方案存在一些问题。

以下是SIMT栈基线方案的主要问题：

  1. 降低SIMD效率：当发生分支发散时，基线SIMT栈会串行执行每个分支目标。在最坏情况下，一个Warp中可能只有单个线程活跃，导致SIMD硬件利用率极低。
  2. 不必要的串行化：这种串行化对于功能正确性并非必需。GPU编程模型并未在标量线程之间强加隐式数据依赖，它们需要通过共享内存和屏障进行显式通信。因此，GPU可以交错执行发散Warp的所有分支目标，以利用SIMT硬件的空闲周期。
  3. 不充分的MIMD抽象：基线SIMT栈强制发散的Warp在编译器定义的重汇聚点重新汇聚，这相当于在每个重汇聚点隐式地强加了一次Warp范围内的同步。虽然这对许多现有应用有效，但这种隐式同步可能带来问题，例如在细粒度锁场景下导致死锁。
  4. 面积开销：每个Warp的SIMT栈需要存储空间（如32x64位）。在典型GPU应用中，分支发散并不常见，这些存储空间本可用于提升应用吞吐率，例如增大缓存或增加ALU单元。

Warp Compaction 的核心思想

GPU通过细粒度多线程来容忍长延迟内存访问。这意味着大量线程同时运行以掩盖延迟。由于这些Warp通常运行相同的计算内核，它们很可能遵循相同的执行路径，并在大致相同的时间点遇到分支发散。

因此，即使单个Warp内部发生发散，跨Warp观察，每个分支目标仍可能由大量线程执行。这些线程只是分散在多个静态Warp中，每个Warp单独处理自身的发散。

Warp Compaction 的关键洞察在于：我们可以重新排列这些来自不同静态Warp、但执行相同指令的线程，将它们打包成新的动态Warp，从而减少活跃线程的分散程度，提高整体SIMD效率。

  • 静态Warp：由GPU硬件在启动内核时，根据连续线程ID形成的Warp。这是程序员直观理解的Warp组织形式。
  • 动态Warp：在运行时通过重新组织线程（例如，将不同静态Warp中执行相同路径的线程组合在一起）形成的Warp。

硬件Warp Compaction技术

以下是一系列利用上述观察来提升遭受分支发散的应用性能的硬件研究。

动态Warp形成 (2007-2009)

这是早期尝试，核心思想是重新排列那些执行相同指令但分散在不同Warp中的线程，形成新的动态Warp，以消除发散。

后续研究（2011年）指出了动态Warp形成的两个主要性能病理：

  1. 贪婪调度策略可能导致线程饥饿，从而降低SIMD效率。
  2. 线程重组会增加非合并内存访问和共享内存体冲突。程序员通常按照线程连续访问内存的假设来优化，而动态重组打乱了这种连续性，可能导致访存模式恶化。

线程块压缩 (2011)

TPC技术建立在动态Warp形成的基础上。它观察到持续不断地重组线程并不能带来额外收益，重组只需要在发散分支之后（发散区域开始）和重汇聚点之前（一致区域开始）进行。

TPC扩展了原有的每Warp SIMT栈，使其管理同一核心内同一线程块的所有Warp，强制它们在发散分支和重汇聚点进行同步和压缩。这相当于将SIMT栈的管理范围从单个Warp扩大到了一个线程块。

然而，在每个发散分支同步所有Warp进行压缩会大大减少可用的线程级并行度。因此，TPC采取折中方案，将压缩限制在单个线程块内。由于GPU应用通常在单个核心上并发执行多个线程块，当一个线程块的Warp为压缩而同步时，其他线程块可以继续执行，从而重叠压缩开销。

大Warp微架构 (2011)

LWM架构也扩展了SIMT栈来管理一组Warp的重汇聚。但它要求组内的Warp完全锁步执行，以便能在每条指令处进行压缩。这比TPC更进一步减少了TLP。

为了补偿锁步执行带来的TLP损失，LWM采用了更复杂的记分牌微架构，以线程粒度跟踪寄存器依赖，允许组内某些Warp略微领先于其他Warp执行。

压缩充分性预测器 (2012-2013)

该工作扩展了线程块压缩，引入了一个预测器。其目标是智能地判断在某个分支点进行线程压缩是否能带来收益，并只在这些预测有益的分支点同步线程。这可以避免因无益的压缩停顿而损失的TLP。研究表明，一个简单的、类似于单级分支预测器的历史预测器就能达到足够的准确度。

内部Warp压缩 (2013)

该技术针对在窄硬件单元上执行宽SIMD执行组的场景。它将一个执行组划分为多个与硬件宽度匹配的子组。一个遭受发散的SIMD执行组可以通过跳过完全空闲的子组来在窄硬件上更快地运行。

为了创建这些完全空闲的子组，需要增加一个混洗机制，在发散点将元素打包到更少的子组中。

同时Warp交错 (2012)

SWI扩展了GPU SIMT前端，使其每个周期能发射两条不同的指令。它通过将Warp大小加倍来补偿增加的复杂度。SWI将一个遭受发散的Warp的指令与另一个发散Warp的指令协同发射，以填补分支发散留下的空缺。

寄存器文件设计的挑战

硬件压缩提案通常在SM核心本地进行，以避免引入额外的核心间通信流量。由于被压缩的线程都位于同一核心并共享同一寄存器文件，因此可以通过更灵活的寄存器文件设计来执行压缩，而无需移动其架构状态。

典型的寄存器文件使用大容量单端口SRAM体来最大化面积效率。同一Warp的线程寄存器连续存储在同一SRAM体中，以便能一起访问。但Warp压缩创建的动态Warp不遵守这种连续排列，并且可能包含来自不同Warp的线程，这些线程本应存储在不同体中。

  • 2007年的工作提出了更灵活的寄存器文件设计，采用更多、端口更窄的体来维持相同带宽，以支持动态Warp的寄存器访问。
  • 2010年关于动态微内核的工作则使用每核心暂存内存作为中转区，通过迁移线程及其架构状态来实现压缩。

软件Warp Compaction技术

以下研究在软件层面实现Warp压缩，无需修改硬件寄存器文件设计，但可能引入额外的内存流量来在线程间迁移数据。

  • 早期流计算 (2000)：在GPGPU时代之前，就将此概念应用于流处理器。它将一个可能发散的核函数在发散分支处根据每个数据元素的分支结果，将其数据流拆分为多个流，分别由不同的核函数处理，最后在控制流发散结束时合并。
  • 并行前缀和实现流压缩 (2009)：使用并行前缀和算法来识别具有相同任务的元素流，并将其压缩为紧凑的子流。该实现利用了GPU片上暂存内存的访问灵活性。
  • 延迟着色技术 (2009)：用于光线追踪，使用流压缩来提高具有多种材质类别的复杂场景中的像素着色SIMD效率。它将击中相似材质类别的光线分组处理。
  • 运行时线程重映射系统 (2010)：提出一个运行时系统，动态地将线程重新映射到不同的Warp中，以改善SIMD效率和内存访问的空间局部性。该系统采用流水线设计，CPU负责动态重映射，GPU负责在重映射后的数据和线程上执行计算。
  • 集体上下文收集 (2015)：一种编译器技术，专注于优化那些每个线程在每个步骤执行不规则计算量的核函数（例如，在不规则图上进行广度优先搜索）。它通过让每个线程处理多个节点，并使用存储在共享内存中的Warp特定栈来卸载任务上下文，实现工作负载平衡。

线程分配的影响

在基线GPU架构中，具有连续线程ID的线程被静态融合形成Warp。这种默认的顺序映射对于大多数工作负载（尤其是规则应用）效果很好，因为相邻线程倾向于访问相邻数据。

但对于不规则应用，情况并非如此。2013年的一项工作指出，顺序映射对于Warp压缩技术是次优的。一个关键限制是：当线程被分配到新Warp时，它们不能被分配到不同的通道，否则寄存器文件状态将不得不移动到向量寄存器的不同通道。

该工作观察到程序结构会使某些控制流路径偏向特定的SIMD通道，这种偏向使得压缩变得困难，因为选择相同路径的线程往往位于同一通道，从而无法合并。因此，他们提出了几种不同的线程映射排列方式，以消除这些程序性偏向，提高压缩率。

另一项2013年（引用了2009年工作）的研究利用了SIMD数据路径宽度并不总是等于Warp宽度的事实（例如，SIMD宽度为16，Warp大小为32）。它观察到，当发生发散时，如果一个连续的SIMD宽度线程组被屏蔽，则该指令可以仅在一个周期内发射，跳过被屏蔽的通道。他们称之为周期压缩。然而，如果被屏蔽的线程不连续，则基本技术不会带来任何性能提升，因此需要类似前面提到的混洗压缩技术。

Warp标量化

2014年的一项工作认为，当Warp内的线程操作相同数据时，SIMT编程模型是低效的。为此提出了多种解决方案，包括在流水线中增加一个标量单元，用于处理编译器或程序员可以预先识别为标量的工作。AMD的GCN架构就包含了用于此目的的标量流水线。最新的NVIDIA架构（如Turing）也引入了标量流水线。

总结

本节课我们一起学习了Warp Compaction这一重要研究方向。我们首先回顾了分支发散带来的问题，然后深入探讨了其核心思想：通过硬件或软件手段，重新组织执行相同路径的线程，以提高SIMD效率。我们介绍了从早期动态Warp形成到后来的线程块压缩、预测器、内部压缩等多种硬件方案，也了解了软件层面的实现策略及其对寄存器文件设计带来的挑战。最后，我们还探讨了线程分配策略和Warp标量化等相关优化。这些研究大多集中在GPU架构的早期发展阶段，如今已成为该领域成熟的技术体系。

下一节，我们将开始探讨第二个重要主题：Intra-Warp Divergent Path Management。

posted @ 2026-03-29 09:14  绝不原创的飞龙  阅读(6)  评论(0)    收藏  举报