Workflow
张量核心架构演进
icon
搜索文档
NVIDIA Tensor Core 从 Volta 到 Blackwell 的演进
傅里叶的猫· 2025-06-23 23:18
性能基本原理 - 阿姆达尔定律指出,对于固定问题规模,通过增加计算资源实现的最大加速比受限于串行部分,即使并行资源无限增加,加速比也只能趋近于1−S,因为串行部分的执行时间无法通过并行化减少 [3][4] - 强缩放指的是在固定问题规模下,通过增加计算资源来缩短执行时间,其加速比由阿姆达尔定律量化 [6] - 弱缩放则是在保持执行时间不变的情况下,通过增加计算资源来处理更大的问题,问题规模和资源同时按比例增加,以维持时间不变 [6] - 数据移动在性能优化中是一个关键瓶颈,被称为 "cardinal sin",现代 DRAM 单元的操作时间为数十纳秒,而晶体管开关速度在亚纳秒级别,这种速度差异导致数据移动本质上更慢 [8] 张量核心架构演进 - Volta 架构作为首个引入张量核心的里程碑,其设计源于深度学习对矩阵乘法硬件加速的需求,2017 年推出的 Volta 架构在 Tesla V100 GPU 中集成张量核心,旨在解决传统 GPU 执行矩阵乘法时指令功耗与计算功耗的失衡问题 [9] - Turing 架构于 Volta 之后推出,其第二代张量核心在 Volta 基础上增加了 INT8 和 INT4 精度支持,进一步拓展了低精度计算能力,同时通过引入深度学习超采样(DLSS)技术,将深度学习应用于游戏图形领域 [10] - Ampere 架构带来了异步数据复制技术,这一创新允许数据直接从全局内存异步加载到共享内存,绕开寄存器中转,解决了 Volta 时代数据加载与 MMA 指令竞争寄存器资源的问题 [11] - Hopper 架构进一步深化了并行计算的层次设计,新增线程块集群(Thread Block Cluster),将多个 SM 分组为图形处理集群(GPC),允许跨 SM 的数据共享与低延迟通信 [12] - Blackwell 架构作为最新一代,针对寄存器压力问题引入张量内存(TMEM),每个 SM 配备 256KB 的 TMEM,以 warpgroup 为单位访问 [13] 结构化稀疏性 - Ampere 架构推出了 2:4 结构化稀疏性,其核心在于对权重矩阵进行修剪,使每 4 个元素中 2 个为零,通过压缩非零元素并利用元数据索引记录位置,理论上可将张量核心吞吐量翻倍 [14] - Blackwell 架构则针对 NVFP4 数据类型引入了 4:8 结构化稀疏性,该模式将 8 个元素划分为 4 对连续元素,要求其中 2 对为非零值、2 对为零 [15] 张量核心规模与内存演进 - 从 Volta 到 Blackwell,张量核心的计算规模呈指数级增长,而内存层次结构则通过容量扩展与架构优化持续适配计算需求 [16] - Volta 架构作为张量核心的起点,单个 SM 配备 8 个张量核心,可实现 1024 FLOP / 周期的 F16 计算能力,支持 m8n8k4 的 MMA 形状 [17] - Ampere 将单 SM 的张量核心计算能力翻倍至 2048 FLOP / 周期,MMA 形状扩展为 m16n8k16 [17] - Hopper 进一步提升至 4096 FLOP / 周期(F16),并引入 F8 格式使计算能力达到 8192 FLOP / 周期,MMA 形状支持 m64n256k16 的更大规模 [17] - Blackwell 则实现了 F16 计算能力 8192 FLOP / 周期、F8 达 16384 FLOP / 周期、F4 达 32768 FLOP / 周期 [17] MMA 指令异步性 - Volta 架构作为初代张量核心,其 MMA 指令采用 warp-scoped 同步执行模式,需 8 线程 quadpair 协作完成 8x8x4 矩阵运算 [20] - Ampere 架构首次引入异步数据复制技术,允许数据从全局内存直接加载至共享内存,绕过寄存器中转 [20] - Hopper 架构实现了 MMA 指令的根本性突破,推出 warpgroup-level 异步 MMA(wgmma),支持 4 个 warp 组成的 warpgroup 协作执行更大规模矩阵运算 [22] - Blackwell 架构将 MMA 异步性推向极致,第五代张量核心的 tcgen05.mma 指令具备单线程语义,无需 warpgroup 协作即可发起 MMA 操作 [23] 数据类型精度演进 - Volta 架构作为张量核心的起点,仅支持 FP16 半精度输入与 FP32 单精度累加 [25] - Turing 架构在此基础上新增 INT8 和 INT4 整数精度支持,首次将低精度整数计算引入张量核心 [25] - Ampere 架构进一步拓展数据类型范围,引入 BF16(脑浮点格式),其 8 位指数与 7 位尾数的设计,在保持与 FP32 相同动态范围的同时,将存储成本减半 [25] - Hopper 架构标志着低精度浮点类型的重大突破,首次引入 FP8 格式(E4M3 和 E5M2),通过 4 位指数与 3 位或 2 位尾数实现更低精度计算 [26] - Blackwell 架构将精度降低推向极致,新增 MXFP 系列微缩放浮点格式(MXFP8、MXFP6、MXFP4),并推出自研的 NVFP4 格式 [26] 编程模型演进 - 早期 CUDA 编程模型遵循高线程占用率原则,通过将多个 CTA 分配至单个 SM,利用线程上下文切换隐藏内存访问延迟 [28] - Ampere 架构首次推出异步数据复制指令,允许线程直接将数据从全局内存加载至共享内存,无需经过寄存器中转 [29] - Hopper 架构进一步深化异步能力,新增线程块集群(Thread Block Cluster),将多个 SM 分组为 GPC,通过协作组 API 暴露硬件执行单元 [29] - Blackwell 架构将异步执行推向全栈支持,第五代张量核心的 tcgen05.mma 指令具备单线程语义,无需 warp 协作即可发起 MMA 操作 [30]