Nvidia TensorCore
输入“/”快速插入内容
Nvidia TensorCore
Nvidia TensorCore
用户542025年11月3日修改
2024年7月6日创建
23618
25953
Introduction
TensorCore 历史概述
序号 | 时间 | 架构 | 异步数据搬运机制 | 编程模型 & 计算指令 | 数据类型 | 里程碑意义 |
1 | 2017 | Volta (V100) | 无异步机制 ld.*/st.* | Warp 级 wmma 指令 | FP16 | 开启混合精度计算，第 1 代 TensorCore |
2 | 2018 | Turing (T4) | 同 Volta 架构 | 同 Volta 架构 | + int8, int4, int1 | 低精度推理，推动 AI 推理普及 |
3 | 2020.05 | Ampere (A100) | 软件管理的异步拷贝 cp.async.* | Warp 级 mma.sync，可编程 Thread 级数据排布 | + fp32, bf16 | 大幅提升 FP32 训练性能，结构化稀疏 |
4 | 2022.03 | Hopper (H100) | 硬件 TMA cp.async.bulk.tensor.* | thread block 级 wgmma & 硬件 TMA 引入 Transformer 引擎 | + fp8 | 引入FP8和动态精度切换，为 LLM 训练优化 |
5 | 2024.03 | Blackwell (B200) | 专用的 Tensor Memory | 跨SM协作的 UTCMMA.2SM & Tensor Memory 第 2 代 Transformer 引擎 | + fp4, fp6 | 专为万亿参数LLM推理和训练设计 |
演进的核心脉络
3 个关键优化点：
1.
内存能力。核心是异步数据搬运机制 & 内存层级设计。AsyncCopy -> 硬件 TMA -> Tensor Memory。
2.
数据格式。越来越丰富，特别是 FP8 / FP4 的引入。
3.
编程模型。提升并行层级：Warp (mma) => Thread Block (wgmma) => 跨 SM (UTCMMA.2SM)。
计算指令的演进
待补充：指令 shape & 单指令算力 flops。
AI Compiler，比如 Triton。不需要支持 wmma 指令。用 mma 足以取代 wmma 的功能。
MMA & WMMA - A100
都是 warp level 的执行模型。
H100 之前，并存的 2 个主要指令。
MMA 指令特性:
•
线程级控制：支持手动控制线程级别的数据排布和寄存器分配。
•
灵活性高，性能好。支持自定义数据布局优化。
•
编程复杂度：高，需要显式管理寄存器分配和数据流水。
WMMA 指令特性
•
封装抽象：封装了完整的 Shared Memory → MMA → Shared Memory 流程。
•
易用性优先：降低编程门槛，适合快速原型开发。
•
灵活性限制：难以规避 SMEM bank conflict，性能优化空间有限，无法自定义数据排布。
尽管 mma 是通过一个 warp 去调度实现的，但还是粒度为线程级别时更容易编写。
以下场景 WMMA 无法实现，但 MMA 支持：
1.
数据排布优化：自定义Shared Memory → MMA的swizzle layout
2.
计算融合：将elementwise操作直接融合到MMA累加寄存器
3.
复杂数据重用：实现更高效的数据局部性优化