这次不看具体模型权重也不讲部署脚本单独把 NVIDIA GPU 里那个叫 Tensor Core 的部件拆开聊清楚。大模型训练和推理的大部分耗时都集中在矩阵乘法上而 Tensor Core 就是专门为矩阵乘加设计的电路。搞懂它为什么快、怎么用、什么时候不生效比记住一堆算力争辩更有用。很多人在本地部署大模型时会看到 FP16、BF16、INT8、FP8 这些概念也会看到 4090、A100、H100、RTX 50 系这样的型号差异。它们之间的真正分水岭往往不只是“显存大”或“带宽高”而是 Tensor Core 的代数、精度支持以及软件库有没有把数据切成适合 Tensor Core 的形状。这篇文章就从“矩阵分块”这个角度把加速原理讲清楚。本文会覆盖Tensor Core 与 CUDA Core 的定位差异、矩阵分块拆解 GEMM 的数学逻辑、为什么 Transformer 基本是一连串 GEMM、用 Python 做一次分块矩阵乘原型、cuBLAS / PyTorch 里的调用姿势最后给出常见误区和实际排查思路。适合对 GPU 原理感兴趣、想理解大模型训练推理差异的后端开发者、算法工程师和本地部署玩家。1. Tensor Core 核心能力速览项目说明硬件定位NVIDIA GPU 中专门的矩阵乘加执行单元从 Volta 世代开始引入核心操作D A × B C以 tile 为单位完成矩阵乘累加与 CUDA Core 的区别CUDA Core 做标量乘加Tensor Core 一次处理一个小分块典型精度FP16、BF16、TF32、INT8、FP8、FP4不同代际支持不同累加器通常使用 FP32 累加降低半精度乘法带来的误差依赖软件cuBLAS、cuDNN、CUTLASS、PyTorch、TensorRT 等算子库或框架是否独立占显存不独立占用显存作为 SM 内的计算单元存在数据来自共享内存和寄存器与大模型的关系线性层、Attention 的 QKV 投影、FFN 计算几乎都由 GEMM 承载加速前提数据分块、精度对齐、库版本匹配缺一点都可能回退到低效路径这张表先给结论。下面逐个拆开。2. 为什么大模型计算几乎等于矩阵乘法大模型的基础模块是 Transformer。不管输入是文本、图片还是语音核心计算路径里一定包含两类矩阵乘第一类是线性投影。每个 token 的向量会乘上权重矩阵得到 Query、Key、Value 和 FFN 中间状态。公式非常简单Q X × W_QK X × W_KV X × W_VFFN 中间结果 X × W_1FFN 输出 中间结果 × W_2第二类是 Attention 的相似度计算。Q 和 K 的转置相乘得到注意力分数再和 V 做矩阵乘法。这部分在长文本场景中尤其消耗算力。把输入批量拼在一起后这些计算就变成了巨型 GEMMX 的形状是 [batch × seq_len, hidden_size]W 的形状是 [hidden_size, output_size]计算结果就是一路 tile 一路 tile 滚出来的大矩阵所以大模型训练和推理的绝大多数浮点操作本质是在不断做 C A × B 或 D A × B C。只要把矩阵乘积算得足够快整个大模型就跑得更快。3. 矩阵分块到底在分什么先看矩阵乘法的原始定义。假设 A 是 M×K 矩阵B 是 K×N 矩阵C 是 M×N 矩阵那么 C[i][j] 需要对 k 求和其中每个输出元素都要做 K 次乘加。如果直接按定义实现会有两个问题内存访问不友好A 的一行和 B 的一列反复从全局内存读取。并行粒度不够没有充分复用已经加载进 GPU 的数据。分块的核心思路是不一次算完整个 C而是把 A、B、C 都按照固定大小切成若干个小块比如以 32×32、16×16、8×8 为粒度。分完以后每个 C 小块只依赖 A 的某几块和 B 的某几块。用 2×2 分块示意A 被切成 A00、A01、A10、A11B 被切成 B00、B01、B10、B11C 被切成 C00、C01、C10、C11那么C00 A00 × B00 A01 × B10C01 A00 × B01 A01 × B11C10 A10 × B00 A11 × B10C11 A10 × B01 A11 × B11这个拆分是精确的不丢失任何信息。只是把原本一个大的计算循环改成了多个小矩阵相乘后相加。为什么要这样做因为 GPU 的存储层级和 CPU 类似。寄存器最快共享内存次之全局显存最慢。如果每个线程都从全局显存里随机取数带宽会被快速打满。分块后每个计算块可以把 A 的子块和 B 的子块先加载到共享内存或寄存器里反复复用后再算下一块访存效率大幅提高。Tensor Core 的工作方式其实就是这种分块思路的硬件化。GPU 不是等软件去把矩阵拆成单个元素循环而是在硬件层面直接接收一个小分块做一次“矩阵块乘加”。4. 用 Python 写一个分块矩阵乘原型为了直观理解矩阵分块这里先用 Python 写一个最简单的分块 GEMM。它不会调用 Tensor Core纯粹用于理解计算顺序和分块边界。import numpy as np def matmul_blockwise(A, B, block_size8): M, K A.shape K2, N B.shape assert K K2, A 的列数必须等于 B 的行数 C np.zeros((M, N), dtypenp.float32) for i in range(0, M, block_size): i_end min(i block_size, M) for j in range(0, N, block_size): j_end min(j block_size, N) for k in range(0, K, block_size): k_end min(k block_size, K) # 取出当前小块 a_tile A[i:i_end, k:k_end] b_tile B[k:k_end, j:j_end] # 累加小块乘积 C[i:i_end, j:j_end] np.dot(a_tile, b_tile) return C # 随机生成一个形状验证分块结果 M, K, N 64, 32, 48 A np.random.randn(M, K).astype(np.float32) B np.random.randn(K, N).astype(np.float32) C_block matmul_blockwise(A, B, block_size8) C_ref np.dot(A, B) print(分块矩阵乘结果最大误差, np.abs(C_block - C_ref).max())这段代码的输出应当是接近 0 的误差。因为采用 float32 时普通矩阵乘内部也可能有不同累加顺序所以只要误差在 1e-4 以内就说明分块逻辑本身正确。真正在 GPU 上优化时block_size 并不是越大越好。分块太大共享内存放不下分块太小数据复用率不够。常见的高性能实现会同时考虑波形、缓存占用、寄存器数量等多个因素。5. CPU、GPU 与 Tensor Core 的职责边界很多人会把“GPU 算得快”简单理解成“GPU 有很多核”。显卡确实有很多 CUDA Core但 CUDA Core 做的是标量乘加。所谓标量就是一次只算单个元素的乘积1 个线程计算 c a * b c并行靠大量线程堆出来Tensor Core 则不同它在硬件层面支持“矩阵块乘加”。线程可以以 warp 为单位一次把一个小矩阵块交给 Tensor Core 处理。常见的 WMMA 系列接口以 16×16×16、16×8×8 之类的小块为粒度。不同架构的 MMA 指令不完全一样但思想一致用专用电路替代多次循环。CUDA Core 做一次乘加需要多个指令周期Tensor Core 用同样时间可以完成整个分块矩阵的乘加。后者不是靠频率取胜而是靠增加单次操作的数据宽度和计算并行度。但 Tensor Core 并不是万能加速器如果计算不是矩阵乘形态Tensor Core 发挥不了优势。如果数据精度不是 Tensor Core 支持的精度可能会回退到普通 CUDA Core 路径。如果矩阵太小调用 Tensor Core 的准备开销可能比实际计算还大最终不一定更快。所以大模型推理和训练的性能优化本质上是在回答一个问题如何把计算最大程度地组织成可交给 Tensor Core 的 GEMM并保证访存不拖后腿。6. 存储层级与数据搬运不搬数据就没有加速Tensor Core 的计算单元在 SM 内部。真正要发挥性能需要把 A 的 tile 和 B 的 tile 搬到 SM 附近的存储位置。以一张常见的 NVIDIA GPU 为例计算过程大致如下全局显存保存完整矩阵 A 和 B。GPU 把 A 的一小块和 B 的一小块读取到共享内存。数据从共享内存进入寄存器交给 Tensor Core 执行矩阵乘加。结果 C 的 tile 暂存到寄存器再写回共享内存最后写回全局显存。这就是分块矩阵乘的硬件落地版。如果没有分块而是让每个线程独立读取大量数据全局显存带宽会成为绝对瓶颈。Tensor Core 算得再快也只能等数据。值得注意的另一个问题是精度和带宽的关系。矩阵乘法的计算量是 M×N×K×2但数据量只有 M×K K×N M×N。如果把权重从 FP32 改成 FP16权重矩阵的数据体积直接减半带宽压力随之下降。这也是为什么大模型训练和推理普遍强调混合精度而不是所有情况都强制使用更高精度。Tensor Core 在硬件上支持 FP16 输入、FP32 累加的设计正好解决了两个问题FP16 权重体积更小访存更划算。半精度乘法带来的累积误差先保存到 FP32 累加器中避免结果漂移过大。这也是为什么 Traning 里常用 BF16 / FP16而不是用 FP32 直接硬算的原因。7. 实际软件路径cuBLAS、PyTorch 怎样调用 Tensor Core普通开发者不会直接写 PTX 指令而是通过高层库触发 Tensor Core。这里最常见的路径有cuBLAS / cuBLASLt提供 GEMM 接口底层根据精度和 GPU 架构选择 Tensor Core 内核。cuDNN卷积和 Transformer 相关原语会调用 Tensor Core 优化后的内核。PyTorchtorch.matmul、nn.Linear、torch.bmm等会链接到不同后端的 cuBLAS 或自定义内核进而使用 Tensor Core。CUTLASSNVIDIA 开源的高性能 GEMM 模板库适合研究人员和框架开发人员二次定制。TensorRT面向推理的优化引擎自动做层融合和精度转换然后生成 Tensor Core 内核。如果在 C 里直接调 cuBLAS比较接近硬件的接口调用方式可能是cublasGemmEx。下面给出一段示意代码具体参数需要按项目中的实际矩阵排布调整#include cublas_v2.h // 这段代码只是为了展示 cuBLAS GEMM 接口的常见形态 // 实际工程中需要根据矩阵维度、存储顺序和指针类型做完整配置。 cublasHandle_t handle; cublasCreate(handle); float alpha 1.0f; float beta 0.0f; // 这里假设 A、B、C 都是已经准备好的 half 类型显存指针 // 形状参数 M、N、K 需要从当前任务中换算 cublasStatus_t status cublasGemmEx( handle, CUBLAS_OP_N, CUBLAS_OP_N, M, N, K, alpha, A, CUDA_R_16F, lda, B, CUDA_R_16F, ldb, beta, C, CUDA_R_16F, ldc, CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP ); if (status ! CUBLAS_STATUS_SUCCESS) { // 处理错误 } cublasDestroy(handle);这段代码展示的是 FP16 输入、FP32 累加、Tensor Core 优先的典型配置。如果换成 INT8还需要准备量化参数和缩放因子如果是 Ampere 系显卡想用 TF32也要在 cuBLAS 里设置对应的 math mode。不同 CUDA 版本对cublasGemmEx的支持细节并不完全一致以本机头文件和官方文档为准。PyTorch 用户通常不需要直接写这个但可以通过开关影响 Tensor Core 行为。一个常见的例子是 TF32 在某些矩阵乘里默认可能开启也可能没有开启取决于框架版本和全局设置。老版本中常见写法如下import torch # 让 PyTorch 允许在 Ampere 及更新架构上使用 TF32 做 FP32 输入矩阵乘 torch.backends.cuda.matmul.allow_tf32 True如果跑矩阵乘时发现 FP32 输入并没有明显变慢大概率就是走了 TF32 Tensor Core 路径。如果 TF32 关闭则可能回退到普通 FP32 计算速度会下降但某些场景精度表现更接近纯 FP32。8. 为什么“显存大”不等于“算力强”本地部署大模型时大家通常先看显存因为显存决定模型能不能装下。但真正生成 token 时显存只是必要条件量产速度取决于另一组指标。矩阵乘法的瓶颈有两种典型情况计算瓶颈每单位数据上的计算量非常大GPU 的 Tensor Core 一直处于高占用状态。访存瓶颈每算一个结果只需要很少的浮点计算数据搬来搬去的时间远大于计算时间。大模型推理的某些阶段很特殊。例如单用户、单 token 流式生成时batch size 可能很小矩阵乘法的一维尺寸很小计算密度不高这时 Token 生成速度往往受带宽和显存读取限制Tensor Core 未必能跑满。而批量推理时多个请求拼在一起矩阵乘法的形状变大Tensor Core 的利用率明显提升单卡吞吐会有很大改善。这就是为什么很多推理优化工具强调 continuous batching 和 paged attention。它们不仅是为了省显存也是在尽可能把零散请求聚合成更大的 GEMM让 Tensor Core 有足够的 tile 可以连续处理。9. 从 Volta 到 BlackwellTensor Core 的代际变化Tensor Core 第一次出现在 NVIDIA Volta 架构的 V100 上当时主要面向 FP16 深度学习训练。到 Turing 架构Tensor Core 增加了 INT8 和 INT4 支持更偏推理。到 Ampere 架构加入 TF32 和 BF16 支持也支持 2:4 结构化稀疏加速。这个阶段 A100 的大模型训练地位很突出。Hopper 架构进一步强化了对 Transformer 的加速H100 加入了 Transformer Engine 等面向 FP8 的设计。到 Blackwell 架构第五代 Tensor Core 对 FP4、FP6、FP8 等低精度推理能力做了进一步扩展。每一代 Tensor Core 的变化并不只是把乘法变快还往往在调整精度支持、累加器设计、共享内存容量和指令形状。普通用户记忆这些代际差异没有太大必要只需要理解Tensor Core 支持的精度和软件库是否适配直接决定了模型能否以低显存、高吞吐的方式运行。这里不建议直接搬规格书上的 TOPS 数字当作实际收益。规格书中的算力往往是在最优数据布局、特定小分块、理想指令流下测出来的。真实的大模型计算里还要考虑内存带宽、中间激活、布局转换和算子切换实际利用率很难达到理论峰值。10. 如何判断 Tensor Core 是否在工作这是很多人的困惑我在代码里用了半精度Tensor Core 就一定会生效吗不一定。可以按下面顺序检查查看 GPU 架构。Tensor Core 需要 Volta 及以后的架构太老的卡不支持。查看输入精度。用 FP32 调用普通cublasGemm在 Ampere 上可能走 TF32 Tensor Core也可能走普通 FP32。用 FP16/BF16 输入更容易触发 Tensor Core。查看算子库日志。PyTorch、TensorRT 有自己的 kernel 选择逻辑可以通过 profiler 确认实际跑的内核名称。查看矩阵维度。Tensor Core 的 tile 内部通常需要对齐例如某一些内核要求 M、N、K 是 8 或 16 的倍数。cuBLAS 会在边界做处理但极端小矩阵可能不会选择 Tensor Core 内核。查看显存和 SM 利用率。如果 SM 利用率很高但 Tensor Core 相关流水线没有动作可能是算子本身不是 GEMM 类型或者走了别的实现路径。NVIDIA Nsight Compute 可以查看实际执行的指令和硬件单元利用率。命令行里可以像下面这样运行一个 Python 脚本进行 profile但具体指标名称和参数需要根据 Nsight 版本调整# 示例用 Nsight Compute 分析一个 Python 训练/推理脚本 # 如果是 PyTorch建议先缩小数据量否则 profile 会非常慢 ncu --set full python train.py运行后重点看矩阵乘相关内核而不是只看 PyTorch 的 Python 调用栈。NDRange 形状、shared memory 使用量、寄存器占用、是否执行了 Tensor Core 指令这些信息能帮助判断热点在哪里。如果不想用这么重的工具也可以用更简单的对照法把同一个矩阵乘分别用 FP32 和 FP16 跑一遍观察耗时和显存带宽差异。如果 FP16 明显更快说明半精度路径大概率生效如果两边差距不大可能模型太小、数据搬运占主导或者 Tensor Core 没有进入预期内核路径。11. 常见问题与排查表问题现象可能原因排查方式解决方案Python 里开了 FP16速度却没有明显提升矩阵太小访存瓶颈算子库没有选择 Tensor Core 内核用 profiler 查看实际内核增大 batch或检查算子是否支持半精度FP32 输入时速度不稳定TF32 可能开启也可能关闭不同框架默认值不同查看 torch.backends.cuda.matmul 配置按模型精度需求显式开关 TF32显存够大但生成速度很慢推理是访存瓶颈Tensor Core 利用率不高观察生成时 GPU 利用率和显存带宽开启批量推理或连续批处理相同卡上训练速度比预期低数据布局、精度、算子实现未匹配 Tensor Core用 Nsight Compute 看 kernel改用混合精度训练低精度量化后效果变差FP8 / INT8 量化范围和缩放参数设置不当比较量化前后输出分布检查量化方法和校准集导入大模型时报算子不支持PyTorch / TensorRT 版本过旧查日志和内核选择结果升级 CUDA、cuDNN、PyTorch跑很多小矩阵乘比预期慢每次 GEMM 启动和准备开销较大观察 GPU 是否持续高占用合并矩阵或使用批处理 GEMMTensor Core 占用看起来为零没有跑 GEMM 类内核或者跑的是不支持 Tensor Core 的实现检查算子类型和库版本调整算子实现或调用更高层算子12. 日常工程中的最佳实践理解 Tensor Core 之后再看大模型训练和部署会更有方向感。实际工程中有几个经验可以复用。第一个经验不要一开始就追求手动控制 Tensor Core。先用框架默认路径配合半精度或适合场景的量化、开启批量推理已经能拿到大部分收益。自己写 GEMM 内核是非常深的坑普通业务场景不值得直接碰。第二个经验区分计算瓶颈和访存瓶颈。同一个 GPU 上跑大 batch 训练和跑小 batch 在线推理瓶颈位置完全不同。前者要重点看算力是否吃满、Tensor Core 是否在忙后者要重点看显存带宽、模型权重加载和 KV cache 访问。第三个经验精度选择要服务于最终效果。Tensor Core 支持半精度不等于任何环节都该用半精度。Embedding、LayerNorm、Softmax 等对精度敏感或计算量占比不高的部分保持 FP32 可能是合理选择。混合精度训练就是让 GEMM 部分跑在半精度或更低精度同时让关键累积步骤保持较高精度。第四个经验数据分块要结合缓存。大模型推理时把不必要的中间激活尽量回收让共享内存能装下更大 tile比单纯堆算力更重要。很多优化策略表面上是省显存实际也改善了分块效率。第五个经验在做模型效果评估或性能优化时要记录精度、batch size、输入长度、GPU 利用率、显存带宽、耗时和内核选择。没有这些记录很难定位问题是出在算子库没选对还是数据分块不合理。面对人脸、声音、版权内容的模型数据也要注意合规问题。Tensor Core 本身只是计算单元不涉及具体数据语义。但如果你用它做模型微调、推理服务或内容生成需要确保数据来源合法、素材使用已获授权、输出内容不侵犯他人权益尤其涉及人物肖像和受版权保护的音频、图像、文本时要有明确授权记录。13. 总结Tensor Core 加速大模型的核心不是“更快的 CPU 核”而是用专用电路去承接完整的矩阵分块乘加。大模型的计算主体是 GEMMGEMM 可以拆成很多小矩阵块Tensor Core 正好就是为这些小矩阵块设计的。理解这一点后再去看大模型架构和推理框架会更容易抓住重点。量化是为了减小数据体积、降低访存压力混合精度是为了让 Tensor Core 用更小数据宽度做矩阵乘批量推理是为了让矩阵乘维度更大、Tensor Core 更容易吃满FlashAttention 等优化则是在管理 tile 的加载和复用让 GEMM 周围的其他开销不拖后腿。第一次动手验证时可以先用小矩阵跑一遍分块矩阵乘原型再跑一次 PyTorch 的 FP32 与 FP16 对照顺便用 Nsight Compute 看一眼实际内核就能有明显感知。对绝大多数使用者来说不需要掌握 PTX 指令重点是能判断自己的代码是否走到了 Tensor Core 路径以及在什么场景下应该调整 batch、精度和算子库版本。接下来可以继续沿着 CUTLASS 源码、FlashAttention 论文或 NVIDIA 混合精度文档深入这些例子都会反复出现同一个关键词矩阵分块。