Ben Pinzone, Compute Developer Technology Engineer
David Clark, Compute Developer Technology Engineer
GTC 2025, March 17th, 2025
简要说明这些概念在性能分析工具(Nsight Compute)中的体现。
概述通用优化策略,并使用具体示例。
有些要点会快速提及或仅引用其他资源,目的在于提高认知。
NVIDIA H100 SXM
该图展示了 NVIDIA H100 SXM GPU 的整体架构,主要组件包括:
- 132 个 SMs (Streaming Multiprocessors) 和 第四代 Tensor Cores。
- PCIe Gen 5:128 GB/s 双向带宽。
- 50 MB L2 缓存。
- 80 GB HBM3:3.3 TB/s 双向带宽。
- 第四代 NVLink:900 GB/s 双向带宽。
Hopper 架构
Hopper 架构的 SM 包含以下特性:
- SM 拥有 4 个子分区 (sub-partitions)。
- 128 个 FP32 单元。
- 64 个 FP64 单元。
- 64 个 INT32 单元。
- 4 个混合精度 Tensor Cores。
- 16 个特殊功能单元 (special function units, 用于超越函数)。
- 4 个线程束调度器 (warp schedulers)。
- 32 个加载/存储 (LD/ST) 单元。
- 64k 个 32-bit 寄存器。
- 256 KiB 统一 L1 数据缓存和共享内存。
- Tensor 内存加速器 (Tensor Memory Accelerator - TMA)。
"编织,第一种并行技术" - CUDA C++ 编程指南
</blockquote>Ceil(每个块的线程数 / 线程束大小)单指令,多线程 (Single-Instruction, Multiple-Thread)
对于某些工作负载,其昂贵的计算由一个轻量级的检查来守护。一个简单的实现可能会因为并非所有线程都能通过该检查而遭受高度的分歧。
解决方案:
该过程分为两个阶段:
1. 侦察阶段 (Scouting phase): 线程块中的线程并行地寻找工作。当某些线程(图中以绿色'x'标记)找到有效工作时,它们会将工作项添加到共享内存中的队列。
2. 处理阶段 (Process phase): 当队列中累积的工作项达到一定数量(Invoke process size)时,整个线程块会切换到处理阶段,共同从队列中取出并执行这些工作项。这确保了在处理阶段所有线程都在执行相同的任务,从而避免了分歧。
以下几页通过一个简化的、逐周期的示例,构建了一个用于理解 Warp 调度器行为和相关性能指标的思维模型。
Warp 状态定义:
- Unused (未使用): 空闲的 Warp 插槽。
- Active (活动): Warp 已被分配到该插槽。
- Stalled (停滞): 活动的 Warp 正在等待某个依赖项(例如,内存读取、指令执行完成),当前周期无法被调度。
- Eligible (符合条件): 活动的 Warp 已准备好执行下一条指令。
- Selected (已选择): 在当前周期,从所有符合条件的 Warp 中被调度器选中并发射指令的 Warp。
周期 N:
在一个给定的周期(Cycle N)中,调度器从所有符合条件的 Warp(图中浅绿色块)中选择一个(图中带斜线的浅绿色块)来发射(Issue)其指令。
周期 N+1:
调度器选择下一个符合条件的 Warp(插槽4)。上一个周期被选择的 Warp(插槽3)现在变为停滞状态,因为它正在等待指令完成,因此不符合被调度的条件。(当然,如果它有其他独立的指令可以执行,情况可能会有所不同)。
周期 N+2:
位于插槽4的 Warp 执行完毕并退出。此时,没有任何符合条件的 Warp,因此调度器的发射单元(Issue Slot)在本周期处于空闲状态,未发射任何指令。
周期 N+3:
新的 Warp 被调度到之前空闲的插槽中。调度器从新的符合条件的 Warp 中选择一个(插槽2)并发射指令。
基于以上4个周期的示例,我们可以计算出一系列聚合的性能指标。
1. cycles_active (活动周期数): 4
- 观察到的
每个调度器一个 warp 的情况下的度量(报告中的数据,已四舍五入):
warps_active (活跃 warps): 1.00warps_stalled (停滞 warps): 0.87warps_eligible (合格 warps): 0.13warps_selected (已选 warps): 0.13上图显示了调度器发出指令的摘要。每个调度器维护一个可以发出指令的 warp 池。理论上,池的上限受启动配置的限制。在每个周期,调度器检查池中已分配 warp 的状态(Active Warps)。未停滞的 warp(Eligible Warps)准备好发出它们的下一条指令。调度器从合格 warp 集合中选择一个 warp 来发出一或多条指令(Issued Warps)。在没有合格 warp 的周期中,发出槽被跳过,没有指令发出。许多被跳过的发出槽表明存在延迟隐藏问题。
GPU通过并行处理大量任务(in-flight work)来掩盖延迟。增加Warp的数量可以隐藏延迟。
下图展示了通过增加Warp数量来填充指令流水线,避免了因单个Warp停滞(Stalled)而导致的“无Warp分发”(No warp issuing)的空闲周期。当Warp 0停滞时,调度器可以选择Warp 1或Warp 2来执行,从而保持硬件的繁忙状态。
利特尔定律(Little's Law)在此处的应用描述了我们需要多少在执行中的指令(instructions in flight)才能避免暴露延迟。
另一种隐藏延迟的方法是增加指令级并行性(Instruction Level Parallelism, ILP)。
如下面的代码和图表所示,通过使用float2代替float,在同一个循环迭代中执行两个独立的操作(result.x += 3.14f; 和 result.y += 3.14f;)。这使得即使在一个Warp内部,当一条指令等待时(例如,等待前一指令的结果),硬件也可以调度执行另一条独立的指令,从而减少停滞,提高执行效率。
如果SM或内存系统资源已经繁忙,则无需担心停滞或未使用的分发槽。
否则,你的程序就是延迟绑定的。需要为硬件提供更多的并发工作。可以尝试以下方法:
占用率是在一个SM上可以并发活动的最大Warp数量。
CUDA内核的可实现占用率将受到以下至少一个因素的限制:
SM资源分配:
块大小。
可以使用NVIDIA Nsight Compute分析CUDA内核的占用率。
</blockquote>占用率可以用一个“发牌”的比喻来理解:线程块(Blocks)就像牌,它们有特定的资源需求(如共享内存、寄存器)。发牌官(SM)在有足够资源服务它们的情况下,将线程块并发地分发出去。
上图案例:给定块大小为128个线程(即4个Warp),所需的共享内存、寄存器和各种硬件限制。
下一步,我们可以通过两种方式提高占用率:
--ptxas-options=-v编译可以报告每个线程的寄存器数量。最大寄存器数可以手动设置:
-maxrregcount标志。__launch_bounds__或__maxrreg__限定符。Hopper架构每个SM拥有64K(65536)个寄存器。
示例:
优化技巧
有限的本地内存使用可能对性能有益。
下图显示了大量的本地内存请求(3.70M Req),这些请求被发送到L1/TEX缓存。
在NVIDIA Nsight Compute (NCU) 中,可以通过查看源代码视图和SASS(汇编代码)视图中的“Live Registers”(活跃寄存器)列来识别寄存器压力大的
使用尽可能轻量级的工具。数学运算的成本从高到低排列如下:
黄色: 除法 (Division)、取模 (mod) 运算符。
constexpr 进行的除法/取模通常没有问题。浅绿色: 乘法 (Multiply)、加法 (add)、减法 (subtract)。
利用快速数学优化 (--use_fast_math)。
__cosf(), __expf(), __frcp_{0}, __fsqrt_ 等。下图展示了默认数学 API 路径与 __expf 或快速数学路径的对比,以及 CUDA 编程指南中关于内在函数误差范围的信息。
无符号整数溢出是已定义行为,编译时需要考虑这一点,可能导致额外的指令。有符号整数溢出是未定义行为,这为编译器生成更快的代码提供了更大的灵活性。
优化技巧:使用有符号整数而不是无符号整数作为循环计数器。
下面的代码和图表展示了使用 unsigned int 和 int 作为循环计数器的性能差异。
给定一个线段数组,计算有多少交点发生。
与C语言除法运算符相关的SASS(汇编)调用了一个注入的辅助SASS(函数)。
FCHK指令之后,CALL.REL.NOINC指令会跳转到慢速路径,而BRANCH指令则会跳过CALL指令(快速路径)。FCHK指令的谓词依赖关系。观察除法结果的使用方式。
新算法:线段分割二维平面。
这是一种轻量级的预检查:在运行任何相交测试之前,检查两个线段的包围盒是否重叠。
设法将输入的浮点数变得更小。
FCHK指令不会标记溢出风险。这种方法依赖于具体应用,可能无法实现。
下表展示了在不同浮点数量级、不同实现和优化策略下的运行时间。
本节讨论在编译时已知常数的情况下,重复计算高阶多项式的优化方法。
通过将 pow 函数调用替换为运行中的指数计算来优化。初始基线使用 pow 函数,运行时间为 40,862 us。通过手动展开计算,避免调用通用函数,性能得到显著提升。
优化技巧:尽可能优先使用更快、更专业的数学函数,而不是更慢、更通用的函数。
使用霍纳方法可以减少指令数量,进一步优化多项式求值。该方法通过重构表达式来减少乘法操作。
优化技巧:检查表达式是否可以被重构以产生更少的指令。
通过使用融合乘加(FMA)指令,可以将乘法和加法操作合并为一条指令,从而提高性能。这可以通过编译器选项 fmad=true 或使用内部函数(intrinsics)如 __fmaf_rn 来实现。
优化技巧:在可能的情况下使用融合乘加指令。
霍纳方法会产生很少的指令级并行性(ILP),因为它创建了一个长的依赖指令链。当需要 ILP 时,可以采用 Estrin 方案。Estrin 方案可以用于添加指令序列化,而开销很小,因为它将多项式分解为两个子多项式,然后用霍纳方法分别求值。
下表比较了在不同并行度(通过每个 SM 的线程块数来体现)下霍nor方法和Estrin方案的性能。
- 在高并行度(8个线程块/SM)下,两种方案性能相近。
- 在低并行度(1个线程块/SM)下,Estrin方案(192 us)由于其更高的ILP,性能优于霍纳方法(326 us)。
优化技巧:如果需要指令级并行度(ILP),请打破依赖指令链。
矩阵乘法是深度学习中的核心运算。在全连接层中,输入(inputs)乘以权重(weights)得到输出(outputs),这本质上是一个矩阵乘法操作。批处理维度(batch dimension)使输入成为一个二维矩阵。
Tensor Core 是专门用于矩阵乘法运算的硬件流水线。
HMMA 1688 是一个 FP16 16x8x8 的 Tensor Core 操作示例。
HMMA 1688 指令可计算 128 个结果(相当于 1024 次 FMA 浮点运算)。D = A x B + C。矩阵维度:
整个 Warp(32个线程)参与计算。
Tensor Core 随着 NVIDIA GPU 架构的演进而不断发展,在数据中心和消费级 GPU 上均有部署。
Ampere (sm_80, sm_86/87):
Hopper (sm_90a): 数据中心 GPU。引入 FP8 支持和共享内存(Shmem)输入,采用 4-warp-wide 操作,性能提升至每个 SM 每时钟周期 4096 FP16 Flops。
注:峰值 FP16 Flops 数据均基于非稀疏计算。
Tensor Core 指令通常是不可移植的,除非它们基于 Ampere SuperMMA API。为了更好的硬件抽象,推荐使用更高层次的 API,如 CUTLASS。
指令集演进:
mma PTXPTX wgmmaPTX TCgen5mma PTX 指令扩展以支持新的数据类型。开发者可以在不同抽象层次上使用 Tensor Core,从低级 API 到高级 API。
低级 API (Low level APIs)
wmma (warp-level), mma (warp-wide), wgmma (warp-group, Hopper), tcgen05.mma (Blackwell)。高级 API (High level APIs)
在大多数情况下,CUTLASS 应该可以满足需求,并提供硬件抽象和出色的性能。
查看相关的 CUTLASS 演讲:
如果您想使用 PTX,请查阅 PTX 文档或 CUTLASS 代码。
mma 指令wgmma 指令tcgen05.mma 指令通用 CUDA
CUDA Python
CUDA C++
开发者工具
与专家交流
多 GPU 编程
性能优化
更多信息请访问:http://nvidia.com/gtc/sessions/cuda-developer