📚 课程目录第2章 GPU计算基础

2、GPU计算基础：GPU架构、线程层次、内存层次、并行计算模型与性能关键指标

GPU（图形处理器）最初为图形渲染而设计，但其高度并行的计算能力使其成为通用高性能计算（GPGPU）的核心硬件。要使用Triton编写高效的GPU算子，首先需要理解GPU的硬件架构、线程组织方式、内存体系以及并行计算模型。这些基础知识决定了算子的设计边界和优化方向。本章将从硬件架构出发，逐步深入到线程与内存层次，最后给出性能评估的关键指标，为后续Triton编程打下坚实的理论基础。

上图展示了现代GPU（以NVIDIA Ampere/Hopper架构为例）的核心组成：多个GPC（图形处理簇）内包含若干SM（流式多处理器），每个SM拥有独立的计算核心、L1缓存和共享内存；所有SM通过统一的L2缓存与HBM显存交互。CPU通过PCIe总线与GPU通信，但数据传输带宽远低于GPU内部带宽，这决定了算子设计必须尽量减少CPU-GPU之间的数据拷贝。

2.1 GPU架构核心原理：从SM到Warp的硬件映射

GPU的硬件设计哲学是“以吞吐量换延迟”——用大量线程的并行执行来隐藏内存访问延迟。理解这一哲学需要从三个硬件层级入手：SM（流式多处理器）、CUDA Core和Warp调度器。

每个SM包含若干CUDA Core（整数/浮点运算单元）、特殊函数单元（SFU）、加载/存储单元（LD/ST）以及寄存器文件。以A100为例，每个SM有128个FP32 CUDA Core，4个Warp调度器，每个调度器每周期可发射一条指令。SM内部还配置了L1缓存和共享内存（共192KB，可动态划分）。线程以Warp（32个线程）为基本调度单位：Warp中的所有线程在同一时刻执行同一条指令（SIMT模式），但可以有不同的数据路径。如果Warp内出现分支发散（if-else），则不同分支会被串行执行，导致性能下降。

硬件线程调度的关键机制是零开销上下文切换。当某个Warp因等待内存数据而阻塞时，调度器立即切换到另一个就绪的Warp执行，切换过程仅需一个周期。因此，GPU需要足够的活动Warp数量（通常每个SM至少16个Warp）来隐藏内存延迟。这就是为什么GPU算子需要创建大量线程——不是每个线程都在“干活”，而是让硬件有足够的“备胎”来填补等待空隙。

硬件层级	组成	关键参数（A100）	对编程的影响
GPC	多个SM + 光栅化单元	8个GPC	决定几何处理能力，计算任务中影响较小
SM	CUDA Core + SFU + LD/ST + 寄存器 + L1/共享内存	108个SM，128 FP32 Core/SM	线程块（Block）被分配到单个SM上执行
Warp	32个线程的集合	4个调度器/SM，每周期发射1条指令	分支发散会串行化，需避免Warp内分歧
CUDA Core	FP32/INT32运算单元	时钟频率1.41GHz	算力峰值 = SM数 × Core数 × 频率 × 2

经验提醒：在Triton中，你不需要直接管理Warp，但Triton编译器会自动将你的“块级”操作映射到Warp上。因此，编写Triton算子时，应尽量保证块内所有元素执行相同的操作路径，避免在块内使用依赖数据值的条件分支。

2.2 线程层次与内存层次：Triton的编程抽象

CUDA的线程层次为Grid → Block → Warp → Thread，而Triton将其简化为Grid → Program两级抽象。在Triton中，一个kernel对应一个Grid，每个program（即一个实例）对应CUDA中的一个Block。Triton编译器负责将program内部的张量操作自动映射到Warp和Thread上，开发者无需关心线程ID的计算。

内存层次方面，GPU从快到慢依次为：寄存器（每个线程私有，容量极小但速度最快）→ 共享内存/一级缓存（SM内部，可编程控制，延迟约20-30周期）→ 二级缓存（芯片级统一缓存，延迟约200周期）→ 全局内存（HBM）（容量最大，延迟约400-800周期，带宽极高）。此外还有只读纹理缓存和常量缓存，但在通用计算中较少使用。

Triton的内存模型与CUDA对应关系如下：

全局内存：通过tl.load/tl.store访问，对应CUDA的__global__内存，所有program可见。
共享内存：Triton中通过tl.zeros或中间结果自动分配，对应CUDA的__shared__内存，同一program内的所有线程可见。
寄存器：Triton中的标量变量和中间张量通常驻留在寄存器中，编译器自动管理。

理解内存层次的核心在于数据复用：如果算子需要多次访问同一数据，应尽量将其加载到共享内存或寄存器中，避免重复访问全局内存。例如，矩阵乘法中，每个输出元素需要读取A的一行和B的一列，如果直接访问全局内存，会产生大量冗余读取；而通过分块（Tiling）将数据块加载到共享内存，可显著减少全局内存访问次数。

核心原理：GPU性能优化的第一法则是“最大化计算/访存比”。对于每个字节从全局内存加载的数据，至少要执行足够多的算术运算来“掩盖”访存延迟。以A100为例，其FP32算力约19.5 TFLOPS，HBM带宽约2TB/s，因此计算/访存比约为10:1。这意味着每加载1字节数据，至少需要执行10次浮点运算才能达到计算瓶颈。

