这次我们不聊具体某一个大模型直接把镜头拉到 GPU 内部搞清楚一个问题Tensor Core 凭什么能让大模型跑得更快。只看显存、只看参数量、只看量化位数很多人在本地部署大模型时容易忽略一个底层事实——每一次推理本质上都是海量的矩阵乘法。Transformer 里的 QKV 投影、注意力分数、FFN前馈网络拆到最底层全是 GEMM通用矩阵乘法。而 Tensor Core 就是 NVIDIA 专门为这类矩阵运算设计的硬件加速单元。它和普通 CUDA 核心走的是完全不同的执行路径。这篇文章用“矩阵分块”这条主线把 Tensor Core 的原理、大模型推理场景下的实际意义、以及本地部署时如何验证 Tensor Core 是否真正生效一次讲清楚。如果你是做模型训练、微调、推理部署或者手上有支持 Tensor Core 的 NVIDIA 显卡想看明白“为什么大模型相关工具都强调 FP16 / BF16”这篇值得读完。1. 核心概念速览能力项说明项目主题Tensor Core 硬件原理与矩阵分块加速核心概念Tensor Core、GEMM、矩阵分块、WMMA、FP16/BF16/FP8解决的问题解释大模型推理和训练中矩阵乘法如何被硬件加速硬件门槛NVIDIA 从 Volta 架构开始引入 Tensor Core消费级 GeForce 20 系以后基本覆盖是否需要 GPU理解原理不需要 GPU验证 Tensor Core 生效需要 NVIDIA GPU软件栈CUDA、cuBLAS、CUTLASS、PyTorch、TensorRT、llama.cpp 等典型收益同样跑 FP16/BF16 的矩阵乘法Tensor Core 性能远高于普通 CUDA 核心局限中小矩阵、低批次、无量化场景收益有限需要库和框架正确调用先说结论Tensor Core 不是靠单核频率取胜而是靠“硬件直接执行分块矩阵乘加”这种数据并行方式。绝大多数时候你不需要直接写 Tensor Core 指令但你要知道什么情况下框架会调用它什么情况下不会。2. 为什么大模型离不开矩阵乘法Transformer 是当前大模型的主流结构。无论 7B、13B、70B 模型还是多模态模型推理和训练时的计算密度集中在几类矩阵运算上。以一次标准 Transformer 前向传播为例典型矩阵运算包括输入嵌入与位置信息组合后乘权重矩阵生成 Query、Key、Value。Q 与 K 做点积得到注意力分数本质是一次批量矩阵乘。注意力分数与 V 再乘一次得到注意力输出。注意力输出经过输出投影矩阵。FFN 全连接层中隐藏状态先后乘两个大权重矩阵中间穿插激活函数。每一层都包含多次 GEMM。模型的层数越深、隐藏层维度越大、上下文越长矩阵乘法的规模就越夸张。在部署场景里KV Cache 和批量推理还会把矩阵规模进一步放大。比如连续处理多个请求时输入从单条变多条矩阵的批量维度 M 增大上下文很长时注意力矩阵的序列维度增大。这些维度一旦变大普通 CUDA 核心逐元素计算矩阵乘法的代价会迅速膨胀。这就是 Tensor Core 存在的意义把矩阵乘加从“用多个标量单元慢慢算”变成“用专用张量单元整块算”。3. 矩阵分块加速的本质手段先说矩阵乘法的标准定义。设矩阵 A 为 M×K矩阵 B 为 K×N计算结果 C 为 M×NC[i][j] Σ A[i][k] * B[k][j]k 从 0 到 K-1如果直接按这个公式三重循环计算每个 C[i][j] 都需要读取 A 的第 i 行和 B 的第 j 列。问题在于A 和 B 的数据量往往远大于高速缓存容量反复从内存搬运的代价会淹没计算时间。分块的核心思路是不要按单个元素遍历而是把 A、B、C 划分成一小块一小块的子矩阵每次只加载一块到片上高速缓存中完成这块需要的乘加后再加载下一块。用计算 C A × B 举例。假设输出矩阵 C 是 1024×1024我们可以把输出分成 128×128 的块。每个输出小块的 C 块只需要 A 中对应的 128 行、以及 B 中对应的 128 列就可以完成累加。// 图示矩阵分块累加结构 for (int m0 0; m0 M; m0 TM) { for (int n0 0; n0 N; n0 TN) { // 初始化输出块大小为 TM x TN float C_block[TM][TN] {0}; for (int k0 0; k0 K; k0 TK) { // 把 A 的一块 TM x TK 载入共享内存 // 把 B 的一块 TK x TN 载入共享内存 load_shared(A_block, m0, k0); load_shared(B_block, k0, n0); __syncthreads(); // 在共享内存中做子矩阵乘加 for (int i 0; i TM; i) { for (int j 0; j TN; j) { for (int kk 0; kk TK; kk) { C_block[i][j] A_block[i][kk] * B_block[kk][j]; } } } __syncthreads(); } } }这段代码是一种简化示意真实的高性能实现还要处理边界、Bank Conflict、Double Buffering、向量化等细节。但核心思想已经体现出来了通过分块让数据在缓存中重复使用减少对全局内存的访问。CPU 上用分块是为了提高 L1/L2 缓存命中率GPU 上用分块是为了把数据搬进 SM流式多处理器内部的共享内存和寄存器供计算单元快速读取。Tensor Core 的硬件设计正好是围绕这一目标展开的。4. Tensor Core 如何执行矩阵分块4.1 一条指令完成一个子矩阵乘加普通 CUDA 核心执行的是一条 FMA 指令完成一次标量乘加d a * b c。Tensor Core 不一样。它把“乘加”提升到矩阵级别。执行 Tensor Core 指令时输入是 A 矩阵的子块、B 矩阵的子块输出更新到累加矩阵 C 和 D。例如 NVIDIA 的 WMMAWarp Matrix Multiply-Accumulate编程模型可以让一个线程束Warp32 个线程共同完成一个较大的矩阵乘加操作而不是每个线程独立做一个元素的乘加。这种模式的好处是一个线程束内的线程通过寄存器协作共同维护一个大的矩阵分块。单个时钟周期内硬件完成大量乘加运算而非只是单个标量运算。数据不需要反复经过通用计算单元由专用矩阵计算单元直接计算。4.2 分块要贴合硬件尺寸软件层面的分块可以任意设定但 Tensor Core 真正能发挥性能时分块尺寸要匹配硬件 WMMA 操作数尺寸。NVIDIA 不同架构的 Tensor Core 指令和推荐矩阵形状有差异比如早期 WMMA 支持 16×16×16、16×8×8 等形状。这意味着如果矩阵很小、不满足张量指令要求的最小分块Tensor Core 也未必能发挥最大效率。所以在大模型部署时框架通常会把连续多个 token、多个 batch 拼在一起凑出足够大的矩阵维度再交给 Tensor Core 执行。4.3 数据流全局内存到共享内存再到寄存器Tensor Core 高性能执行的关键不只是计算单元而是数据调动路径。典型路径是把 A 和 B 的分块从全局内存加载到共享内存Shared Memory。从共享内存把数据分发到线程的寄存器中形成 Tensor Core 需要的矩阵片段。调用 Tensor Core 指令完成子矩阵乘加。把累加结果写回共享内存再刷回全局内存。继续遍历后续分块直到完成整个输出矩阵。因此高效启动一个 GEMM 时GPU 内部有大量线程在协同工作一部分负责数据搬运一部分负责计算一部分负责写回结果。所有的工作都围绕分块展开。Tensor Core 把“计算密度高”的那一步加速了如果数据搬运没有跟上整体性能依然会被内存带宽卡住。5. Tensor Core 与显存占用、模型规模的关系5.1 权重精度决定显存下限大模型参数是固定的但存储精度可以选择。模型权重以 FP32 存储一个参数占 4 字节以 FP16/BF16 存储占 2 字节以 INT8 量化存储占 1 字节以 INT4 量化大约占 0.5 字节。按 7B 参数量估算存储精度每参数字节7B 模型权重显存估算FP324约 28GBFP16 / BF162约 14GBINT81约 7GBINT40.5约 3.5GB这只是权重本身。实际运行还需要加载 KV Cache、激活值、临时缓冲区。所以本地跑模型时显存是否够用不能只看权重文件的大小。注意一点这里所说的 FP16/BF16/INT8和 Tensor Core 直接相关。Tensor Core 对半精度和低精度矩阵运算有专门优化。也就是说用 FP16/BF16 跑模型不仅显存比 FP32 省一半计算速度还可能因为 Tensor Core 而大幅提升。这也是为什么主流推理框架默认用 FP16/BF16而不是 FP32。5.2 大模型推理中的典型 GEMM 形状自回归生成模型分两个阶段Prefill预填充一次输入一段 Prompt矩阵的序列长度、batch 都较大。此时 GEMM 的 M、N、K 都很可观Tensor Core 利用率较高。Decode逐 Token 生成每次只生成一个新 Tokenbatch 往往较小矩阵形状较“瘦长”。如果单请求推理M 很小Tensor Core 利用率相对不佳。要提升利用率通常把多个请求合并成一个 batch增大矩阵的 M 维度。所以批量推理不只是提高吞吐量也是为了让矩阵分块更饱满更好喂给 Tensor Core。5.3 显存带宽可能是另一条腿大模型推理还有强内存带宽需求。当模型无法完全放进显存时每算一步都要把权重从系统内存搬运到显存速度会严重变慢。Tensor Core 优化的是计算密集度不能解决权重搬运问题。因此判断大模型部署性能时要同时看计算能力和显存带宽。一个模型在 GPU 上能跑多快等于“计算怎么分块”和“权重怎么搬”共同决定的结果。6. Tensor Core 在不同部署方式中如何被调用6.1 PyTorch 模型推理与训练PyTorch 是常见的大模型训练和推理框架。PyTorch 的矩阵乘法最终会调用 cuBLAS 等底层库而 cuBLAS 会在可用时选择 Tensor Core 的实现路径。对 FP16/BF16 输入这一点基本是自动完成的。对 FP32 输入则可能使用 TF32Tensor Float 32。TF32 是一种特殊精度格式它在使用 Tensor Core 计算时保留一定范围的指数位但尾数精度低于标准 FP32因此性能和精度需要权衡。在 PyTorch 中可以通过环境变量或代码控制 TF32 开关。import torch # 查看默认设置 print(torch.backends.cuda.matmul.allow_tf32) print(torch.backends.cudnn.allow_tf32) # 允许矩阵乘法使用 TF32 torch.backends.cuda.matmul.allow_tf32 True # 允许 cuDNN 卷积使用 TF32 torch.backends.cudnn.allow_tf32 True如果训练脚本对精度要求较高可能需要关闭 TF32如果只是推理TF32 通常能显著提升矩阵计算速度视觉质量影响有限。6.2 llama.cpp 与 GGUF 量化推理本地部署大模型时llama.cpp 是绕不开的项目。它通过 GGUF 模型格式和量化工具把模型权重压缩到 INT8、INT4 等精度。当底层使用 CUDA 后端时会调用 cuBLAS 等库只要输入精度和矩阵尺寸满足条件就会走到 Tensor Core 路径。常见的运行命令不展开展开LLaMA.cpp 的 CMake 和运行参数经常会提示是否启用了 CUDA日志中能看到类似设备信息。如果编译时禁用了 GPU 支持则只能走 CPU 推理Tensor Core 没有意义。6.3 vLLM 等推理服务引擎vLLM、TensorRT-LLM 等服务级推理引擎会把连续请求做动态 batching尽量把矩阵的 batch 维度填满。同时会使用量化感知内核和融合算子让算子尽可能以矩阵分块形式执行。对于这些引擎Tensor Core 通常是默认启用的不需要显式指定。不过具体能否发挥 Tensor Core 性能还取决于模型是否被正确转换到服务引擎的格式。FP16/BF16 权重通常直接利用 Tensor Core而某些自定义量化实现可能走 CUDA Core 路径性能差异较明显。7. 如何验证 Tensor Core 是否在生效很多人以为只要开了 CUDATensor Core 就一定在工作。实际上Tensor Core 的启用取决于三个条件GPU 硬件支持。矩阵计算指令选择了 Tensor Core 路径。精度与数据类型匹配。在本地环境中可以逐步验证。7.1 确认 GPU 支持 Tensor Core用命令或工具查看 GPU 的架构代号和计算能力。nvidia-smi输出的 Name、CUDA Version 能帮你确认识别到的 GPU。要看更详细的计算能力可以在 CUDA 环境下查询或者通过 PyTorchimport torch print(torch.cuda.get_device_name(0)) # 返回类似 (9, 0) 的元组代表 compute capability print(torch.cuda.get_device_capability(0))通常来说NVIDIA GeForce 20 系以后、以及很多专业卡、数据中心卡的架构都提供 Tensor Core。不同架构支持的数据精度范围不同比如较新的架构支持 BF16、FP8而较早架构对 BF16 的支持有限。7.2 观察矩阵乘法的实际数据类型Tensor Core 对 FP16/BF16/TF32/INT8 等输入都有相应执行路径。如果代码里强行把模型、输入转成 FP32并且关闭 TF32那 Tensor Core 基本帮不上忙。在做推理时可以打印模型参数和输入张量的 dtypefor name, param in model.named_parameters(): print(name, param.dtype) break如果看到 FP32而你又希望模型运行在 FP16/BF16 下可以在推理阶段用自动混合精度或者手动半精度。import torch model model.half().cuda() input_ids input_ids.half()7.3 使用分析工具观察 Tensor Pipe Activity只靠nvidia-smi看到 GPU 利用率高不等于 Tensor Core 在运转。CUDA 核心执行大量标量运算时GPU 利用率也可以很高。要精确定位是否使用了 Tensor Core需要使用 Nsight Compute 这类 profiling 工具。查看 SM 内部 Tensor Pipe 的活跃周期是最直接的判断方式。比如关注pipe_tensor相关的硬件计数器统计到的活跃周期越高说明 Tensor Core 参与度越高。命令行运行 Nsight Compute 的通用方式ncu --metrics sm__pipe_tensor_op_hmma_cycles_active.avg.pct_of_peak_sustained_elapsed python run_inference.py这只是典型指标名示例实际指标名称需要根据 CUDA 版本和 GPU 架构做调整。重要的判断标准是Tensor Core 是否执行了 HMMA半精度矩阵乘加、IMMA整型矩阵乘加、FMA、FP8 等指令。分析报告里能看到 kernel 名和对应的指令分布。7.4 通过同尺寸 FP32 与 FP16 推理对比如果不想用复杂的 profiling 工具可以用一个粗略办法把同样的模型推理任务分别用 FP32 和 FP16/BF16 跑一遍比较耗时。如果硬件支持 Tensor CoreFP16 推理通常会有明显优势。如果两者差异很小可能是算子没有走到 Tensor Core 路径或者任务规模太小、数据搬运成为瓶颈。需要提醒这个测试不是标准基准结果会受到模型结构、batch、上下文长度、驱动版本影响。更严谨的做法是使用 PyTorch 自带的 benchmark 脚本或者 vLLM 的性能压测工具。8. 资源占用与性能观察建议8.1 计算密度小的场景先别指望 Tensor Core在本地跑 7B 模型、只做单轮对话、上下文长度很短时单次解码的矩阵形状很小GPU 的计算单元未必能满负荷。这时观察到的性能瓶颈往往在显存带宽、内存拷贝和算子启动开销上。想要感受 Tensor Core 的差异建议用批量请求或长 Prompt。批量增大后矩阵的 M 维度变大计算密度才上得去。8.2 精度越低Tensor Core 优势越明显FP32 计算如果被 CUDA Core 执行性能要打很大折扣。FP16/BF16 和 INT8/INT4 是让 Tensor Core 发挥威力的常见精度。这就是为什么大模型训练普遍使用混合精度而推理服务普遍使用量化模型。量化的目标不只是省显存更是在精度可接受的前提下利用低精度运算单元缩短计算时间。8.3 观察显存占用可以用nvidia-smi实时看显存使用情况watch -n 1 nvidia-smi如果只用nvidia-smi看到显存占用高、GPU-Util 高但不能看到具体算子在干什么。要做算子级观察应该结合 PyTorch Profiler 或 Nsight Systemsnsys profile -o profile_output python run_inference.py分析Nsight Systems生成的时间线可以看到每个 kernel 的启动时间和 GPU 上的耗时占比可以帮助定位哪一类 GEMM 是主要瓶颈。8.4 如何降低显存占用使用量化精度更低的 GGUF、AWQ、GPTQ 模型。减小最大上下文长度控制 KV Cache 占用。使用 Flash Attention 等融合注意力实现减少临时矩阵的显存开销。减少批量大小或关闭多余缓存。这些措施可以间接决定显存里还能放多少数据也会影响矩阵分块的形状。并不是一定降低显存就更好有时显存降低但分块过碎矩阵乘法的效率也会下降。9. 常见问题与排查方法问题现象可能原因排查方式解决方案模型加载后仍很慢权重仍是 FP32 且 TF32 关闭打印模型 dtype用 FP16/BF16 推理或开启 TF32GPU 利用率高但吞吐低Tensor Core 没有参与CUDA Core 在跑标量计算使用 Nsight Compute 查看 kernel 指令分布检查算子实现看是否调用了 cuBLAS 等库的高效路径GPU 显存不足模型权重超过显存容量或上下文过长查看模型权重耗用和 KV Cache 占用降低量化位数、缩短上下文、增大 batch 前先估算FP16 跑起来没有更快矩阵较小搬运开销占主导增大 batch 或上下文再测使用 vLLM 等支持动态 batching 的服务框架BF16 模型无法运行GPU 架构不支持 BF16查看 GPU 计算能力转 FP16 或使用旧架构可用格式更换显卡后性能提升不大驱动/CUDA 版本未正确匹配查看驱动和 CUDA 版本更新驱动、重新编译推理框架量化模型输出质量下降量化精度损失较大对比 FP16 基准输出尝试更低压缩强度或更高精度格式这类问题在大模型本地部署中很常见。重点不是背命令而是形成一套排查路径先确认硬件支持再确认精度类型再确认算子库是否被调用最后通过 profiling 工具观察实际指令。10. Tensor Core 相关最佳实践10.1 在模型侧保持 FP16/BF16 优先能跑 FP16/BF16就不要跑 FP32。不仅省一半显存还能匹配 Tensor Core 的高效计算路径。很多大模型的官方权重默认就是 BF16 或 FP16不要再手动转回 FP32。10.2 批量请求合并本地起一个轻量推理服务时尽量把多个请求合并成 batch。这样 GEMM 的 M 维度够大Tensor Core 的利用率会明显提升。单独起线程反复请求不如一次性打包。# 典型批量推理示例多条 prompt 拼成一个 batch prompts [你好, 讲一下 Tensor Core, 什么是矩阵分块] inputs tokenizer(prompts, return_tensorspt, paddingTrue).to(cuda) outputs model.generate(**inputs, max_new_tokens128)实际部署中vLLM 这类框架内部会做连续 batching。10.3 模型、输入、输出分目录管理本地部署大模型时建议把权重文件、输入测试数据、输出结果分开目录存放。模型文件往往较大离线下载后放在固定目录批量推理脚本只负责读输入、写输出避免文件相互覆盖。models/ llama-7b-q4.gguf inputs/ prompt.txt outputs/ result-001.txt logs/10.4 理解框架日志中的设备信息运行 llama.cpp 或 vLLM 时仔细观察启动日志里的设备编号、显存总量、模型加载精度。很多性能问题在启动日志阶段就会有提示比如没有检测到 GPU、没有以半精度加载、后端错误地回退到 CPU。10.5 谨慎对待精度开关有些框架参数默认关闭混合精度或 TF32另一些默认打开。在训练和微调时TF32 关闭与否会影响精度和收敛结果在纯推理场景TF32 通常可以放宽。建议在测试阶段对比开关前后的输出差异。11. 总结与下一步Tensor Core 加速大模型的核心不是单个硬件有多神秘而是它精准吃透了“大模型 大量矩阵乘加”这一计算特征。矩阵分块负责把大矩阵拆成适合片上缓存处理的子矩阵Tensor Core 负责把子矩阵乘加变成一条专用硬件指令。两者配合FP16/BF16 的大模型推理才能跑出比传统 CUDA Core 高得多的吞吐。建议你做三件事用torch.cuda.get_device_capability确认自己的 GPU 架构。用 FP16 和 FP32 跑同一个模型任务对比推理耗时建立 Tensor Core 是否生效的直接感知。无论用 Ollama、vLLM 还是 llama.cpp多看启动日志和 profiling 工具确认算子真正走了 Tensor Core 路径而不是只听“GPU 利用率 100%”。下一步可以从 Flash Attention 或 CUTLASS 入手进一步看注意力算子的分块设计。矩阵分块只是入口里面的数据布局、Bank Conflict、双缓冲机制才是真正拉开性能差距的细节。这篇内容是原理和部署排查向的入门路线参考实际性能受显卡型号、驱动版本、CUDA 版本、模型架构、量化方式影响应以本机测试为准。