“你买的 4090 多少钱？”、“H100 性能真厉害！” ，GPU 的价格性能一直是大家乐于谈论的话题，作者也经常可以在茶余饭后听到这样的讨论。在热火朝天地谈论性能指标、价格以外，本着”知其然也要知其所以然“的道理，作者学习整理了 GPU 本身的工作原理，编程模型，架构设计，在这里将我的学习笔记整理成文与大家分享，希望与大家一起 “知其然也要知其所以然”。
1. 引言
Why GPU?
为什么要使用 GPU？很多同学的第一反应就是“快”，这当然没错。而一个更严谨的说法是，GPU 兼顾了“通用性”与“高效性”，才使得其一步步成为高性能计算的首选。
针对计算性能，1974 年 Dennard 等人提出了 Dennard 缩放比例定律（Dennard Scaling）。
Dennard 缩放比例定律 (Dennard Scaling) ：当晶体管特征尺寸缩小时，其功率密度保持恒定。具体表现为电压随特征尺寸线性下降，电流密度保持稳定，使得单位面积的功耗与晶体管尺寸成比例关系。
一言以蔽之：晶体管越小越省电。
推导到芯片设计领域：晶体管缩小，芯片能塞入的晶体管更多，同时保持整体能耗稳定，推动计算机性能持续提升。 在计算机发展的前四十年间，基于 Dennard 定律的晶体管微缩是提升性能的主要路径。但在 2005-2007 年间，随着晶体管进入纳米尺度，量子隧穿效应引发的漏电流呈指数增长，阈值电压难以继续降低，最终导致该定律失效。此时，工艺微缩带来的性能增益已无法抵消功耗的快速增长，著名的"功耗墙"问题开始显现。
单纯依靠缩小晶体管尺寸来提升性能的方法不再可行，部分工程师开始转向专用硬件，即专门为了某种或某几种计算设计的计算硬件，例如 Google 的 TPU（Tensor Processing Unit，张量处理器），就是一款专为加速机器学习任务而设计的专用硬件。然而，专用计算硬件只能聚焦于某一类或者某几类特定的计算任务，在处理其他任务时则可能力不从心。
而 GPU 则是向通用性演进的典型代表。虽然其最初设计目标是为图形渲染加速，但高度并行的 SIMT（单指令多线程）架构意外契合了通用计算的演进需求，无论是基于 CUDA 的深度学习训练，还是通过 OpenCL 加速的流体仿真，都能通过高度并行获得远超 CPU 的计算性能。
GPU 的“快”
为什么快？
高计算并发：与 CPU 相比，GPU 将更大比例的芯片面积分配给流处理器（如 NVIDIA 的 CUDA 核心），相应地减少控制逻辑（control logic）所占的面积，从而在高并行负载下获得更高的单位面积性能。
低内存延迟：内存访问导致的延迟也是影响性能的一大因素，GPU 通过在其每个核心上运行大量线程的方式，来应对并掩盖因全局内存访问导致的延迟。这种设计使得 GPU 即使在面临较慢的内存访问时，也能维持高效的计算性能。具体来说，每个 SIMT 核心同时管理多组线程（多个 warp，一个 warp 32 个线程），当某个 warp 因为等待内存数据而暂停时，GPU 可以迅速切换到另一个 warp 继续执行。这种快速切换使得 GPU 能够在等待内存数据返回的同时，保持高利用率，从而有效地“隐藏”了内存访问延迟。
特化内存与计算架构：GPU 通常配备高带宽的显存（如 GDDR6 或 HBM），能够快速读取和写入数据。如 NVIDIA A100 使用 HBM2e 显存最高可达到 1.6TB/s 带宽，是普通 DDR5 内存（51.2GB/s）的 31 倍。计算架构方面，GPU 集成专用计算单元实现硬件级加速，例如，NVIDIA 的 Tensor 核心针对结构化稀疏计算做专门设计，在低精度损失的情况下，可以极大得提升计算性能。
有多快？
理论算力计算：GPU 算力常以 FLOPS（Floating-Point Operations Per Second，每秒浮点运算次数）来表示，通常数量级为 T(万亿)，也即是大家听到的 TFLOPS。最常见的计算方式为 CUDA 核心计算法。
# CUDA核心计算法
算力（FLOPS）= CUDA核心数 × 加速频率 × 每核心单个周期浮点计算系数
# 以A100为例
A100的算力（FP32单精度）= 6912（6912个CUDA核心） × 1.41（1.41GHz频率） × 2（单周期2个浮点计算） = 19491.84 GFLOPS ≈ 19.5 TFLOPS实测性能评估：通过计算只能得到纸面上的理论算力，如果同学们手上真的有 GPU，那么实测性能评估则可以直接让你获取你的 GPU 的性能。此处为大家提供几种最常见的实测方式和思路。
首先推荐一个非常实用的工具 GPU-Z，它是一款免费工具，可提供计算机中显卡的详细参数信息，支持实时监控 GPU 负载、温度、显存使用情况等关键数据，是排查显卡性能问题或计算故障的实用诊断工具。GPU-Z 是监控工具，而 3DMark 则是最流行的性能测试工具，通过模拟高负载游戏场景评估电脑图形处理能力（在 steam 平台即可购买，电脑上有 GPU 的同学不妨买来跑个分试试）。
最后再介绍一下 GEMM（General Matrix Multiplication，通用矩阵乘法），这是一种经典的并行计算领域的计算密集型应用，与跑分工具这样的封装好的峰值性能测试工具相比，GEMM 的重点反而不是进行性能测试，而是不断调整优化逼近理论峰值的过程。GEMM 通过执行时间 T 和总操作数（M×K 与 K×N 的两矩阵相乘）计算实测算力：
算力 = 总操作数 / 执行时间 = A(M, K) × B(K, N)/ T = 2 × M × N × K / T 如果实测算力低于 GPU 理论峰值算力，则表明可能存在低效内存访问、计算资源利用率低、未充分利用硬件加速单元等问题，这些问题均可通过逐步优化来解决，以逼近理论峰值，当然也有温度/功耗问题和显存带宽瓶颈等硬问题，但影响较小。对实际操作进行 GPU 编程有兴趣的同学可以选择深入了解 GEMM，学习实现的比较好的 GEMM 库是如何优化以逼近理论峰值的，在这个过程中深入理解 GPU 计算和编程。
GPU 架构概述
在这里作者要做一个简单的说明，现代的 GPU 架构，先不论不同厂家，仅 NVIDIA 一家就有数十年的架构迭代史，其中涉及的各种优化改进，限于篇幅，本文不可能一一介绍。但是，要想完整了解整个 GPU 架构的发展，作者认为可以分两步走：以 NVIDIA 为例，就是“从 0 到 Fermi“，和”从 Fermi 到 Blackwell“。Fermi 架构是现代通用 GPU 架构的基石，其中许多核心设计思想传承至今，而此后直到作者撰文的 2025 年最新的 Blackwell 架构，都可以看做在基础上的一路迭代。本文介绍的重点为两步走里的第一步，即讲解现代通用 GPU 中的基石级的通用技术与设计，读者迈好第一步，就可以以此为基础广泛探索。
第一张图为 Fermi 架构图（来自Fermi 架构白皮书），完整的 Fermi 架构 GPU 由 4 个 GPC 组成（黄色框），每个 GPC 有 4 个流式多处理器 SM (Streaming Multiprocessor, 红色框），每个 SM 又有 32 个 CUDA Core，此外还有 L1、L2 Cache、共享内存、显存等组件。而每个 SM、每个 CUDA Core 的结构则可见第二张图。这样看还是过于复杂，为了更清晰的从原理上了解通用 GPU 机构，本文将根据以下的简化通用 GPU 架构图讲解，介绍 GPU 架构使用的术语也将倾向于学术界常见的通用术语：
SIMT 核心（SIMT Core）是 GPU 的核心计算单元，类似于 CPU 的多核集群，负责协调和管理大量线程的并行执行，对应 NVIDIA 架构中的 SM。SIMT（Single Instruction, Multiple Threads，单指令多线程），是 GPU 的核心执行模型，其本质是通过统一指令指挥多个线程并行处理不同数据。后文将做单独讲解。多个 SIMT 核心组成 SIMT Core Cluster，对应 NVIDIA 的 GPC，每个 Cluster/GPC 可以看做是一个可完整运作的 mini GPU，而实际的 GPU 由多个 GPC 组成，也就是大家常说的“多核”。
在同一个 SIMT 核心内运行的线程可以通过共享内存（Shared Memory）来进行彼此通信同步，SIMT 核心内还包含一级指令和数据缓存，用来减少与低级内存的交互次数从而提高性能。而 SIMT Core Cluster 之间通过 Interconnection Network 通信。
除 SIMT 核心外，另一重要部分是内存和内存管理，在图中即简化为 Memory Partition 和 GDDR 部分。Memory Partition 部分管理显存的访问，跨 SM 的 L2 全局一致性缓存也位于此处。GDDR，即为大家常常提到的显存，其是位于 GPU 芯片外部的专用内存，用于存储图形数据等，相比于 CPU 的普通内存通常针对访问延迟和带宽进行优化。
2. GPU 编程
本章将介绍如何编写程序使用 GPU 完成非图形类的计算，介绍重点在于揭示 GPU 的通用编程模式，以及程序执行的流程，并非专门的 GPU 编程教学。
程序如何执行？以 SAXPY 为例
SAXPY，即将向量 X 的元素乘以 A，再加上向量 Y。以下是用 C 语言实现的 CPU 计算 SAXPY 的代码：
// SAXPY 函数实现
void saxpy(int n, float a, float *x, float *y) {
for (int i = 0; i < n; i++) {
y[i] = a * x[i] + y[i];
}
}
int main() {
float a = 2.0;
int n; // 向量长度
float *x; // 向量 x
float *y; // 向量 y
// 此处省略内存分配、元素赋值、长度指定
// ...
// 调用 SAXPY 函数
saxpy(n, a, x, y);
return 0;
} 针对上述 CPU 计算代码，将代码改写为使用 CUDA 编写的在 GPU 上运行 SAXPY：
**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];
}
}
int main() {
float a = 2.0;
int n; // 向量长度
float *hx; // host 向量 x
float *hy; // host 向量 y
// 此处省略内存分配、元素赋值、长度指定
// GPU 内存分配
int vector_size = n * sizeof(float); // 向量数据大小
float *dx; // device 向量 x
float *dy; // device 向量 y
cudaMalloc(&dx, vector_size);
cudaMalloc(&dy, vector_size);
// 将 host 向量内容拷贝到 device 向量
cudaMemcpy(dx, hx, vector_size, cudaMemcpyHostToDevice);
cudaMemcpy(dy, hy, vector_size, cudaMemcpyHostToDevice);
// 执行 saxpy
int t = 256; // 每个 thread block 的线程数
int blocks_num = (n + t - 1) / t; // thread block 数量
saxpy<<<blocks_num, t>>>(n, a, dx, dy);
// 将 device 向量 y 内容（计算结果）拷贝到 host 向量 y
cudaMemcpy(hy, dy, vector_size, cudaMemcpyDeviceToHost);
// ... (剩余逻辑)
return 0;
} 设备侧与主机侧
GPU 编程的思维是将 GPU 当作 CPU 的协同外设使用，通常 GPU 自身无法独立运行，需要 CPU 指定任务，分配数据，驱动运行。第一行的__global__关键字，表示这段函数是内核函数（kernel，注意与 Linux 内核无关），是交给 GPU 执行的，而 main 函数则无此标识，由 CPU 执行。通常，将交 GPU 执行的代码部分称为设备(device)代码，而交给 CPU 执行的代码部分称为主机(host)代码。host 与 device 是 CUDA 编程惯用的风格，CPU 称为 host 侧，而 GPU 称为 device 侧。
main 函数中的 cudaMalloc、cudaMemcpy，是 CPU 操作 GPU 内存的操作，在分离式 GPU 架构（也就是独显）中，CPU 分配内存用于 GPU 计算，再将数据传输到分配的内存空间，然后在 GPU 上启动内核函数。GPU 执行的内核函数只能从分配的 GPU 内存空间读取数据。代码中的 host 向量对应 CPU 内存的数据，而 device 向量则代表 GPU 内存的数据。
值得一提的是，近年来统一内存（unified memory）在 GPU 的应用中逐渐流行，统一内存是指一种允许 CPU 和 GPU 共享同一段地址空间的内存架构，这种架构下可以实现 CPU 和 GPU 之间数据交换的自动化，开发者不需要手动管理数据在 CPU 到 GPU 之间的传输。
线程组织
完成内存分配和数据拷贝后，CPU 触发 GPU 执行 saxpy 内核函数。触发时同时指定了执行内核函数的线程的组织形式。在 CUDA 编程中，线程以 thread，thread block，grid 的层级结构进行组织，如上图所示：
● 线程（thread，绿色部分）：最基本的执行单元。线程包含独立寄存器状态和独立程序计数器。
● 线程块（thread block，黄色部分）：由多个线程组成的集合，支持一维、二维或三维结构。线程块内的线程可以通过共享内存进行通信，线程块之间无法通过共享内存通信，但可通过全局内存进行数据交互。
● Warp（蓝色线框）：硬件底层概念，GPU 实际运行时将 32 个线程组成一个 warp，同一 warp 内的线程同步执行相同的指令。
● 线程块与 warp 的关系：warp 是底层概念，NVIDIA 的 warp 固定包含 32 个线程，warp 是线程硬件调度的最小粒度。线程块是软件概念，线程块有多少个线程组成由代码指定。在运行时，硬件会将线程块中的线程 32 个为一组打包成多个 warp 进行调度，因此，线程块里的线程数最好为 32 的整数倍，以避免为拼凑完整 warp 而自动分配无效线程造成资源浪费。
● 网格（grid，总体）：网格是所有线程块的集合，支持一维、二维或三维结构，覆盖整个计算任务的运行范围。
thread，thread block，grid，warp 是 NVIDIA 的术语，而对于 AMD，四者又有其独特的称呼，因为本文使用的例子为 CUDA 编程，GPU 编程部分的讲解也将使用 NVIDIA 的术语体系，下表为术语对照表：
区别于 NVIDIA，AMD 的一个 wavefront 由 64 个 work item 组成。线程块有时也被称为 CTA（Co-operative Thread Array)。
代码执行 saxpy 部分：
// 执行saxpy
int t = 256; // 每个thread block的线程数
int blocks_num = (n + t - 1) / t; // thread block数量
saxpy<<<blocks_num, t>>>(n, a, dx, dy); 此处指定线程块为一维的，一个每个线程块（thread block）有 256 个线程（thread）。又计算得到了线程块的数量 block_num，指定网格（grid）也为一维，一个网格中有 block_num 个线程块。最后，用<<< >>>三个尖括号包含网格的线程块数、线程块的线程数，指定一个 grid 有 block_num 个线程块，一个线程块有 256 个线程。
线程块数量的计算
一个线程块由多少个线程组成可以指定，与此不同的是，线程块本身的数量则是由计算规模决定的，这段代码根据向量的长度计算了线程块的数量：
int blocks_num = (n + t - 1) / t; // thread block数量 这样计算的目的是保证线程数量足够，即每一个计算单元都有一个线程负责计算。
例如，如果向量长度 n=250，则 block_num = (250 + 256 - 1) / 256 = 1，每个线程块有 256 个线程，那么要保证每个向量元素有一个线程负责计算，1 个线程块就够了。又例如，如果向量长度 n=257，则 block_num = (257 + 256 - 1) / 256 = 2，需要两个线程块才能提供足够的线程，当然，本例子中的两个线程块足以提供 512 个线程，有很多线程实际上是闲置了。
总结上述计算方式，可以得到计算线程块数量时最常见的向上取整编程范式：
// B：线程块数，N：问题规模，T：线程块内线程数
B = (N + T - 1) / T 指定线程执行内核函数指令
最后，我们来关注 saxpy 内核函数本身，main 函数中分配的每个线程都会并发地执行这段代码：
__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];
}
} 此处为每个线程分配了一个其所属的向量元素，然后驱动线程分别完成计算。
首先计算 i，i 为线程的编号，blockIdx 是 block 在 grid 上的坐标，blockDim 则是 block 本身的尺寸，threadIdx 为 thread 在 block 上的坐标。此前提到我们的 grid、block 都是一维的，因此只需要取其 X 维度，因此 block 的编号就直接取 blockIdx.x，而一个 block 有 blockDim.x 个线程，线程编号为 threadIdx.x。
假设当前线程是第二个线程块上的第 10 个线程，即第 266 个线程，则其 index 应为 265：
i = blockIdx.x * blockDim.x + threadIdx.x = 1 * 256 + 9 = 265 得到线程编号 i 后，第 3 行判断 i 是否落在[0,n]区间内，n 为线程总数。如果为否，则该线程就是前面提到的多分配的闲置线程，不调度。而对于需要调度的线程，则根据自己的线程编号，读取源向量不同位置的元素，执行计算，并将结果写入结果向量的不同位置。这样，我们就为不同线程安排了独立的工作，让他们并发地完成工作。
多维线程组织结构
截止到这里我们提到的 grid、thread_block 都是一维的，实际可以支持一维、二维、三维，这里再举一个三维的例子：
// 主机端调用代码
void launch_kernel_3d() {
// 三维数据尺寸
int dimX = 64 int dimY = 32 int dimZ = 16;
// 定义三维线程块（Block）和网格（Grid）
dim3 blockSize(8, 4, 4); // 指定每个块包含8x4x4=128个线程
dim3 gridSize(
(dimX + blockSize.x - 1) / blockSize.x, // X方向块数
(dimY + blockSize.y - 1) / blockSize.y, // Y方向块数
(dimZ + blockSize.z - 1) / blockSize.z // Z方向块数
);
// 启动内核函数
kernel_3d<<<gridSize, blockSize>>>(d_data, dimX, dimY, dimZ);
} 使用 dim3（CUDA 数据结构）来承载三维 grid、thread block 的尺寸。grid 为三维，因此要计算 X、Y、Z 三个维度上 thread_block 的数量，仍套用前文提到的向上取整计算方法。而如果是三维的 grid、block，其计算线程编号时就需要取 X、Y、Z 三个维度：
// 核函数定义（处理三维数据）
**global** void kernel_3d(float* data, int dimX, int dimY, int dimZ) {
// 计算三维索引
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
if (x < dimX && y < dimY && z < dimZ) {
// 处理三维数据（例如：三维矩阵元素操作）
int idx = x + y * dimX + z * dimX * dimY; // 线程编号
data[idx] *= 2.0f; // 示例：每个元素翻倍
}
} SIMT
前文提到：SIMT（Single Instruction, Multiple Threads，单指令多线程），由 NVIDIA 提出，是现代通用 GPU 的核心执行模型，甚至可以说正是 SIMT 的出现，使得 GPU 从一种处理图形计算的专用硬件，进化为处理各类计算的通用处理器。SIMT 的本质是通过统一指令指挥多个线程并行处理不同数据，结合上述例子，此处展开讲解。
SIMT 本质上是一种并行计算的范式，要彻底理解 SIMT，以及 SIMT 存在的意义，就必须从另一种更基础的并行计算的范式——SIMD 讲起。因为 SIMT 是对 SIMD 进行“线程级抽象”得到的，或者说，SIMT 是“基于 Warp 的 SIMD”。
SIMD（Single Instruction Multiple Data，单指令多数据），即：在同一时刻向多个数据元素执行同样的一条指令。SIMD 范式常见的一种实现是 CPU 的向量化运算，将 N 份数据存储在向量寄存器里，执行一条指令，同时作用于向量寄存器里的每个数据。可见 SIMD，特别是向量化运算，是一种偏硬件底层的并行计算优化，而 SIMT 范式则是通过线程编程模型隐藏了底层 SIMD 的执行细节。
在向量化运算实现的 SIMD，有 N 个这样流程并发执行：“指令+操作数 → 结果”，而 SIMT 的设计思想，则将“指令+操作数”抽象成了“线程”，线程可以看做是打包了指令和操作数的一个执行单元：线程包含独立寄存器状态（操作数）和程序计数器（指令）。在软件编程时，程序以线程为单位进行调度，编程者只需要关注安排多少线程执行哪些指令，而无需过多考虑底层细节。这使得编程模型更接近多线程 CPU，降低开发者适配难度。
SAXPY 例子中的内核函数，就是以 SIMT 模型进行编程的，安排所有线程执行相同的指令，但每个线程执行指令时的指令操作数均不同，这便是 SIMT：
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
// 每个线程都执行这条指令，每个线程读取不同元素执行相同计算
y[i] = a * x[i] + y[i];
} 而在各线程实际运行时，硬件层面便会回归 SIMD 范式。继续以 SAXPY 为例，实际执行时 GPU 硬件会将其组织为 warp，warp 中的每个线程基于唯一索引 i，访问不同的内存位置，以不同的数据执行相同的指令，这便是 SIMD：
// 一个Warp中每个线程的执行流程（线程0-31）
//【指令 + 操作数 = 结果】的SIMD范式
i = 0 → y[0] = a*x[0] + y[0]
i = 1 → y[1] = a*x[1] + y[1]
...
i = 31 → y[31]= a*x[31]+ y[31] 传统的 SIMD 关注的是一条条指令本身的执行方式，而 SIMT 则将 SIMD“包了一层”，底层实现 SIMD，表面上提供线程级编程模型，让编程者很大程度上可以从串行的角度思考，而屏蔽了很多并行角度的执行细节。
这种编程便利最好的体现就是在出现分支（如 if-else）时：Warp 执行每个 Branch Path，执行某个 path 时，不在那个 path 上的线程闲置不执行，线程活跃状态通过一个 32 位的 bitmask 标记，分支收敛时再对齐汇总到下一段指令等等。后文将对这一过程作详细讲解，而在这里读者只需要理解到，如果只有底层 SIMD，那么这一切复杂流程都要编程者自己思考+编排，而在 SIMT 编程模型下编程者只需要编写分支代码，把这些编排交给硬件底层即可。
指令集与编译
刚才我们讲解了 CUDA C 语言编写的 SAXPY，到这里，只是到了高级语言层面，众所周知，高级语言需要转换为机器码才能被机器执行，本节将简单介绍 CUDA C/C++的程序的编译流程，以及 CUDA 的 PTX、SASS 指令集。
指令集：SASS、PTX
SASS（Streaming Assembly）是 GPU 的机器指令集，是实际在 GPU 上执行的指令。SASS 指令集直接对应 GPU 架构（Maxwell、Pascal 等），虽然不是严格的一一对应，但通常每个 GPU 架构有专属的 SASS 指令集，因此需要针对特定架构进行编译。
PTX（Parallel Thread Execution）是一种中间表示形式，位于高级 GPU 编程语言（如 CUDA C/C++）和低级机器指令集（SASS）之间。PTX 与 GPU 架构基本无耦合关系，它本质上是从 SASS 上抽象出来的一种更上层的软件编程模型，PTX 的存在保证了代码的可移植性（同一份 PTX 分发到不同架构上转为对应 SASS）与向后兼容性（可将 PTX 代码转为最新 GPU 架构对应的 SASS）。
PTX 是开发者可编程的最底层级，而 SASS 层则是完全闭源的，这也是 NVIDIA 的“护城河”之一。
编译流程
CUDA 程序的编译由 NVCC（NVIDIA CUDA Compiler）完成。
首先，NVCC 完成预处理；随后分类代码为设备代码和主机代码，NVCC 驱动传统的 C/C++编译器主机代码的编译和汇编；对于设备代码，NVCC 将其编译针对某架构的 SASS，编译过程中涉及 C --> PTX --> SASS 的转化，但通常不显式表现出来，生成的 PTX/SASS 码也会被直接嵌入最终的可执行文件。
运行期，GPU 会优先查找可执行文件中是否有适合当前架构的 SASS，如有则直接执行。若无，则 GPU 驱动（driver）会使用 JIT（Just-In-Time）编译手段，将 PTX 码编译为当前架构对应的 SASS 再执行（前提是可执行文件必须包含 PTX）。
3. SIMT 核心架构
前面两章，我们主要从总体概述和软件编程的角度了解了 GPU。相信不少同学在校园课程中，曾学习过 CPU 的核心架构，我一直以为，在了解了底层硬件是如何运作之后，我们看待处理器/硬件的视角才会有本质上的转变，从一个用户（这是执行我代码的黑盒）转变为一个专业技术人员（这是中央处理器）。因此，我们将更进一步，从更偏硬件的视角进一步了解 GPU 架构。
软硬分界线
前文提到 SIMT 核心也就是 NVIDIA 的 SM，也给出了来自 Fermi 白皮书的SM 结构图。但是，线程以 Warp 为单位在 SM 上执行，具体如何执行，执行的流程是什么，每个组件发挥什么作用，单单从结构体是看不出来的，因此我们需要引入 SM 的指令流水线结构图来进行讲解：
如图所示，SIMT 核心流水线从运行的处理阶段可以分为 SIMT 前端和 SIMD 后端两个部分：
- SIMT 前端：主要负责指令的获取、译码和发射、分支预测、以及线程的管理和调度。这部分设计的组件对应SM 结构图中的蓝色、橙色部分（Warp Scheduler、Register File）。
- SIMD 后端：主要负责完成计算。这部分设计的组件对应SM 结构图中的绿色部分（Core）。
SIMT 前端与 SIMD 后端的划分本质上是控制流与数据流的解耦，SIMT 前端关注指令流/控制流，而 SIMD 后端关注单个指令执行/数据流。
SIMT 前端在硬件运行时“落实”了程序对线程的调度：SIMT 前端以 warp 为单位调度线程，其包含的指令缓存（I-Cache）、解码器和程序计数器 PC 组件集中管理线程的指令流，并使用 SIMT 堆栈等技术实现线程间的条件分支独立控制流。SIMD 后端主要负责执行实际的计算任务。在 SIMT 前端确定了 warp 要执行的指令后，指令发射，SIMD 后端负责高效地完成一条条指令。具体的数据计算单元 ALU，以及存取计算数据的寄存器访问（Operand Collector）、寄存器文件（Register File）、内存读写（Memory）位于此处。
说到这里，这么多组件、组件之间有各种配合，不少同学估计已经要绕晕了。下面本文如果平铺直叙地直接深入一个个组件的细节，就会变得难以理解。因此，下面本文将采取一种“三步走”的讲解策略，先构建一个能执行计算任务的“最小系统”流水线，然后逐步向其中添加优化与功能，最终经过三步，构建出上图中完整的流水线架构。
第一步：最小可用系统
如上图，我们将 SIMT 内核的架构做了最大可能的简化，构成了一个“最简 GPU”。这个最小可用系统由 6 部分构成，此 6 个组件相互配合，使得我们的最简 GPU 可以做到最简的指令执行功能：即顺序执行每一条指令，一条指令执行完再执行下一条：
- Fetch：取指令
- Decode：指令解码
- SIMT Stack：SIMT 堆栈，管理线程束的分支执行状态，下文讲解
- Issue：指令发射
- ALU：算数逻辑单元，代表执行计算的组件
- MEM：存储器访问单元，代表对 L1 Cache、共享内存等各层级内存访问的管理。
其中 1、2、4、5、6 部分是在 CPU 上久而有之的“老面孔”了，本文不多做解释。本节将重点介绍 GPU 独有的“新面孔”：SIMT 堆栈。
分支发散：哪些线程执行哪条指令？
在 GPU 并行计算的发展历程中，SIMT 堆栈是早期架构解决线程分支管理问题的核心机制。
现实中的计算任务常包含大量条件分支（if-else、循环等）。在遇到条件分支发散（Branch Divergence）当线程束内线程选择不同执行路径时，会产生线程发散（Thread Divergence）：
如上图，起初有 5 个线程执行相同的指令，直到分支发散处，根据 SIMT 的特性：多线程执行相同指令，但每个线程有自己独立的数据，假设此处是一个 if-else，有不同数据的线程将得到不同的条件判断结果，2 个线程进入 if 分支，3 个线程进入 else 分支，进入不同分支的线程执行的指令流自然不同。
此处便出现了线程发散，即同一 warp 内的线程要执行不同指令，单由于线程以 warp 为最小单位调度，同一时钟周期内同一 warp 内的线程必须执行相同的指令，那么不同执行分支的线程就需要分开调度，例如一个时钟周期调度该 warp 执行 if 分支（if 分支的线程活跃），下个时钟周期再调度该 warp 执行 else 分支的线程（else 分支的线程活跃）。也就是说，以 warp 为单位调度不代表每次调度 warp，其中全部 32 个线程都活跃，也可以只有部分线程活跃，其余线程闲置。
分支发散带来的复杂性不仅是线程指令流的发散，还有调度顺序。如上图，if-else 分支发散后，分支聚合，5 个线程执行红色部分，但依赖 if 和 else 分支线程的运行结果，那么就要求蓝色部分和黄色部分先执行完，再执行红色部分。
为解决分支发散时的线程调度，NVIDIA 于 2008 年在 Tesla 架构中首次引入 SIMT 堆栈，并作为 2010 年 Fermi 架构的核心技术，其核心思想是：
- 路径跟踪：当线程束遇到分支时，通过堆栈记录所有可能执行路径的上下文（如程序计数器 PC、活跃线程掩码）。
- 串行化执行：依次调度 warp 中每个分支路径上的线程，其他线程暂时闲置。
- 重新收敛：在所有路径执行完毕后，恢复完整 warp 的并行执行。
SIMT 堆栈
为了介绍 SIMT 堆栈的工作原理，我们引入一个稍复杂一点的分支发散例子，如下图中的左图，是一个程序的分支流，其中有两层嵌套的 if-else。而下方右图则用表格的形式展示了左图程序执行过程中 SIMT 堆栈的情况：表格最下行为栈顶，三行分别为聚合点 PC、下条指令 PC 和活跃掩码（Active Mask）。
聚合点 PC，即分支聚合点的指令指针，例如，对于 B、F 这一分支发散，其聚合点 PC 就是 G。
下条指令 PC，顾名思义，就是当前指令的下一条指令的 PC，如 A 的下条指令 PC 为 B。
活跃掩码（Active Mask），代表了哪些线程执行这条指令，本例子中假设有 4 个线程，而活跃掩码就有 4 位，每一位分别对应一个线程，这一位为 0，则线程不执行这条指令，为 1 则执行，例如，指令 B 的活跃掩码为 1110，代表前三个线程执行 B，而第四个线程执行 else 分支的 F（因此 F 的活跃掩码为 0001）。
观察执行 A、B、C、D 时的 SIMT 堆栈，可以得到 SIMT 堆栈的运行方式：在遇到分支发散时，先将分支聚合点压入堆栈，随后压入各分支的指令，各分支指令执行完毕后，回到聚合点，执行聚合点的指令。
我们跟着例子走一遍：
- 执行指令 A，发现有分支发散。此时先将分支聚合点 G 压栈，再将两分支 F、B 先后压栈。
- 执行栈顶的 B，发现又有分支发散。此时先将聚合点 E 压栈，再将两分支 D、C 压栈。
- 执行栈顶的 C、D，回到聚合点 E。后续按弹栈顺序，再执行 F、G，完成执行。
通过以上调度策略，保证了存在依赖时的正确性，例如，如果执行 E 依赖执行 A、B、C、D 的执行结果，SIMT 栈刚好保证了 E 在 ABCD 后执行。
至于压入各分支时的压栈顺序，如压入 C、D 时的顺序，因为 C、D 二者之间不存在依赖关系，从正确性角度而言，CD 或者 DC 顺序都可以，此时通常从性能角度出发，优先压入有更少线程执行的指令（线程少离栈顶远，线程多离栈顶近），从而保证有更多线程执行的指令先弹栈执行，这样做有助于尽量减少栈的层数，提高性能。
SIMT 堆栈的问题
尽管 SIMT 堆栈在早期 GPU 架构中实现了分支管理能力，但其设计本质上面临多重硬件与效率瓶颈，难以适应现代计算任务（光线追踪、AI 训练推理等）对复杂控制流的需求：
- 传统方案依赖固定深度的硬件堆栈，每个线程束需独立维护堆栈，导致寄存器占用率攀升。
- 堆栈通常只有 4-8 级最大深度，这就意味着如果程序控制流过于复杂，例如，在训练 Transformer 模型时，自注意力机制可能触发数十/上百层条件判断，远超堆栈容量。
- 每次分支发散时，硬件需执行压栈，并在路径切换时弹栈。例如，一个包含 5 层嵌套 if-else 的着色器，需至少 10 次堆栈操作（进入和退出各一次）。随着程序变得复杂，此类操作越来越多，会造成显著的流水线延迟。
- 最后，由于堆栈的严格后进先出（LIFO）特性要求分支路径必须按嵌套顺序执行，很容易造成负载失衡甚至死锁。例如，在光线追踪中，部分线程可能因等待材质纹理读取而停滞，而其他线程已完成计算，但受限于堆栈顺序无法提前推进。
独立线程调度
在 Volta 之前的架构（如 Pascal、Fermi）中，在分支线程调度上，由 SIMT 完成调度，而 Warp 作为基本调度单元，所有线程共享统一的 PC 和活动掩码，当 Warp 内线程执行不同分支路径时，需按路径顺序串行执行。例如：线程 0-3 执行分支 A 的指令，线程 4-31 执行分支 B 的指令，则必须排队执行，一部分线程先执行分支 A 的指令，另一部分线程必须等待。
而从 Volta 架构开始，引入了独立线程调度（Independent Thread Scheduling）。每个线程拥有独立的程序计数器(PC)和执行状态寄存器，允许同一 Warp 内的线程在不同分支路径上并行执行指令流。但硬件层面仍以 Warp 为基本调度单元。
无堆栈分支收敛
同时，也是从 Volta 架构开始，随着独立线程调度的引入，传统 SIMT 堆栈被弃用，分支收敛机制也升级到了无堆栈分支重新收敛（Stackless Branch Reconvergence）机制，通过收敛屏障（Convergence Barriers）技术来低成本解决分支代码执行调度问题，独立线程调度为无堆栈分支重新收敛提供了硬件支持。
无堆栈收敛屏障机制的核心手段之一是屏障参与掩码（Barrier Participation Mask）与线程状态协同管理，其核心思想可以通过 ADD 和 WAIT 操作来展示：
- ADD（屏障初始化）：当 Warp 执行到分支发散处前，通过专用 ADD 指令，活跃线程将其标识位注册到指定收敛屏障的 32 位掩码中，标记参与该屏障的线程组。
- WAIT（屏障同步）：在预设的收敛点（如分支汇合处），硬件插入 WAIT 指令。到达此处的子线程组将线程状态标记为“阻塞”，并更新屏障状态寄存器。当所有参与线程均抵达屏障后，调度器才重新激活完整线程束。
为了便于理解，下面用一个图表示一个简单的的 ADD，WAIT 的例子：
另外，通过新增的 syncwarp()函数，开发者也可手动指定分支后的同步点，强制线程在特定位置重新收敛。
相比于 SIMT 堆栈，收敛屏障只需要使用仅需位掩码和状态寄存器，对于一个 Warp（32 个线程），一个屏障只需要 32bit（每个 bit 对应一个线程），操作成本和硬件资源占用均极低，且不会再有堆栈深度限制，可以支持任意深度的条件分支嵌套。这一设计使得现代 GPU（如 NVIDIA Volta+架构）在复杂控制流场景下仍能保持高吞吐量，成为实时光追、AI 推理等应用的关键支撑。
第二步：动态指令调度以提高并发
在第一步构建的最小可用系统中，采用的是“一条指令执行完再执行下一条”的最简执行策略。前文提到过，GPU 为了隐藏内存访问的延迟，需要在内存访问指令为执行完前，先分配 warp 去执行其他指令。这里的策略其实就是动态指令调度，根据指令依赖关系和执行单元可用性，动态决定指令发射顺序。
但此处有一个重要条件，就是先分配执行的这个其他指令，不能依赖于未完成指令的结果，否则无法执行。因此，需要先判断指令之间是否存在依赖关系，才能选择出不依赖未完成指令的指令进行执行。为了分析指令之间的依赖关系，以支持乱序执行，第二步为我们的系统增加了 I-Cache、I-Buffer 和 ScoreBoard 三个组件，并且 ALU 和 MEM 又多了一个指向 ScoreBoard 的“回写”操作。
I-Cache（指令缓存）、I-Buffer（指令缓冲区）：缓存从内存中读取的指令，和解码后的指令。此二者将一系列指令存放在一起，用于进行依赖分析，并在分析结束后快速读取指令进行乱序执行。I-Cache 和 I-Buffer 为指令依赖分析提供了数据，ScoreBoard（计分牌）则是实际执行依赖分析操作的组件。
GPU 计分板的核心目标是检测指令间的数据依赖关系（如 RAW、WAR、WAW），并控制指令发射顺序以避免冲突。数据依赖关系反映到硬件层面体现为对寄存器的读写依赖关系，因此，GPU 的计分板被设计为一个 bitmap，其记录了每一条未完成指令的目标寄存器，即如果这条指令要写寄存器 R1，则将 R1 对应的 bit 置为 1。在指令完成后，再将 R1 对应的 bit 写回 0。该流程如下图所示：
由于寄存器是线程私有的，需要为每个线程分配足够的寄存器，因此 SIMT 核心中的寄存器数量是很大的，即便做到一个寄存器只需要一个 bit 表示状态，ScoreBoard 也会变得过大。因此，实际设计中，每个 warp 维护一个自己的 ScoreBoard，由于每个 warp 同一时间只能执行同一条指令，一条指令能访问的寄存器也是有限的，因此每个 warp 的 ScoreBoard 有 3-4bit 即可，每一个 bit 称为一个表项（entry）。
在判断一条指令是否能执行时，将该指令的源/目标寄存器与其所属 warp 的计分板表项做比较（计算 AND），生成依赖位向量（Dependency Bit Vector）。如果依赖位向量有任何一位为 1，则说明存在数据冲突（依赖），该指令不能执行，反之如果全部为 0，则可以发射执行。
第三步：提高并发指令的数据供给效率
并发指令数据访问
寄存器是处理器内部的高速存储单元，用于临时存放指令执行过程中所需的操作数、中间结果和地址信息。在 GPU 中，每个 SIMT 核心都拥有独立的寄存器文件（寄存器的集合体，本质上是一组寄存器组成的存储阵列）。
第二步引入的计分板（ScoreBoard）机制，解决了时序维度上的数据依赖问题，从而支持发射无依赖指令进行延迟隐藏，除了时序上的复杂性，指令并行还会带来空间上的复杂性，即大量并发指令同时尝试访问寄存器文件获取指令数据，寄存器文件必须支持多 warp 并发访问。
简单粗暴：多端口寄存器文件
端口（port），是读写存储单元的接口。每多一个端口，存储单元就可以多支持一个并发读写操作，单端口的情况下，同时只能支持一个读或一个写，若一个读操作与一个写操作并发，则只能串行执行，而增加一个端口，称为双端口，则此时的一读一写就可以并发完成。
因此，为了支持大量 warp 并发访问寄存器数据，一个简单粗暴的做法是，为寄存器文件设计足够多的端口，来容纳所有并发读写操作。
尽管多端口设计在理论上可行，但其硬件代价呈指数级增长，包括导致芯片面积暴增，同时，动态功耗会随端口数平方增长、高访问延迟等。因此，简单的硬件堆料是低效且不可取的。
单端口内存库
寄存器文件与共享内存的并发访问冲突，本质上源于一个根本矛盾：存储单元的物理端口数量与程序所需的并发访问量之间的不匹配。若将多端口设计比作“拓宽车道”，则单端口内存库（Single-Ported Memory Bank）更像是“优化交通规则”——通过精细化调度，在有限硬件资源下挖掘最大效率。
在计算机存储体系结构中，banking 是一种将存储体分成多个独立的部分（bank），每个 bank 可以独立访问，从而提高并行访问能力的技术。如图所示，单端口内存库将寄存器文件分成多个 bank 后，每个 bank 可以独立进行读写操作，每个 bank 只有一个端口，如果同一时刻只有同一个线程访问，则可以成功，但如果有一个以上线程并发访问，则将产生访问冲突。
如此以来，分为多个 bank 的寄存器文件，一定程度上模拟了多端口寄存器文件的行为，即支持了跨 bank 的并发读写操作。单端口内存库也是 GPU 核心架构中最常见的片上存储单元微架构，除寄存器外，其同样应用于共享内存。
同时，为了进一步优化，有很多通过寄存器布局优化减少 Bank 冲突的机制。其中常见的有：
- 交错寄存器布局（Interleaved Register Allocation）。让不同的 warp 的同编号寄存器分配到不同 bank 上。如 warp0 的 R0 分配到 bank0，而 warp1 的 R0 则分配到 bank1。这种布局方式在 warp 均匀调度发射指令（常见调度模式，大量 warp 轮流执行）时可以有效地防止冲突。
- 动态 Bank 分配（Dynamic Bank Allocation）。根据指令的寄存器访问模式，动态调整逻辑寄存器到物理 Bank 的映射关系，避免静态固定映射导致的冲突。
- 编译器驱动的寄存器分配优化。编译器在代码生成阶段，通过智能分配寄存器，减少 Bank 冲突。
- 在以上基础上，发展出了混合 Bank 设计（Hybrid Banking）。将寄存器文件划分为不同特性的 Bank 子集，采取不同的布局分配机制，针对不同访问模式优化。
还有冲突：Operand Collector
接前文的例子，不论是“拓宽车道”还是“优化交通规则”，总会有车道争抢的问题，那么也就总是需要“路口红绿灯”来居中协调。
针对单 bank 的并发操作还是会引发数据冲突，这时就需要引入 Operand Collector（操作数收集器）进行指令的统一调度。Operand Collector 是 GPU 流水线中的一个关键硬件模块，负责在指令执行前收集所有必需的操作数（即寄存器或内存中的数据）。它的核心目标是解决寄存器文件（Register File）的 Bank 冲突问题，并通过动态调度最大化寄存器访问的并行性，从而提升指令吞吐量。
当指令进入寄存器读取阶段（Register Read Stage）时，系统为其分配一个收集单元（Collector Unit），每个收集单元为一条指令服务，负责缓存该指令所需的所有源操作数（如 ADD R1, R2, R3 中的 R2 和 R3）。收集单元向寄存器文件发送读请求，获取源操作数。例如，指令 ADD R1, R2, R3 需要读取 R2 和 R3。
当不同指令出现数据冲突时，Operand Collector 将动态调度这些冲突的请求，将冲突请求分配到不同周期排队执行。若进入排队状态，收集单元暂存已就绪的操作数，直到所有操作数准备完毕，指令拿到操作数发射执行。
4. 总结
讲到这里，已历上万字，我们从引言中“Dennard Scaling”的失效开始，引入 GPU 出现的背景，又介绍了 GPU 的通用性，以及高并发、低延迟保证的高计算速度。随后，我们以最常见的 CUDA 为例，介绍了 GPU 编程的基础，SIMT 与 SIMD，编译链接的过程。最后，我们深入硬件层面，分为三步走，先用最简系统“run 起来”，然后分别解决了指令依赖问题，以及并发执行中的数据访问冲突问题，构建并了解了一个通用 GPU 核心的架构。
本文介绍的 GPU 知识，只是对各厂商、各架构设计做“求同存异”后，得到的主干性的、通用性的基础知识，而 GPU 作为当代最为炙手可热的科技产品之一，其发展是日新月异的。笔者希望这些基础知识可以作为有兴趣的读者的“指路牌”，指引读者在本文建立起来的基础视野上，进一步探索。
END
作者：leowwlwang
文章来源：腾讯技术工程
推荐阅读
更多腾讯 AI 相关技术干货，请关注专栏腾讯技术工程 欢迎添加极术小姐姐微信（id:aijishu20)加入技术交流群，请备注研究方向。
