Why GPUs for Deep Learning?GPUs and CPUs represent fundamentally different hardware philosophies. Understanding this difference explains why LLM training is 100–1000× faster on GPUs.philosophyfəˈlɒsəfin. 哲学哲学体系思想体系人生哲学生活工作准则某一知识或经验领域的理论fundamentallyˌfʌndəˈmentəliadv. 根本地基础地用来加强语气从根本上说multiplicationˌmʌltɪplɪˈkeɪʃ(ə)nn. 乘法大量增加embarrassinglyɪmˈbærəsɪŋliadv. 使人尴尬地令人难堪地teraˈterəpref. 太拉兆兆万亿tera-embarrassingly 此处不翻译为 “尴尬地”是英文固定修辞形容任务拆分极其简单、各子任务几乎无数据交互、几乎不需要同步通信简单到让人觉得 “这并行太容易都有点不好意思”因此简化为「易并行」GPUs 在吞吐量方面进行了优化——它们能够并行执行数千个线程每个线程都负责执行简单的任务。现代 GPU 拥有数千个“核心”执行单元这些核心被组合成流式多处理器。中央处理器CPU是专门为应对延迟latency问题而优化的它们能够尽可能快速地执行多个线程拥有较大的缓存、分支预测功能以及乱序执行能力。现代 CPU 拥有 8 到 96 个coreGPUs 在吞吐量throughput方面进行了优化——它们能够并行执行数千个线程每个线程都负责执行简单的任务。现代 GPU 拥有数千个“core”执行单元这些 core 被组合成流式多处理器Streaming Multiprocessors流式多处理器在深度学习的工作负载中矩阵乘法matrix multiplication运算占据主导地位对O(n^2)个数据执行O(n^3)次运算这种运算方式具有出色的并行性embarrassingly parallel。对于一个拥有 700 亿个参数70B的模型单个变换器的前向传播forward pass过程每生成一个 token 需要约 140 TFLOP 的计算资源140 万亿次浮点运算——这完全符合 GPU 的处理能力FLOPFloating Point Operation浮点运算一次加减乘除算 1 个 FLOPTFLOPTera FLOP万亿次浮点运算1 TFLOP 10^12FLOPs注意性能单位常写 FLOPS带 s每秒这里 TFLOP 是总运算量不带 sNVIDIA GPU Microarchitecture Generations微架构怎么选择训练 70B 以上的模型使用 H100/B200 节点并配备 NVLink 连接需要快速的连接以实现张量并行 tensor parallelism 处理。每个实例至少应配备 8 个 H100 节点推理过程对延迟敏感latency-sensitive对于高带宽BW需求建议使用 H100/H200 系列芯片对于内存限制memory-bound严重的应用场景如需要大量 KV 缓存MI300X 系列芯片更为合适在 7B 到 13B 的范围内进行微调A100-80GB 的型号具有成本效益cost-effective。使用单块 GPU 并结合 LoRA 技术来实现优化预算对于 7B 模型上的 LoRA A100-40GB 甚至 A1024GB的配置已经足够满足需求了Tensor Core张量核心专门优化矩阵乘法LLM 注意力层、全连接层核心运算是大模型训练提速的关键硬件单元每代升级支持更低精度FP16→BF16→FP8→FP4算力成倍提升Transformer EngineTransformer 专用硬件引擎Hopper 初代、Blackwell 第二代针对 LLM 的注意力、归一化、层 Norm 做硬件级流水线优化大幅降低per token计算开销NVLinkNVIDIA 多卡高速互联总线用于多卡分布式训练时交换模型梯度、KV 缓存带宽越高多卡集群通信瓶颈越小HBM高带宽堆叠显存相比普通 GDDR 显存带宽高出数倍LLM 推理 / 训练需要超大带宽吞吐 KV 缓存是大模型刚需TFTFLOP万亿次浮点运算表格中BF16 TF代表该卡 BF16 精度下的峰值算力区分 TFLOP总运算量与 TFLOPS每秒算力硬件指标MIGAmpere 架构引入的 GPU 多实例虚拟化单张 A100 可切分为多张小 GPU用于小推理任务资源复用RLHF基于人类反馈的强化学习大模型对齐流程需要大量显存存放奖励模型、策略模型硬件代际核心趋势总结显存容量持续暴涨从 V100 32GB → B200 192GB单卡可容纳更大参数模型降低多卡通信成本低精度精度迭代FP16 → BF16 → FP8 → FP4用更低位宽换取更高算力与显存吞吐专用 LLM 硬件单元从通用 Tensor Core 升级为独立 Transformer Engine针对性优化大模型计算互联与显存带宽翻倍增长B200 NVLink 带宽、HBM 带宽相比 H100 直接翻倍支撑万亿参数超大规模集群训练硬件代际芯片厂商按一套全新底层架构推出的一整代硬件产品线一代就是一个完整技术周期包含统一设计的核心电路、配套显存、专用加速单元、互联接口。每更新一代会同步升级计算核心、存储、互联、精度指令集性能、能效、功能相比上一代有跨越式提升跨代代际迭代质变架构名称更换底层电路重构新增关键硬件单元算力 / 带宽翻倍支持全新精度Ampere → Hopper新增 FP8、Transformer 引擎、HBM3算力从 312TF→990TF属于换代同代小改版量变不算新代际架构不变仅扩容显存、小幅提带宽核心计算单元完全一致H100 (80GB) → H200 (141GB)都是 Hopper 架构属于同代扩容版本不产生新硬件代际HBMHigh Bandwidth Memory 高带宽内存专门给 AI 训练卡、超算 GPU 用的堆叠式高速显存核心解决大模型的「内存墙」瓶颈3D 垂直堆叠把多层 DRAM 芯片像积木一样上下叠在一起用硅通孔 TSV 垂直互联替代游戏卡平铺在 GPU 四周的 GDDR 方案超宽总线单栈位宽 1024bit 起步普通游戏卡 GDDR 仅 256/384bit靠超大位宽实现爆炸级带宽2.5D 先进封装显存堆叠块通过中介层直接贴紧 GPU 核心走线极短、延迟低、功耗更优FP16 / BF16 / INT8 / FP8 / FP4位宽数字占用多少个 bit位宽越小单个数字占用显存越少、单次硬件可并行计算数量越多 → 算力越高、显存吞吐越大浮点 FP分符号位、指数位、尾数位适合 AI 模型权重 / 激活值数值范围跨度大整数 INT全是整数刻度无指数推理场景专用训练几乎不用FP16半精度浮点16bit存储拆分1 符号位 5 指数位 10 尾数位代表架构Pascal2016首次支持优缺点✅ 相比 32bit FP32 显存占用减半算力翻倍❌ 指数范围极小数值稍大就溢出训练不稳定需要额外 FP32 主权重保存定位初代混合精度训练基础Volta Tensor Core 原生加速早期大模型训练标配BF16Brain Float 1616bit存储拆分1 符号位 8 指数位 7 尾数位代表架构Ampere2020新增 BF16 Tensor Core和 FP16 核心差异指数位从 5 扩充到 8和 FP32 指数范围完全一致训练不会溢出不用额外保存 FP32 权重代价是尾数精度略降但神经网络对尾数不敏感定位现在 LLM 训练标准精度平衡稳定与显存开销INT88 位整数整数量化非浮点存储纯 8bit 整数无指数通过缩放因子映射浮点区间代表架构Turing2018首次 INT8 推理加速优缺点✅ 位宽直接砍到 8bit显存占用仅 FP32 的 1/4推理算力极高❌ 无指数数值动态范围差必须做量化校准训练完全不能用仅推理定位线上推理轻量化方案适合低延迟、小显存场景不参与训练迭代链FP88 位浮点Hopper 时代核心分两种子格式E4M34 指数 3 尾数精度高尾数越高精度越高但是对应的指数部分所表示的最大范围就越小适合激活值E5M25 指数 2 尾数范围大适合权重代表架构Hopper H1002022新增 FP8 Tensor Core突破点同时兼顾浮点的大范围和8bit 低位宽训练、推理都能用显存占用是 BF16 的一半算力直接翻倍配合 Transformer Engine 大幅降低per token计算量定位当前 70B / 万亿参数大模型训练主流精度FP44 位浮点Blackwell 下一代4bit 超低位宽浮点指数 尾数压缩到仅 4bit代表架构Blackwell B2002024支持 FP4优缺点✅ 显存占用仅 BF16 1/4算力再翻倍KV 缓存、权重极致压缩❌ 精度损失更大需要配套硬件引擎做误差补偿定位下一代超大模型推理、低开销微调专用精度小节训练浮点迭代FP16 → BF16 → FP8 → FP4解决训练显存不足、算力低的痛点全程浮点保证训练收敛稳定Pascal/VoltaFP16能用但易溢出AmpereBF16训练标准解决溢出HopperFP88bit 浮点训练推理通用算力翻倍BlackwellFP44bit 超低显存占用极致吞吐推理整数分支 INT8不属于训练迭代链是独立轻量化方案Turing 架构单独推出只做推理不参与训练和上面浮点路线并行互不替代BF16 - Brain Float 16 和 FP16 区别BF16 优先推荐大模型训练代码简单、不用缩放、数值稳Hopper/Blackwell GPU 原生硬件加速FP16 老旧 A100 及以下常用动态范围小必须加 GradScaler梯度裁剪有严格调用顺序容易出现 NaN两者激活显存均减半权重显存看是否保留 FP32 主副本Grad underflow 梯度下溢梯度数值过小FP16 无法表示直接归零导致参数不更新Overflow / NaN 溢出数值超出 FP16 表示范围产生非法计算值unscale_ 解缩放把放大后的梯度还原为真实梯度master FP32 copy优化器保存的高精度权重副本用于稳定参数更新brainbreɪnn. 脑非正式智力智慧聪明的人智者非正式电脑供食用的动物脑髓BF16Brain Float 16仅需用 autocast(dtypetorch.bfloat16) 上下文包装即可不需要梯度缩放器 GradScaler代码更简洁数值稳定性更强底层原因BF16 指数位和 FP32 一致动态范围大训练时几乎不会出现梯度下溢 / 溢出无需缩放补偿FP16标准半精度 Float16必须搭配 GradScaler 梯度缩放器用来防止梯度下溢Scaler 会动态维护一个放大系数前向 / 反向时把梯度放大避免极小梯度被 16 位精度截断变成 0如果检测到溢出出现 NaN 非法值会跳过本轮参数更新并降低缩放系数防止后续持续溢出FP16 梯度裁剪gradient clipping强制规范如果同时使用梯度裁剪 clip_grad_norm_必须先执行 scaler.unscale_(optimizer) 再裁剪缩放后的梯度scaled gradient数值被放大直接裁剪会使用错误阈值先解缩放还原真实梯度再做裁剪阈值才符合预期显存节省逻辑两者都把激活值用 16bit 存储能大幅降低激活显存占用权重显存占用分两种情况保留 FP32 主权重副本权重显存几乎无节省全程 16bit 权重权重显存减半什么是权重神经网络里的权重weight/parameter就是模型每一层的参数矩阵比如 Transformer 的 QKV 矩阵、FFN 权重、LN 偏置、词嵌入表全部统称权重模型训练时这些参数会长期保存在 GPU 显存里这部分占用就是权重显存占用工业使用FP32、FP16、BF16情况模式 A保留 FP32 主权重副本绝大多数训练框架默认显存里存两份权重高精度主副本FP3232 位浮点用于优化器更新、保证收敛稳定计算副本FP16/BF1616 位用于前向 / 反向矩阵乘加速显存开销FP32 一份 16bit 一份总占用 ≈ 原始 FP32 单份的 1.5 倍结论开启半精度后权重显存几乎不节省甚至更费显存模式 B全程只用 16bit 权重仅推理 / 部分极简微调不保存 FP32 高精度副本权重全程只用 FP16/BF16 16bit 存储权重显存直接减半缺点优化器更新精度不足大模型训练容易发散、loss 震荡训练场景极少用GPU Internal Architecture – The Streaming Multiprocessor (SM)warpwɔːpv.使扭曲使变形使反常使有偏见用绞船索牵曳船纺织用语把纱排列成经淤灌土地n.织布中的经纱经线变形扭曲小说中或假想的空间旅行靠扭曲时空实现A GPU is organized as an array of Streaming Multiprocessors (SMs), each of which is an independent processor with its own register file, shared memory, and execution units. Understanding SMs is key to understanding GPU performance.GPU 是由一系列流式多处理器 SMs 组成的 array每个流式多处理器都是一个独立的处理器processor拥有自己的寄存器文件、共享内存以及执行单元。理解流式多处理器是掌握 GPU 性能的关键。左图单颗 A100 流式多处理器SM内部结构 —— 包含 64 个 FP32 CUDA core、4 个张量coreTensor Core、4 个线程束Warp Scheduler调度器、256KB 寄存器堆、192KB 可配置共享内存 / L1 缓存。左侧边缘标注的带宽数值直观体现数据从寄存器到 HBM 显存的带宽断崖式下跌右图完整 A100 芯片整体架构总计 108 个 SM共享 40MB 二级缓存L2 Cache搭载 80GB HBM2e 显存把 GPU 比作工厂SM 独立生产车间Tensor Core 是车间里专门做大件货物的高速机床寄存器 机床手边工作台存取最快共享内存 车间公共物料台L2 缓存 全厂中央仓库HBM 显存 工厂外远程大仓库往返运输速度极慢内存墙Memory Wall问题就是机床速度极快但原料要频繁跑远路调取机床大量时间闲置等待物料一个SM从上到下详解SM 是 GPU 最小独立计算单元所有模型运算都在 SM 内完成64 FP32 CUDA Cores 分为 4 组 ×16 个 FP32 核心负责通用浮点运算、基础矩阵计算FP32 即 32 位全精度浮点是模型梯度更新、数值校验的基础单元4 Tensor CoresTC0~TC34×4×4 MMA 矩阵乘累加硬件单元专门加速 LLM 注意力、全连接层的大规模矩阵乘法A100 TC 原生支持 FP16/BF16 混合精度是大模型训练提速核心硬件32 Load/Store Units 加载/存储单元负责数据读写、显存 / 缓存数据搬运4 SFU special function unit特殊函数单元计算激活函数tanh、gelu、log、exp 等大模型前向推理必备GPC Graphics Processing Cluster图形处理簇整片 A100 包含 7 个 GPC每个 GPC 内排布多组 SM全卡合计108 个 SM张量Core总量108 SM × 4 TC 432 个 Tensor Core决定整卡混合精度峰值算力40MB L2 Cache 全局二级缓存所有 SM 统一共享缓解 SM 与 HBM 之间的数据吞吐压力减少访问慢速 HBM 的次数HBM2e 显存子系统80GB 容量总双向带宽 2 TB/s分为 5 个显存堆叠栈存放全部大模型权重、梯度、KV 缓存是性能瓶颈 “内存墙”NVLink 高速互联12 条链路单链路 50GB/s总互联带宽 600GB/s多卡分布式训练时传输梯度、参数用于 LLM 多卡集群并行内存墙Memory Wall图左侧竖排带宽数值直观展示带宽断崖图中那个是弄反了应该是 19TB/s 100TB/s寄存器读写19 TB/sSM 共享内存 / L1100 TB/sL2 缓存2 TB/sHBM 全局显存2 TB/s数据越靠近计算核心带宽越高、延迟越低一旦需要读取 HBM 显存带宽直接下跌一个数量级显存读写会拖慢全部计算单元这就是 AI 领域常说的内存墙瓶颈大模型优化算子融合、KV Cache 量化、重计算、共享内存分块全部目标就是尽可能减少访问 HBM 显存的次数匹配之前提到的 140 TFLOP per token 算力开销HBM High Bandwidth Memory 高带宽显存2 第二代标准 HBM2e enhanced增强版。HBM2e 第二代增强型高带宽内存完整 HBM 迭代路线HBM1 → HBM2 → HBM2e → HBM3 → HBM3e → HBM4HBM2初代第二代标准带宽、频率较低用于 V100HBM2eHBM2 的超频增强改版A100Ampere专属配套显存CUDA CoresCUDA 核心Compute Unified Device Architecture 统一计算设备架构Scalar ALUs for FP32/INT32 operations. 64 per SM on A100. Used for element-wise ops, reductions, and non-matrix operations.标量算术逻辑单元负责 FP32/INT32 通用运算。A100 每个 SM 内置 64 个 CUDA 核心用于逐元素运算、规约、以及所有非矩阵类计算。标量单元一次只处理 1 个数值不擅长大规模矩阵乘法LLM 场景用途激活函数逐点计算、梯度累加、归一化等细碎运算和 Tensor Core 形成互补矩阵交给 TC零散标量计算交给 CUDA CoreTensor Cores张量核心Specialized matrix-multiply-accumulate (MMA) units. Each performs a 4×4×4 fused multiply-add per cycle. 4 per SM on A100, delivering 16× throughput over CUDA cores for supported precisions.专用矩阵乘累加MMAmatrix-multiply-accumulate硬件单元。每个 TC 每个时钟周期可完成一次 4×4×4 融合乘加运算A100 每个 SM 配备 4 个张量Core在支持精度下计算吞吐是普通 CUDA 核心的 16 倍。MMA 矩阵乘法 累加是 Transformer 注意力层、全连接层的核心计算4×4×4同时完成 4×4 矩阵 × 4×4 矩阵 4×4 偏置矩阵16 倍吞吐低精度矩阵运算场景算力碾压普通 CUDA 核心是大模型训练提速的关键硬件。Register File寄存器堆Fastest storage (1 cycle latency). Shared among all active threads. Spilling to L1 causes significant slowdown.GPU 内速度最快的存储仅 1 个时钟周期延迟由 SM 内所有活跃线程共享。一旦数据溢出转存至 L1 缓存会造成严重性能下降延迟最低计算单元取数无需等待是算力吞吐的基础Spill寄存器溢出线程用到的变量超出寄存器容量时被迫写到 L1 缓存带宽、延迟大幅恶化是算子优化重点规避问题Shared Memory / L1共享内存 / 一级缓存On-chip SRAM explicitly managed by the programmer. The key to Flash Attention’s performance (tiles fit entirely in shared memory).片上静态内存需要开发者手动管理是 Flash Attention 高性能的核心保障注意力分块可完整存入共享内存片上 SRAM速度远高于 L2、HBM 显存可自由划分「共享内存」/「L1 缓存」比例Flash Attention 核心原理将注意力矩阵切分为小块tile全部放进高速共享内存读写避免反复访问慢速 HBM 显存完美缓解内存墙瓶颈Warp Schedulers线程束调度器Each SM has 4 warp schedulers (A100). A warp 32 threads executing in lockstep (SIMT model). Schedulers hide memory latency by switching between warps.A100 每个 SM 配备 4 个线程束调度器一个 Warp线程束包含 32 个同步执行的线程遵循 SIMT 单指令多线程模型。调度器通过切换不同 Warp掩盖显存读写延迟SIMT同一 Warp 内 32 线程执行完全相同指令延迟隐藏机制某一组 Warp 等待显存数据时调度器立刻切换另一组就绪 Warp 继续计算避免 CUDA 核心空转调度器A100 相比 V100 调度能力翻倍可并发调度更多线程束硬件利用率更高组件层级性能速记从快到慢寄存器堆 共享内存 / L1 缓存 L2 缓存 HBM 全局显存Flash Attention、算子融合、寄存器复用等所有大模型优化目标都是尽可能把计算数据留在前两级高速存储减少访问 HBMGPU 采用单指令多线程执行模式SIMTSingle InstructionMultiple Threads。在一个 warp 单元一个线程束中包含 32 个线程所有线程执行相同的指令但操作的是不同的数据。当线程执行不同路径时例如在 if/else 条件语句中两条路径会被串行处理——这种现象被称为 warp 发散warp divergence。因此GPU 内核需要尽量减少分支操作minimize branching。对于大语言模型的工作负载来说主要的运算操作如 GEMM、注意力机制、softmax 等在多个线程间具有统一的uniform控制流程这使得它们非常适合采用 SIMT 执行模式一张 A100 里最小执行单位是 Warp线程束固定捆绑 32 个线程同一个 Warp 里32 条线程必须同步执行完全相同的指令步调完全一致lockstep一个 Warp 内 32 个线程一半进if、一半进else两条路径逻辑不一样if(threadIdx.x%20){// 偶数线程走这条分支大量矩阵计算}else{// 奇数线程走这条分支简单加法}但硬件不允许 32 个线程同时跑两套不同代码只能串行执行第一轮屏蔽所有奇数线程只让偶数线程执行 if 分支代码第二轮屏蔽所有偶数线程只让奇数线程执行 else 分支代码原本 1 轮就能跑完的任务硬生生拆成 2 轮串行执行算力直接折半这个「同一 Warp 线程分裂到不同代码分支、被迫串行执行」的现象就叫 warp divergence 线程束发散把一个 Warp 当成 32 人小组所有人必须同步做同一个动作正常无发散所有人都做加法一步完成出现 if/else 发散一半人要扫地、一半人要擦桌子。规则不允许同时干两件事只能先全体待命一半人扫地再全体待命另一半擦桌子。中间大量等待时间整体效率暴跌GPU Memory Hierarchy and Bandwidth每个层级的速度大约都是上一层速度的 10 倍慢而容量规模则大 100 到 1000 倍。A100 具有 312 TFLOP/s 的 BF16 张量计算性能但 HBM 带宽仅为 2 TB/s这意味着从 HBM 中读取每一字节的数据后在下一字节到达之前可以执行约 156 次浮点运算。如果你内核每字节能处理的浮点运算次数少于 156 次那么你的计算系统就存在内存限制问题计算单元会处于空闲状态无法及时处理数据Attention is Memory-Bound; FFN is Compute-BoundAttention 自注意力层访存受限 memory-boundmemory-bound访存瓶颈定义GPU 算力充足但需要频繁读写 HBM 显存数据搬运速度跟不上计算速度Tensor Core 大部分时间在等数据算力被显存带宽锁死长文本痛点n 越大n^2显存占用暴涨KV 缓存、打分矩阵反复读写内存墙问题极度严重FFN / MLP 前馈网络计算受限 compute-boundGEMM通用矩阵乘法是 Tensor Core 最擅长的运算arithmetic intensity计算访存比极高读取少量权重就能完成海量浮点计算compute-bound计算瓶颈定义显存读写压力很小GPU 计算单元算力打满性能上限由 Tensor Core、CUDA Core 的峰值算力决定不存在大量等待数据的情况对应两类优化的适用场景Flash Attention显存优化算法只能加速 Attention 层对 FFN 几乎无提升量化压缩减小权重存储体积对 FFN 的提速效果远大于 AttentionFlash Attention 解决的是 Attention 的访存爆炸把 n×n 打分矩阵切小块存入高速共享内存规避慢速 HBM 读写但 FFN 本身不缺显存带宽因此优化无效量化压缩权重FFN 权重体积巨大d×4d升维、降维低位宽量化能大幅减少权重读写量而 Attention 的瓶颈是n^2序列缓存权重占比低量化收益微弱