一、TensorCore前世今生：从Fermi到Hopper架构演进
1.1 为什么要有TensorCore？
说起TensorCore，得先聊聊我入行那会儿的经历。2012年我刚接触CUDA时，用的还是Fermi架构的GTX 580。那时候做深度学习训练，一个简单的卷积层要跑好几分钟。你想想看，现在一个batch几毫秒就完事了，这差距有多大。
其实GPU一开始并不是为深度学习设计的。它最早是为了图形渲染而生。但后来大家发现，GPU做矩阵乘法特别快——这不就是神经网络的核心运算吗？
我简单梳理一下架构演进的几个关键节点：
| 架构 | 发布年份 | 关键特性 | 对深度学习的影响 |
|---|---|---|---|
| Fermi | 2010 | 首个完整CUDA架构，ECC内存 | 开始支持通用计算 |
| Kepler | 2012 | GK104，动态并行 | AlexNet训练成为可能 |
| Maxwell | 2014 | 低功耗，共享内存优化 | 推理场景开始普及 |
| Pascal | 2016 | NVLink，FP16支持 | 混合精度训练萌芽 |
| Volta | 2017 | 第一代TensorCore | 深度学习专用硬件诞生 |
| Turing | 2018 | INT8/INT4 TensorCore | 推理加速新高度 |
| Ampere | 2020 | 第三代TensorCore，TF32/BF16 | 训练推理全面加速 |
| Hopper | 2022 | 第四代TensorCore，DPX指令 | Transformer优化专用 |
嗯，这里要注意。Volta架构的V100是第一块真正为AI设计的GPU。我记得当时拿到V100的测试卡，跑同样的ResNet-50，比P100快了将近4倍。这可不是简单的频率提升，而是硬件层面做了专门的矩阵运算单元——TensorCore。
1.2 TensorCore硬件原理
说白了，TensorCore就是一个专门做矩阵乘加运算的硬件单元。它跟CUDA Core最大的区别是什么？
CUDA Core做的是标量运算：一次算一个数。TensorCore做的是矩阵运算：一次算一个4x4的小矩阵。
我画个图帮你理解：
你看，一个TensorCore在一个时钟周期内，能完成一个4x4矩阵的乘加运算。这相当于64个乘法和48个加法同时完成。而一个CUDA Core一个周期只能做一个乘法。
关键数据：V100上有640个TensorCore，每个SM有8个。理论上，TensorCore的矩阵运算吞吐量是CUDA Core的16倍。我实际测试下来，在混合精度训练场景下，大概能拿到8-12倍的加速比。
1.3 混合精度训练概念
说到混合精度训练，很多人第一反应就是「用FP16代替FP32嘛」。其实没那么简单。
我刚开始做混合精度训练时踩过一个坑。当时直接把所有参数都转成FP16，结果模型死活不收敛。后来才发现，梯度更新那一步必须用FP32，否则精度不够。
混合精度训练的核心思路是：
- 前向传播和反向传播：用FP16计算，速度快，省显存
- 权重更新：用FP32计算，保证精度
- 损失缩放：防止梯度下溢
这里我贴一段实际用的混合精度训练代码：
// 混合精度训练核心逻辑（伪代码）
// 1. 前向传播：FP16
__half* input_fp16 = cast_to_half(input_fp32);
__half* weight_fp16 = cast_to_half(weight_fp32);
__half* output_fp16 = matmul_tensorcore(input_fp16, weight_fp16);
// 2. 损失计算：FP32
float loss = compute_loss(cast_to_float(output_fp16), target);
// 3. 反向传播：FP16
__half* grad_fp16 = cast_to_half(compute_gradient(loss));
// 4. 权重更新：FP32（关键！）
float* grad_fp32 = cast_to_float(grad_fp16);
float* weight_fp32 = get_master_weights(); // 保留FP32副本
weight_fp32 -= learning_rate * grad_fp32;
// 5. 损失缩放
float scaled_loss = loss * loss_scale;
// 反向传播时使用 scaled_loss
// 更新权重前除以 loss_scale我的经验：刚开始做混合精度训练时，建议先用PyTorch的AMP（Automatic Mixed Precision）库。它自动帮你处理了损失缩放和精度转换。等熟悉了再手写CUDA版本。
1.4 TensorCore与CUDA Core的区别
这个问题我面试时经常问候选人。很多人答不上来，或者答得很模糊。
其实区别就三点：
| 维度 | CUDA Core | TensorCore |
|---|---|---|
| 运算粒度 | 标量（一个数） | 矩阵（4×4） |
| 适用场景 | 通用计算、控制流 | 矩阵乘法、卷积 |
| 精度支持 | FP32/FP64/INT32 | FP16/BF16/TF32/INT8 |
| 编程方式 | 标准CUDA C++ | 需调用wmma API或cuBLAS |
| 吞吐量 | 基准 | 8-16倍于CUDA Core |
说白了，CUDA Core是通用CPU的替代品，什么都能算。TensorCore是专用加速器，只做一件事——矩阵乘法，但做得极快。
我打个比方：CUDA Core就像一把瑞士军刀，什么都能干。TensorCore就像一把专业的厨师刀，只切菜，但切得飞快。你让厨师刀去拧螺丝，肯定不行。但切菜时，瑞士军刀完全比不了。
注意：不是所有运算都能用TensorCore加速。只有满足以下条件的运算才能利用TensorCore：
- 矩阵维度是8的倍数（Ampere架构放宽到16的倍数）
- 数据类型是FP16、BF16、TF32或INT8
- 运算类型是矩阵乘法或卷积
我曾经见过有人试图用TensorCore做向量加法，结果性能反而更差。因为TensorCore启动开销大，小矩阵用反而不划算。
1.5 从Volta到Hopper：TensorCore的进化
每一代TensorCore都在进化。我简单说说关键变化：
- Volta（第一代）：只支持FP16输入，FP32累加。矩阵大小固定为4x4。
- Turing（第二代）：加入INT8和INT4支持，开始支持推理场景。
- Ampere（第三代）：引入TF32格式，无需改代码就能用。支持BF16。矩阵大小扩展到16x16。
- Hopper（第四代）：加入DPX指令，专门优化Transformer。支持FP8。
我个人觉得，Ampere的TF32是个里程碑。你想想看，以前用FP32训练，换成TF32后，代码一行都不用改，直接获得2-3倍加速。这种「免费午餐」在工程领域可不多见。
Hopper的DPX指令更是为Transformer量身定制。现在大模型这么火，这个优化来得正是时候。
一句话总结：TensorCore就是NVIDIA为深度学习打造的「硬件加速器」。它牺牲了通用性，换来了极致的矩阵运算性能。用好它，你的训练速度能提升一个数量级。
好了，这一章就讲到这里。记住我说的：TensorCore不是万能药，但在矩阵乘法这个场景下，它就是最快的选择。下一章我们聊聊如何用CUDA编程调用TensorCore。