来源： Youtube | Tom Vman | 2024-02-05 分类： 其他 原文发表： Feb 05, 2024 纪要生成： 2026-03-05
本期嘉宾Tom Vman是技术博客learn apparat作者，经营独立咨询公司，即将加入Lightning AI团队。本次是CUDA模式读书会第四讲，围绕《大规模并行程序设计》教材第4、5章内容，讲解GPU计算与内存基础，帮助开发者编写高性能CUDA核函数。内容涵盖GPU硬件架构、线程调度规则、内存层级、分块优化等实操知识点，适合PyTorch等框架的底层性能优化从业者学习。
本节重点 - 对比CPU与GPU核心架构差异，介绍流多处理器（SM）的硬件构成 - 说明消费级GPU与数据中心GPU的核心差异，指出FP64算力的性能短板 - 讲解RTX 3090等主流消费级GPU的SM数量、缓存配置参数
详细精要
近年Volta架构之后的GPU已为每个线程配置独立程序计数器，修正了旧教材的过时表述
消费级GPU硬件参数特征：以RTX 3090为例，介绍主流消费级GPU的硬件配置细节
若代码意外使用64位浮点常量，会自动切换到FP64计算模式，导致性能大幅下降
单SM硬件构成：单个RTX 3090的SM分为4个独立子单元，具备独立的寄存器、调度模块
💬 精华片段（中文）
"the fp64 rate is 1/64 of the fp32 T flop right um and this is because it just doesn't it has two compute unit units for the for 64bit floating points just to make the programs run at all but they're not going to be fast"
本节重点 - 讲解线程块与SM的绑定规则，说明线程块大小选择对GPU占用率的影响 - 解释线程束的构成与线性编号规则，对比Nvidia与AMD的线程组术语差异 - 介绍线程束分歧的产生原因与性能损耗，说明新架构的优化逻辑
详细精要
多线程块调度时，SM以线程束为单位切换执行，寄存器无需清空恢复，切换开销可忽略
线程束的构成与编号规则：每个线程块会被自动拆分为多个32线程的线程束（Warp），是GPU调度的最小单位
开发者可通过Shuffle指令验证线程束编组规则，避免因索引假设错误导致的性能问题
线程束分歧的产生与影响：同一线程束内线程执行不同分支或不同循环迭代次数时，会产生分歧导致硬件闲置
💬 精华片段（中文）
"if you have thread Divergence this part here there's part of the GPU sitting idle whenever there's one of those W spots and so uh you want to avoid thread warp Divergence if you can"
本节重点 - 讲解GPU占用率的定义，说明影响占用率的核心约束条件 - 介绍寄存器、共享内存使用对线程调度的影响，以及launch_bounds编译指令的作用 - 说明PyTorch与CUDA API的硬件参数查询方法，避免硬编码Magic Number
详细精要
应避免在消费级GPU上使用FP64与INT64计算，嘉宾曾因误将索引改为INT64导致PyTorch二项式核函数性能大幅下降
资源使用对占用率的约束：寄存器、共享内存的使用量会直接限制单个SM可承载的线程数量
早期使用Excel表格计算占用率，现在可通过NVIDIA Nsight Compute工具直接分析
硬件参数查询方法：所有硬件配置参数均可通过API查询，无需硬编码
💬 精华片段（中文）
"previously there was an Excel sheet for the occupancy calculation now this is part of inside compute and I uh Venture that we're going to see some of that in the profiling uh sessions that we're going to take"
本节重点 - 分析PyTorch程序的时间分布，说明GPU利用率不足的常见排查方向 - 讲解核函数融合的性能提升原理，介绍PyTorch历代融合器的发展与适用场景 - 对比GPU各级内存的带宽与容量差异，说明减少全局内存访问的核心优化思路 - 说明核函数融合带来的数值精度差异原因，明确精度误差的正常范围
详细精要
GPU计算耗时分为核函数启动固定开销、内存访问耗时、实际计算耗时三个部分
核函数融合的优化原理：核函数融合可减少全局内存的读写次数，是最常用的性能优化手段之一
PyTorch最早的融合器支持逐点算子融合，帮助LSTM性能接近CNN水平；第二代融合器支持规约操作，其中NVIDIA Fuser仍在活跃迭代；Inductor、Triton等方案支持更复杂的算子融合
GPU内存层级差异：GPU采用金字塔式内存层级，从寄存器到全局内存带宽逐步下降、容量逐步上升
嘉宾以近似Gelu算子为例，手写融合核函数的性能比原生PyTorch逐算子实现快7-8倍，甚至略快于PyTorch内置实现
浮点数精度差异说明：核函数融合可能带来微小的数值差异，根源是浮点数加法不满足结合律
💬 精华片段（中文）
"this avoiding accesses to Global memory also as a core ingredient of flash attention"
本节重点 - 讲解理论性能上限（光速）的测算方法，对比实际运行性能判断优化空间 - 介绍屋顶线模型的核心逻辑，帮助开发者区分计算 bound 与内存 bound 核函数 - 说明GPU各级内存的适用场景与声明方式
详细精要
嘉宾测试的优化后核函数实际运行耗时约26微秒，达到理论上限的75%，优化空间极小
屋顶线模型核心逻辑：通过计算强度区分核函数的瓶颈类型，指导优化方向
线程束调度可隐藏内存访问延迟，内存访问与计算可并行执行，因此性能为两者的最大值，形成屋顶形状的性能曲线
GPU各级内存的使用规则：不同类型内存有不同的声明方式与适用场景
💬 精华片段（中文）
"this is jokingly called the speed of light because it's the theoretical limit uh for the speed"
本节重点 - 讲解矩阵乘法的访存特征，说明分块优化的核心逻辑 - 介绍分块矩阵乘法的实现细节，以及线程同步的注意事项 - 说明非对齐尺寸矩阵的填充适配方法，以及分块优化的性能收益
详细精要
分块大小为T时，每个输入元素的全局内存读取次数可降低到n/T次，访存量减少T倍
分块矩阵乘法的实现要点：分块矩阵乘法需要手动管理共享内存加载与线程同步
完成当前分块的计算后需再次同步线程，避免部分线程提前加载下一分块覆盖共享内存数据
非对齐矩阵适配与性能收益：分块优化可适配非对齐尺寸矩阵，带来稳定的性能提升
💬 精华片段（中文）
"the idea here is to uh of tiling is to read these parameters once and put them into shared memory and then try to reuse the same value red from the uh Global memory"
本节重点 - 汇总本次课程的核心知识点，梳理GPU性能优化的核心方向 - 预告下节课内容为合并内存访问优化，介绍扩展练习方向
详细精要
可通过屋顶线模型、理论上限测算定位性能瓶颈，针对性优化
后续内容预告：下节课将讲解合并内存访问的规则与优化方法，进一步提升内存访问效率
💬 精华片段（中文）
"in the next chapter we can see how to organize read and write so they are all consecutive and aligned to Global memory location so we have qualis memory access and we'll see a lot about how this works in detail"
| 术语 | 解释 |
|---|---|
| Streaming Multiprocessor (SM，流多处理器) | GPU的核心计算单元，集成大量ALU、寄存器、调度单元，可并行运行数千个线程，是GPU并行计算的核心硬件模块 |
| Warp（线程束） | Nvidia GPU的最小调度单位，固定为32个线程，同一线程束内的线程共享调度资源，执行相同指令时性能最优 |
| Wavefront（波前） | AMD GPU对线程束的称呼，默认大小为64线程，可通过编译器参数修改为32线程 |
| Warp Divergence（线程束分歧） | 同一线程束内的线程执行不同分支或不同循环迭代次数，导致部分计算单元闲置，性能下降的现象 |
| Occupancy（占用率） | GPU硬件单元的实际利用率，越高说明硬件资源利用越充分，性能通常越好 |
| Roofline Model（屋顶线模型） | 一种性能分析模型，通过计算强度区分核函数是计算 bound 还是内存 bound，指导优化方向 |
| Tiling（分块/切片） | 一种内存优化方法，将大矩阵/大张量拆分为小分块，加载到高速共享内存中复用，减少全局内存访问次数 |
| __syncthreads() | CUDA内置线程同步函数，同步同一个线程块内的所有线程，确保所有线程完成前面的操作后再继续执行 |
| FP32/FP64 | 32位单精度浮点数、64位双精度浮点数，消费级GPU的FP64算力远低于FP32，应尽量避免在消费级GPU上使用FP64 |
| Shared Memory（共享内存） | GPU的高速片上内存，每个SM独立配备，带宽是全局内存的10倍以上，容量较小，由同一个线程块内的所有线程共享 |
| Global Memory（全局内存） | 通常所说的GPU显存，容量大但带宽较低，是核函数输入输出的主要存储介质 |
| Kernel Fusion（核函数融合） | 将多个独立的小核函数合并为一个核函数，减少中间结果的全局内存读写，提升性能的优化方法 |
| launch_bounds | CUDA编译指令，用于告知编译器核函数的线程块大小、最大线程块数量等参数，辅助编译器优化寄存器分配 |