1. 张量加速核概述：什么是张量加速核、为什么需要它、发展历史与现状

1.1 什么是张量加速核？

张量加速核，英文叫 Tensor Core。说白了，它是一种专门为矩阵乘法设计的硬件单元。

你可能会问：CPU 和 GPU 不是也能做矩阵乘法吗？没错，但它们做得不够快。尤其是当你面对深度学习模型里那些动辄几百乘几百、甚至上千乘上千的矩阵时，传统计算单元就显得力不从心了。

我个人的理解是：张量加速核就像一条高速公路，而普通计算单元就像乡间小路。同样是从 A 点到 B 点，高速公路能让你在相同时间内运送更多货物。

具体来说，张量加速核能在一个时钟周期内完成一个 4×4 矩阵的乘加运算。这在传统 CUDA 核心上需要多个周期才能完成。

核心定义：张量加速核是一种专为矩阵乘法（GEMM）优化的硬件执行单元，支持混合精度计算（如 FP16 输入、FP32 累加），能在单周期内完成小规模矩阵的乘加操作。

1.2 为什么需要张量加速核？

这个问题其实很直接。我们来看看深度学习计算的特点。

深度学习模型，尤其是 Transformer 架构，计算量主要集中在矩阵乘法上。以 BERT-Large 为例，一次前向传播需要执行大约 10^12 次浮点运算。如果用传统 GPU 核心跑，速度会非常慢。

我在项目中遇到过这样的情况：用 V100 跑一个 12 层的 Transformer 模型，训练一个 epoch 要 3 天。后来切换到支持 Tensor Core 的 A100，同样的模型，同样的精度要求，时间缩短到了 1 天以内。差距就是这么明显。

为什么会这样？因为传统计算单元在处理矩阵乘法时，存在几个瓶颈：

指令开销大：每个乘加操作都需要取指令、译码、执行，效率低
数据搬运频繁：矩阵元素需要反复从内存搬运到寄存器
并行度受限：单个核心能同时处理的操作数有限

张量加速核通过硬件直接实现矩阵乘法的数据流，把这些开销降到了最低。

我的建议：如果你在训练大模型，一定要检查你的框架是否启用了 Tensor Core。很多情况下，默认配置并没有开启，性能会差很多。

1.3 张量加速核的发展历史

嗯，这里我们简单梳理一下时间线。

年份	产品	关键特性
2017	NVIDIA Volta V100	第一代 Tensor Core，支持 FP16 输入、FP32 累加
2020	NVIDIA Ampere A100	第三代 Tensor Core，支持 TF32、BF16、INT8 等多种精度
2022	NVIDIA Hopper H100	第四代 Tensor Core，支持 FP8、Transformer Engine
2023	AMD MI300X	Matrix Core，对标 Tensor Core
2024	Intel Gaudi 3	集成矩阵加速单元

从这张表你能看到，张量加速核的发展速度非常快。几乎每两年就有一次重大升级。

我记得 2017 年第一次看到 V100 的 Tensor Core 时，心里想的是：这东西真的能用吗？结果一测试，性能提升确实惊人。从那以后，我就开始深入研究它的原理。

1.4 张量加速核的现状

现在，张量加速核已经成为 AI 芯片的标配。不管是 NVIDIA、AMD 还是 Intel，都在自己的产品中集成了类似的硬件单元。

目前的主流趋势有几点：

精度越来越低：从 FP16 到 BF16，再到 FP8，甚至 INT4。精度越低，计算越快，但需要配合量化技术
规模越来越大：单个芯片上的 Tensor Core 数量从几十个增加到几百个
编程越来越简单：早期需要手动调用 wmma API，现在编译器能自动识别矩阵乘法并映射到 Tensor Core

我曾经踩过一个坑：在 A100 上用 FP16 跑一个模型，结果精度损失严重。后来发现是累加精度不够，改成 FP16 输入、FP32 累加就好了。所以，精度选择一定要根据模型需求来，不是越低越好。

注意：张量加速核虽然快，但并不是所有矩阵乘法都能加速。只有满足特定尺寸（如 4 的倍数）和特定数据布局的矩阵才能被高效处理。如果你的矩阵尺寸不对，性能反而可能下降。

1.5 核心知识体系

下面这张图展示了张量加速核的核心知识体系，我把它画成了 SVG 结构图，方便你理解。

这张图把张量加速核的知识分成了三大块：硬件架构、计算原理和编程接口。后面我们会逐一深入讲解。

1.6 一个简单的代码示例

为了让你有个直观感受，我写一段 CUDA 代码，展示如何用 Tensor Core 做矩阵乘法。这段代码用的是 CUDA 的 wmma API。

#include <cuda_fp16.h>
#include <mma.h>

using namespace nvcuda;

// 使用 Tensor Core 进行 16x16 矩阵乘法
__global__ void tensorCoreMatMul(half *a, half *b, float *c) {
    // 声明 warp 级矩阵分片
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

    // 初始化累加器
    wmma::fill_fragment(c_frag, 0.0f);

    // 加载矩阵数据
    wmma::load_matrix_sync(a_frag, a, 16);
    wmma::load_matrix_sync(b_frag, b, 16);

    // 执行矩阵乘加：C += A * B
    wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

    // 存储结果
    wmma::store_matrix_sync(c, c_frag, 16, wmma::mem_row_major);
}

这段代码看起来简单，但背后涉及了很多硬件细节。比如 wmma::mma_sync 这个函数，它会被编译器映射到一条 PTX 指令，然后由 Tensor Core 硬件直接执行。

我在刚开始接触这个 API 时，总觉得它像个黑盒子。后来深入研究了 PTX 指令集，才真正理解了它的工作原理。这部分内容我们会在后面的章节详细展开。

小提示：如果你想在代码里用 Tensor Core，建议先检查你的 GPU 是否支持。用 cudaDeviceGetAttribute 查询 cudaDevAttrTensorCore 属性即可。

1.7 小结

这一章我们讲了张量加速核的基本概念、为什么需要它、以及它的发展历史和现状。你可能会觉得内容有点多，但这些都是后续章节的基础。

记住几个关键点：

张量加速核是专门做矩阵乘法的硬件单元
它能大幅提升深度学习训练和推理的性能
目前主流 AI 芯片都集成了类似单元
编程接口越来越友好，但底层原理仍然重要

下一章，我们会深入硬件架构，看看 Tensor Core 在芯片内部到底长什么样。

📚 课程目录

1. 张量加速核概述：什么是张量加速核、为什么需要它、发展历史与现状

1.1 什么是张量加速核？

1.2 为什么需要张量加速核？

1.3 张量加速核的发展历史

1.4 张量加速核的现状

1.5 核心知识体系

1.6 一个简单的代码示例

1.7 小结