Tensor Core架构演进 - Tensor Core是推动GPU计算能力远超摩尔定律的核心技术，已成为现代AI和机器学习的基石[1] - 从Volta到Blackwell共经历五代架构演进：Volta(第一代)、Turing(第二代)、Ampere(第三代)、Hopper(第四代)、Blackwell(第五代)[11] - 每代架构在MMA(矩阵乘加)指令执行方式、数据精度支持、内存层次结构等方面均有重大创新[11][18][30][39][46] 性能工程原理 - 阿姆达尔定律量化了并行计算的最大加速比，性能提升受限于串行部分执行时间[5] - 强扩展通过增加计算资源解决固定规模问题，弱扩展通过增加资源解决更大规模问题[6] - 数据移动成本远高于计算成本，现代DRAM速度比晶体管开关速度慢两个数量级，形成"内存墙"[10] 编程模型演变 - PTX编程模型采用线程网格-CTA-Warp的三级线程层次结构，对应寄存器-共享内存-全局内存的内存层次[13][14] - SIMT执行模式以Warp(32线程)为单位发出指令，与SIMD不同在于指定单线程行为而非向量宽度[15] - SASS是PTX底层指令集，但文档不完善因NVIDIA对竞争对手保密[17] 各代Tensor Core特性 Volta(第一代) - 引入HMMA指令执行8x8x4矩阵乘法，需8线程四对协作完成[22][25] - 支持FP16输入/FP32累积，符合混合精度训练需求[26] - 每个SM含8个Tensor Core，每周期1024 FLOP[22] Turing(第二代) - 增加INT8/INT4精度支持，引入Warp级同步MMA[27] - 首次将深度学习应用于游戏图形(DLSS技术)[27] Ampere(第三代) - 引入异步数据复制，直接从全局内存到共享内存，缓解寄存器压力[29] - Warp级同步MMA指令，完整32线程参与运算，每SM每周期2048 FLOP(Volta两倍)[30] - 支持BF16格式，提供FP32级别动态范围且无需损失缩放[32] Hopper(第四代) - 新增线程块集群概念，CTA可跨SM协作访问分布式共享内存[33] - 引入张量内存加速器(TMA)，批量异步复制全局内存到共享内存[35] - Warpgroup级异步MMA(wgmma)，4个Warp(128线程)协作执行更大规模矩阵运算[39] - 支持8位浮点(E4M3/E5M2)和22位定点累加[41] Blackwell(第五代) - 新增Tensor Memory(TMEM)专用存储，256KB容量/SM，更靠近计算单元[43] - 第五代MMA指令完全脱离寄存器，操作数驻留共享内存/TMEM[46] - 支持CTA对级MMA(MMA.2SM)，两个SM协作执行[45][49] - 引入MXFP8/6/4和NVFP4等微缩放浮点格式[51][52] 架构演进趋势 - Tensor Core规模扩展速度远超数量增加，MMA形状从Volta的8x8x4扩大到Blackwell的256x256x16[59][60] - 共享内存容量持续增加(Volta 96KB→Blackwell 228KB/SM)，寄存器文件保持256KB[64][65] - 操作数存储位置从寄存器逐步转向共享内存/TMEM，提升数据局部性[67] - MMA指令从同步逐步转向异步执行，提高流水线效率[69][71] - 数据类型持续向低精度发展，从FP16到4位格式，同时缩减高精度支持[73][74] 结构化稀疏性 - Ampere引入2:4稀疏模式(每4元素含2零)，理论可双倍提升吞吐量[54] - Blackwell为NVFP4引入4:8成对稀疏模式，要求更严格[57] - 实际应用中因剪枝难度和优化不足，稀疏性优势未充分体现[55]

性能基本原理 - 阿姆达尔定律指出，对于固定问题规模，通过增加计算资源实现的最大加速比受限于串行部分，即使并行资源无限增加，加速比也只能趋近于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]