CUDA 性能调优实战指南：从系统视角理解GPU优化
现代 GPU 提供了惊人的计算吞吐量，但要真正在 CUDA 上实现高性能，往往远不止“启动更多线程”那么简单。认真对待 CUDA 的开发者很快会发现，语法知识与性能推理之间存在鸿沟。教程要么在介绍完基础示例后就戛然而止，要么直接跳到一些零散的优化技巧，却很少解释它们为何有效（或何时使用）。本 CUDA 指南旨在填补这一空白。
本指南从系统思维和架构感知的角度来探讨 CUDA。目的是让你能够推理你的 CUDA 代码如何实际映射到真实的 GPU 硬件、性能瓶颈如何产生，以及如何利用现代性能分析工具在优化时做出明智决策。一旦你学会了如何推理执行、内存移动和调度，优化就不再是猜测，而成为一个有章可循的工程过程。
关键要点
- CUDA 性能关乎推理，而非语法。理解线程、线程束、线程块和流多处理器如何映射到硬件，比了解 API 调用更重要。
- 内存行为主导实际性能。大多数内核是内存受限的；合并访问、缓存和数据重用通常比原始浮点运算更重要。
- 线程束是重要的执行单元。SIMT 执行和线程束分支直接影响效率——围绕线程束行为设计控制流和线程块大小。
- 占用率是手段，而非目标。需要足够的线程束来隐藏延迟，但追求 100% 占用率往往会因寄存器压力和溢出而损害性能。
- 性能分析驱动优化。Nsight Systems 识别时间花在哪里；Nsight Compute 解释原因。先测量，再优化，始终重新测量。
什么是 CUDA？
CUDA 代表计算统一设备架构。CUDA 是 NVIDIA 开发的并行计算平台和编程模型，用于在图形处理单元上进行通用计算。它允许开发者编写程序，将并行工作负载从 CPU（主机）卸载到一个或多个 GPU（设备）上执行。
在 CUDA 编程术语中，GPU 被视为加速器。CUDA 既不是操作系统，也不是独立的语言。CUDA 应用程序使用标准 C/C++ 以及一些额外的 CUDA 关键字（例如，用于 CUDA 内核的 __global__）和库进行编程，使你的代码能够在 GPU 上执行。（CUDA 包含低级的 Driver API 和高级的 Runtime API——例如 cudaMalloc、cudaMemcpy 等函数——以帮助抽象 GPU 任务。）
CUDA 支持异构计算，这意味着 GPU 拥有自己的设备内存，与 CPU 的主机内存是分开的。使用 CUDA 时，你需要显式管理主机和设备之间的内存（例如，使用 cudaMemcpy）。线程在程序中被启动以执行数据并行操作。
CUDA 执行模型：线程、线程束、线程块和网格
本节构建推理 CUDA 性能所需的思维模型。CUDA 暴露了一个分层执行模型——线程 → 线程束 → 线程块 → 网格——将你的内核映射到 GPU 硬件上。
协作规则：线程能做什么和不能做什么
CUDA 程序将在 GPU 上执行的任务实现为内核，内核是编译后在 CUDA 设备上运行的函数。当从主机启动一个内核时，它会指定一个线程块网格，其中每个线程块由许多并行执行的线程组成。网格可以概念上视为一个线程块数组（一维、二维或三维）。
每个线程独立执行内核的实例，拥有自己的寄存器和本地内存。CUDA 确保一个线程块内的线程可以通过同步和共享数据（使用屏障同步和共享内存）进行协作；然而，不同线程块中的线程没有直接的方法进行同步或共享内存。
明确地说，CUDA 编程模型假设不同线程块之间没有依赖关系；线程块可以以任何顺序执行，程序仍必须产生正确的结果。
线程束执行：SIMT、分支和线程块大小选择
线程块内的线程被组织成线程束。一个线程束由 32 个线程组成，代表一个硬件调度单元：线程束以 SIMT（单指令，多线程）方式执行；一个线程束的所有线程同时执行相同的指令（在不同的数据上）。
如果线程束的每个线程都到达某个从/向内存加载/存储的指令，那么该线程束的所有加载/存储操作都会并行执行。另一方面，如果线程发生分支（由于条件分支），底层硬件必须序列化该分支的执行：不执行该分支的线程被屏蔽，其他线程可以继续。
线程束分支会导致高性能损失，因为它未充分利用 GPU 的并行通道。在实践中，当线程束的所有线程都采用相同的执行路径时，将获得最佳性能。
线程块大小通常应为线程束大小的倍数（32、64、128、256、512）。这通常可以消除最后一个线程束利用率不足的缺点，并使线程束调度粒度与线程块形状相匹配。
CUDA 执行层次和线程束级 SIMT 行为
让我们考虑下图：
在上图中，内核启动指定了一个线程块网格。每个线程块包含许多线程，这些线程被调度为包含 32 个线程的线程束。线程块内的线程可以通过共享内存和屏障同步（例如，syncthreads()）进行协作。线程块不能直接同步或共享内存；它们是独立的。线程块的执行顺序没有保证。
SIMT 编程模型应用于线程束内部；线程束分支导致并行路径的屏蔽执行和分支路径的序列化。
CUDA 的层次结构可以总结为：
- 网格 -- 内核启动的所有线程的集合
- 线程块 -- 网格被细分为线程块
- 线程束 -- 线程块被细分为 32 个线程的线程束
- 线程 -- 每个线程都有自己的索引（blockIdx, threadIdx），它用这个索引来确定应该计算哪些数据。
这种三层层次结构使内核能够处理任意大的数据集，因为线程块可以由可用硬件以任何顺序调度。编程模型的抽象也隐藏了线程束如何调度到 GPU 硬件上。
GPU 架构基础（影响性能）
要优化内核，你必须首先建立一个以硬件为先的思维模型，了解性能从何而来：工作如何调度到流多处理器上、哪些片上资源是有限的（寄存器、共享内存/L1、线程束发射槽），以及内存层次结构如何影响延迟和带宽。下面，我们将这些硬件约束直接映射到实际结果，如占用率、溢出和延迟隐藏。
流多处理器要点：线程束执行引擎
现代 NVIDIA GPU 拥有流多处理器。流多处理器是负责调度/执行线程束的 GPU 计算硬件块。它包含许多资源，如功能单元（浮点/整数 ALU、特殊功能单元等）、大型寄存器文件、片上共享内存和可配置的 L1 数据缓存。
所有流多处理器都可以访问共享的 L2 缓存，以及片外全局 DRAM 内存。用 NVIDIA 编程指南中的话来说：“每个流多处理器包含一个本地寄存器文件、一个统一数据缓存和功能单元。”统一缓存包括 L1 缓存和共享内存空间（通常可以由程序员调整两者的分配比例）。
关键的硬件资源在每个流多处理器上是有限的。例如，一个 NVIDIA Ampere 或 Hopper GPU 的流多处理器可能拥有大约 64–128 KB 的寄存器和 64–96 KB 的可配置 L1/共享内存。此外，可能只有少数线程束调度器（通常为 2 或 4 个）可以在每个周期从每个流多处理器发射指令。
资源限制和占用率：寄存器、共享内存和溢出
然而，这些限制的实际表现与占用率有关：过多的线程（或请求过多寄存器的线程）可能会使流多处理器的寄存器或共享内存饱和。
例如，在计算能力 7.0（Volta）上，我们有一个拥有 65,536 个 32 位寄存器和 2048 个并发线程（64 个线程束）的流多处理器。如果每个线程使用 32 个寄存器，那么流多处理器可以调度高达 100% 的占用率（2048 线程 * 32 寄存器 = 65,536）。但是，如果每个线程使用 64 个寄存器，则只能容纳总共 1024 个线程（65,536/64 = 1024 线程），这会降低我们的占用率。
在实践中，我们经常限制寄存器或调整线程块大小以获得“足够好”的占用率，同时避免溢出到本地内存。
流多处理器上的线程束调度：轮询发射、延迟隐藏和占用率上限
一个硬件特性是线程束调度。流多处理器将以轮询（双发射）方式循环遍历其每个活动线程束，并尝试发射指令。
总之，需要记住的主要硬件单元是：
- 流多处理器 – 可以在一个流多处理器上同时执行许多线程块/线程束。每个流多处理器的最大线程块数由寄存器、共享内存和最大线程数决定。
- 寄存器 – 最快的内存，线程私有。如果每个线程使用太多寄存器，那么可以并发执行的线程束数量就会减少。
- 共享内存 / L1 缓存 – 每个流多处理器的片上暂存器，如果访问没有存储体冲突，延迟非常低。存储体通常为每个流多处理器 32 个（每个 32 位字对应一个存储体）。共享内存访问模式必须围绕避免存储体冲突来设计。
- L2 缓存 – 片上（更大、更慢）的缓存，由所有流多处理器共享。
- 全局（设备）内存 – 片外 DRAM，带宽高但延迟也显著更高。必须优化全局内存访问模式（合并访问）。
CUDA 内存层次结构
CUDA 有多个内存空间，每个空间具有不同的延迟、带宽和作用域。从最快到最慢：
- 寄存器：每个线程都有自己的私有寄存器文件。寄存器基本上零延迟，带宽最高。在适当的情况下，编译器会将频繁访问的局部变量放在寄存器中。（寄存器在代码中不显式存在；由编译器处理。）
- 本地内存：如果局部变量需要的内存超过寄存器可用量（例如，溢出的变量），编译器会在本地内存中分配此内存。物理上，此内存位于全局内存（设备 DRAM）中。访问此内存与直接访问全局内存一样慢，因此避免溢出很重要。
- 共享内存：一个显式管理的片上内存空间，由线程块内的线程共享（此作用域生命周期有限）。线程块内的线程可以快速读写（约几十个周期）共享内存区域。共享内存被划分为存储体以支持同时访问。如果线程束中的 32 个线程各自访问不同的存储体，则所有加载/存储操作在一个周期内完成。
- L1 缓存：在许多 GPU 上，从全局内存的加载通常会经过 L1 缓存（有时称为共享内存，取决于程序员的配置方式）。（在较新的架构，如 Ampere/Hopper 中，数据和纹理/常量有独立的“L1”缓存。）每个流多处理器的总 L1 缓存大小通常在 128–192 KB 左右。L1 缓存有助于隐藏线程束或线程块内重用数据的延迟。
- L2 缓存：还有一个统一的 L2 缓存（通常几 MB），由所有流多处理器共享。所有全局内存请求都经过 L2 缓存。L2 缓存在线程块之间缓存数据，如果你的程序在线程束/流多处理器之间重用数据，可以帮助减少片外流量。
- 全局（设备）内存：这是设备上最大的内存，通常位于片外 DRAM 中。全局内存具有最高的带宽（约数百 GB/s），但也有很高的延迟（约数百个周期）。对全局内存的糟糕访问模式很容易浪费内存带宽。
内存合并
全局内存以 32、64 或 128 字节块（取决于架构和访问大小）进行访问。CUDA 内存子系统试图将线程束发出的内存访问合并为尽可能少的事务。
例如，如果线程束中的连续线程访问全局内存中连续的 4 字节字，则所有访问都可以通过获取几个 128 字节段来满足。我们称之为合并访问。当内存访问未合并时（例如，线程以较大跨度访问内存），许多字节被浪费。
对于从全局内存获取的每 32 字节，只有一小部分数据被线程使用。NVIDIA 工具会显示，在这种情况下，每次加载只使用了 32 字节中的 4 字节。这浪费了 87.5% 的带宽！
为了获得高全局内存吞吐量，应在连续线程中对齐和访问内存。跨步/分散的内存访问会导致事务中大部分数据为空；这会降低有效带宽。
在实践中，这意味着要塑造数据结构和循环索引，使得线程 tid 访问 array[tid]，或每个线程访问一小段连续块，而不是跨步访问模式。
共享内存存储体冲突
共享内存逻辑上被划分为存储体（现代 GPU 上有 32 个，每个 32 位字对应一个）。每个内存地址根据其地址索引映射到一个存储体。如果线程束中的线程访问 N 个存储体中的内存地址，那么这 N 个内存访问可以并行执行。这将实现单存储体访问带宽的 N 倍。
然而，如果多个线程访问同一个存储体（但不同地址），就会发生存储体冲突：硬件必须发出多个子请求，将它们序列化。下图说明了这一点：
其他内存空间
- 常量内存：可由所有线程访问的缓存只读内存。适用于可以广播到每个线程的少量数据。
- 纹理内存：针对 2D 空间局部性优化的缓存只读内存。适用于图形或数据插值。
- 统一虚拟内存：CUDA 运行时功能，允许 CPU 和 GPU 共享一个共同的地址空间。页面由系统按需迁移。允许更简单的编程模型，但如果使用不当，可能会遭受隐藏的页面错误。
由于大多数现实世界的内核是内存受限的，最大的收益来自合并访问、缓存和最大化数据重用（在寄存器或共享内存中）。
工程师使用的性能模型
这些模型为你提供了一种快速方法来确定你的内核是计算受限、内存受限还是延迟/占用率受限。这样，你可以将优化重点放在最有效的地方。
计算受限 vs. 内存受限（算术强度和屋顶线思维）
CUDA 工程师将性能问题框定为计算受限与内存受限。计算受限意味着你的内核每访问一个字节有足够的算术操作来饱和 GPU 的 ALU（浮点单元）。大型稠密矩阵乘法内核是计算受限的经典例子。这些内核很容易接近 GPU 的峰值浮点运算速率（理论峰值的 90–95%）。
内存受限的内核表现出低算术强度（低浮点运算/字节），当它们饱和内存带宽时达到峰值性能。例如，归约（求和）内核每两次读取执行一次加法，因此它几乎永远不会以高于峰值浮点运算速率很低百分比的速度执行；相反，它会饱和内存带宽。
理解内核是内存受限还是计算受限可以帮助你理解优化工作的重点：内存受限的内核可以从仔细的数据布局/缓存中获得更多收益，而计算受限的内核可能受益于使用 Tensor Core 或调整指令级效率。
你可以简单地将其视为屋顶线模型：性能 ≤ min(计算峰值, 算术强度 × 带宽)
这里不需要公式，但从概念上讲，如果测量到的性能远低于这两个限制，你可能存在一些瓶颈（如糟糕的内存访问模式）。
占用率 vs. 延迟隐藏（为什么“100% 占用率”不是目标）
另一个关键术语问题是占用率与延迟隐藏。占用率是活动线程束与每个流多处理器最大可能线程束的比率。高占用率确保每个流多处理器上有足够的驻留线程束来隐藏内存延迟：当一个线程束在等待内存时，另一个可以发射和执行。
然而，高占用率并不自动等于高性能。
你会发现，如果占用率超过某个适度值（比如 20–40%），那么向线程块添加更多线程的回报就会递减。如果你使用更高的占用率但每个线程有更多的寄存器，你甚至可能看到性能下降，尽管线程数更少（占用率更低）。
因此，不要以 100% 占用率为优化目标，而是针对真正的瓶颈：如果你是内存受限的，利用占用率来隐藏延迟。如果你是计算受限的，确保你正在喂饱你的 ALU。
性能分析工作流程
在优化之前，先测量时间和停顿来自哪里。NVIDIA 提供了两个主要工具：
- Nsight Systems 是一个系统范围的时间线分析器。使用 Nsight Systems，你可以在单个时间线上可视化 CPU 线程、GPU 内核、内存复制和主机活动。Nsight Systems 允许你查看应用程序，了解内核和 CPU 任务在哪里重叠、应用程序在哪里空闲，以及 GPU 和 CPU 是否得到良好利用。
- Nsight Compute 是一个侧重于设备指标的内核分析器。它收集每个内核的性能计数器（流多处理器占用率、发射的指令、内存事务等），并将它们汇总到表格中进行分析。使用 Nsight Compute，你启动你的内核（通常多次）并接收指标，如实现的占用率、内存吞吐量、线程束停顿的原因、加载/存储指令的利用率等。
通常，你可以先启动 Nsight Systems 来识别慢速内核/阶段，然后启动 Nsight Compute 来检查一个热点内核的指标。
可重现的性能分析手册
性能分析会产生开销并干扰计时，因此确保测量的可重现性非常重要。如果应用程序序列化其内核启动、在多次运行之间刷新缓存并“固定”GPU 时钟，Nsight Compute 将报告确定性指标。然而，由于热节流、驱动程序可变性以及与其他进程的并发性，测量噪声仍然存在。
为了最小化可变性：
- 在测量前预热内核。确保 GPU 处于稳定时钟且缓存已预热。
- 尽可能固定时钟频率。在 Linux 上，运行 nvidia-smi -lgc <min, max> 并启用持久模式。
- 将内核固定到流并避免并发分析；Nsight Compute 将序列化内核。这对于其他分析工具不成立。
- 在测量运行之间刷新缓存，或使用 --cache-control 选项来控制是否在运行之间刷新缓存。
- 使用固定的输入大小和随机种子。即使问题大小有微小变化，也可能改变算术强度和内存访问模式。
- 重复测量并报告均值和方差。如果你不使用 Nsight 工具进行计时，请使用 cudaEvent 计时器或 nvprof --metrics 进行低开销计时。
优化手册：症状 → 原因 → 修复
基于性能分析和经验，我们可以概述常见症状及相应的行动：
症状：实现的内存带宽低（且浮点运算速率低）
可能原因：未合并或低效的全局内存访问。过多的全局内存流量（太多加载/存储）。数据重用差（未利用片上内存/缓存）。
推荐修复（决策规则）：合并全局访问：确保线程束线程访问相邻地址（例如，线程 i 读取元素 i 或 i+1，而不是 i*stride）。减少事务：考虑更小的数据类型或向量化加载（在安全且自然的情况下）。增加重用：将重用数据暂存到寄存器/共享内存（或依赖 L1）以减少全局流量（例如，在共享内存中进行矩阵乘法子块）。最小化每浮点运算的全局 I/O：减少冗余读取/写入；将中间值保留在片上。
症状：Nsight Compute 中停顿率高（内存依赖 / 执行依赖）
可能原因：内存依赖：长延迟加载；缓存未命中；冗余加载；合并差。执行依赖：长指令依赖链。线程束分支：分支导致线程束内序列化执行。共享内存存储体冲突：序列化的共享加载/存储。
推荐修复（决策规则）：对于内存依赖：修复合并；减少冗余加载；考虑使用 __ldg() 处理只读数据（如果它能改善缓存）。对于执行依赖：展开循环；将工作分解为独立指令以增加指令级并行性。对于分支：重构分支（例如，先分组标记元素）；使用线程束同步模式保持线程束在同一路径上。对于存储体冲突：填充共享数组，使线程映射到不同的存储体（例如，向 2D 共享块添加填充列）。
症状：启动了许多小内核（或每次启动的计算量不足）
可能原因：内核启动开销占主导。每次启动的并行工作不足。工作分散在许多微小内核中。
推荐修复（决策规则）：融合内核：将多个步骤合并到一个内核中以减少启动次数。批处理工作：避免在循环中每行/元素启动；启动一个处理整个工作负载的内核。使用并发工具：使用 CUDA 流或 CUDA Graphs 来批处理独立任务并隐藏启动延迟。
症状：算术密集型内核的计算吞吐量（浮点运算速率）低
可能原因：内存停顿使 ALU 饥饿。指令吞吐量差（低效的指令组合）。指令级并行性不足。
推荐修复（决策规则）：先检查内存：确认内核确实是计算受限的（不是因内存而停顿）。使用计算特性：在适用时优先使用内部函数（例如，MAD / __fmul_rn）并考虑 Tensor Core（混合精度）。向量化加载：确保编译器在对齐/布局允许时可以使用向量加载（例如，float4）。增加指令级并行性：展开循环；重构以保持多个操作在流水线中。验证线程束可用性：使用占用率分析确保有足够的线程束保持流水线繁忙（不需要 100% 占用率）。
症状：占用率有限（由于寄存器/共享内存）
可能原因：每个线程使用太多寄存器。每个线程块使用太多共享内存。线程块大小/资源使用阻止更多线程块适合每个流多处理器。
推荐修复（决策规则）：如果占用率非常低（<10–20%），减少寄存器压力（重构代码或谨慎使用 -maxrregcount）。如果共享内存占用阻止了驻留，则减少每个线程块的共享内存占用。调整线程块大小以增加每个流多处理器的线程块数（同时注意溢出和存储体冲突）。使用 CUDA 占用率计算器或 Nsight Compute 的计算器来探索权衡（如果高寄存器能避免溢出，仍然可能胜出）。
并发性：流和重叠
下表试图捕捉 CUDA 流提供的不同并发类型。
概念：默认行为（单流）
含义：在同一流中发出的操作（内核启动、内存复制）顺序执行。
重要性（实际影响）：简单的顺序和正确性，但可能让 GPU 或复制引擎空闲（无重叠）。
说明 / 约束：良好的基线；当传输和计算可以重叠时，对于吞吐量通常不是最优的。
概念：多流（并发）
含义：使用多个流允许 GPU 在可能时重叠独立任务（计算和/或复制）。
重要性（实际影响）：通过并行运行工作来提高利用率（例如，在计算当前批次时复制下一批次）。
说明 / 约束：只有当操作真正独立且硬件资源可用时，重叠才会发生。
概念：计算/传输重叠（流水线）
含义：当内核在流 0 中运行时，你在流 1 中启动 cudaMemcpyAsync 来传输下一个输入。
重要性（实际影响）：通过在复制引擎上进行复制而流多处理器进行计算，来隐藏 PCIe/NVLink 传输延迟。
说明 / 约束：现代 GPU 可能拥有多达两个复制引擎（主机到设备和设备到主机），当分区良好时可以实现更充分的重叠。
概念：并发内核（多流）
含义：在不同流中启动多个内核，以便如果流多处理器资源剩余，它们可以同时运行。
重要性（实际影响）：当单个内核未完全占用 GPU（未使用的流多处理器容量）时，可以增加吞吐量。
说明 / 约束：真正的内核并发通常有限；硬件可能对内核进行时间分片，而不是完全并行执行。
概念：流的实际主要价值
含义：流对重叠计算与数据移动影响最大。
重要性（实际影响）：通常，最大的端到端加速来自隐藏传输，而不是同时运行许多内核。
说明 / 约束：始终在时间线分析器（如 Nsight Systems）中验证重叠是否发生。
结论
在本指南中，我们从 CUDA 语法中退后一步，从系统视角看待问题。我们探讨了内核如何在物理 GPU 硬件上执行，以及理解执行模型、内存层次结构和调度行为——加上正确使用性能分析工具——如何让你能够准确识别真正的瓶颈并自信地应用有意义的优化。
有效的 CUDA 性能来自推理、测量和迭代。如果你让架构特性和你的数据指导优化，而不是盲目跟随直觉，你会发现 CUDA 编程是一门工程学科，而不是一系列临时调整的集合。
常见问题解答
为什么启动更多线程不足以获得良好的 CUDA 性能？
因为性能受限于内存访问模式、线程束执行和硬件资源约束——而不仅仅是并行性。
在实践中，最常见的 CUDA 性能瓶颈是什么？
全局内存访问效率低（合并差、重用低、缓存未命中）是大多数实际内核中的主要瓶颈。
高占用率总是好的吗？
不。占用率高于约 20–40% 通常就足够了。超过这个值，额外的线程束很少有帮助，并且可能因寄存器溢出而降低性能。
线程束分支何时真正有害？
当不同执行路径具有相当的工作量并为许多指令序列化时，分支有害；短或统一的分支通常没问题。
我应该先使用哪个分析器：Nsight Systems 还是 Nsight Compute？
从 Nsight Systems 开始，找到慢速阶段和重叠问题，然后使用 Nsight Compute 深入分析一个热点内核。
参考资料
原文链接：CUDA Guide: Workflow for Performance Tuning |
发表于：