2.3 并行计算模型与性能关键指标

GPU并行计算模型遵循SPMD（单程序多数据）模式：所有线程执行同一个内核函数，但处理不同的数据。在Triton中，这种模型被进一步封装为“块级编程”——你编写的是对整个张量块的操作，而不是单个元素的操作。Triton编译器负责将块操作分解为线程级操作。

评估GPU算子性能的关键指标包括：

吞吐量（Throughput）：单位时间内处理的数据量（如GFLOPS、GB/s）。这是最直接的性能指标。
占用率（Occupancy）：活跃Warp数占SM最大可容纳Warp数的比例。高占用率有助于隐藏延迟，但过高的占用率可能导致寄存器溢出。
内存带宽利用率：实际带宽与峰值带宽之比。对于访存密集型算子，这是主要瓶颈。
计算效率：实际算力与理论峰值算力之比。对于计算密集型算子，应接近100%。
延迟（Latency）：单个请求从输入到输出的时间。在推理场景中，延迟与吞吐量需要权衡。

判断一个算子是计算密集型还是访存密集型，可以使用算术强度（Arithmetic Intensity）：总浮点运算次数 ÷ 总字节访问量。如果算术强度大于硬件计算/访存比（约10），则为计算密集型；否则为访存密集型。例如，矩阵乘法（C = A×B）的算术强度随矩阵维度增大而线性增长，大矩阵乘法是计算密集型；而向量加法（y = x + y）的算术强度仅为1，是典型的访存密集型。

指标	定义	优化目标	测量工具
GFLOPS	每秒十亿次浮点运算	接近理论峰值	NVIDIA Nsight Compute
带宽利用率	实际带宽/峰值带宽	≥80%	Nsight Compute / nvprof
占用率	活跃Warp/最大Warp	50%-75%为宜	Nsight Compute
算术强度	FLOPs / Bytes	≥10（A100）	手动计算

2.4 案例拆解：向量加法的性能分析

我们以最简单的向量加法（y = x + y）为例，演示如何分析GPU算子性能。该算子算术强度为1，属于访存密集型，理论性能上限由内存带宽决定。

场景与目标：在A100（峰值带宽2TB/s）上，对长度为1亿（400MB）的float32向量执行加法，目标达到80%以上的带宽利用率。

初始实现与问题：使用Triton编写朴素实现，每个program处理1024个元素，grid大小为97656。初步测试带宽利用率仅45%，远低于预期。

@triton.jit
def vector_add_kernel(X, Y, N, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < N
    x = tl.load(X + offsets, mask=mask)
    y = tl.load(Y + offsets, mask=mask)
    tl.store(Y + offsets, x + y, mask=mask)

分析过程：使用Nsight Compute分析发现，内存事务效率仅为55%，存在大量非合并访问。原因在于BLOCK_SIZE=1024时，每个program访问的地址范围跨越了多个内存页，导致TLB（页表缓存）未命中率升高。此外，grid过大导致调度开销增加。

优化步骤：

将BLOCK_SIZE调整为2048，减少program数量，提高每个program的内存访问连续性。
使用tl.max_contiguous和tl.multiple_of提示编译器内存对齐信息，帮助生成更高效的向量化加载指令。
启用L2缓存持久化（cudaAccessPolicyWindow），确保数据在L2中驻留更久。

结果对比：优化后带宽利用率从45%提升至82%，耗时从1.2ms降至0.65ms。进一步增加BLOCK_SIZE到4096没有明显提升，因为已经接近带宽上限。

复盘经验：访存密集型算子的优化重点是合并访问（相邻线程访问相邻地址）和减少内存事务。Triton的块级抽象天然支持合并访问，但需要合理选择块大小。此外，mask操作会降低效率，应尽量让N是BLOCK_SIZE的整数倍。

常见误区：初学者常认为“线程越多越好”，但过大的grid会导致调度开销和TLB压力。对于访存密集型算子，应优先保证内存访问的连续性和对齐性，而不是盲目增加线程数。另外，不要忽视CPU-GPU之间的数据传输开销——如果数据需要频繁往返，PCIe带宽（约32GB/s）会成为真正的瓶颈。

2.5 总结复盘：GPU算子的设计原则

本章的核心知识可归纳为以下设计原则：

理解硬件映射：Triton的program对应CUDA Block，编译器自动处理Warp调度。开发者应关注块大小和寄存器使用，而非线程ID。
内存层次利用：优先使用共享内存进行数据复用，减少全局内存访问。Triton中可通过tl.load到局部变量实现类似效果。
计算/访存比：根据算术强度判断算子类型，针对性地优化。访存密集型追求带宽饱和，计算密集型追求指令级并行和ILP。
避免分支发散：在Triton中，块级条件分支会导致整个块串行执行，应尽量使用掩码（mask）操作代替分支。
性能分析工具：NVIDIA Nsight Compute是必备工具，它提供占用率、内存吞吐、指令效率等详细指标，帮助定位瓶颈。