企业官网建设流程全解析

这次我们不聊具体某一个大模型，直接把镜头拉到 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 指令，但你要知道什么情况下框架会调用它，什么情况下不会。

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×N：

C[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 的 WMMA（Warp Matrix Multiply-Accumulate）编程模型，可以让一个线程束（Warp，32 个线程）共同完成一个较大的矩阵乘加操作，而不是每个线程独立做一个元素的乘加。

这种模式的好处是：

• 一个线程束内的线程通过寄存器协作，共同维护一个大的矩阵分块。
• 单个时钟周期内硬件完成大量乘加运算，而非只是单个标量运算。
• 数据不需要反复经过通用计算单元，由专用矩阵计算单元直接计算。

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 高性能执行的关键不只是计算单元，而是数据调动路径。典型路径是：

1. 把 A 和 B 的分块从全局内存加载到共享内存（Shared Memory）。
2. 从共享内存把数据分发到线程的寄存器中，形成 Tensor Core 需要的矩阵片段。
3. 调用 Tensor Core 指令完成子矩阵乘加。
4. 把累加结果写回共享内存，再刷回全局内存。
5. 继续遍历后续分块，直到完成整个输出矩阵。

因此，高效启动一个 GEMM 时，GPU 内部有大量线程在协同工作：一部分负责数据搬运，一部分负责计算，一部分负责写回结果。所有的工作都围绕分块展开。Tensor Core 把“计算密度高”的那一步加速了，如果数据搬运没有跟上，整体性能依然会被内存带宽卡住。

5. Tensor Core 与显存占用、模型规模的关系

5.1 权重精度决定显存下限

大模型参数是固定的，但存储精度可以选择。模型权重以 FP32 存储，一个参数占 4 字节；以 FP16/BF16 存储，占 2 字节；以 INT8 量化存储，占 1 字节；以 INT4 量化大约占 0.5 字节。

按 7B 参数量估算：

这只是权重本身。实际运行还需要加载 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 生成）：每次只生成一个新 Token，batch 往往较小，矩阵形状较“瘦长”。如果单请求推理，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 输入，则可能使用 TF32（Tensor 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 是否在生效

很多人以为只要开了 CUDA，Tensor Core 就一定在工作。实际上，Tensor Core 的启用取决于三个条件：

1. GPU 硬件支持。
2. 矩阵计算指令选择了 Tensor Core 路径。
3. 精度与数据类型匹配。

在本地环境中，可以逐步验证。

7.1 确认 GPU 支持 Tensor Core

用命令或工具查看 GPU 的架构代号和计算能力。

nvidia-smi

输出的 Name、CUDA Version 能帮你确认识别到的 GPU。要看更详细的计算能力，可以在 CUDA 环境下查询，或者通过 PyTorch：

import 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 基本帮不上忙。

在做推理时，可以打印模型参数和输入张量的 dtype：

for 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 Core，FP16 推理通常会有明显优势。如果两者差异很小，可能是算子没有走到 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 Systems：

nsys profile -o profile_output python run_inference.py

分析Nsight Systems生成的时间线，可以看到每个 kernel 的启动时间和 GPU 上的耗时占比，可以帮助定位哪一类 GEMM 是主要瓶颈。

8.4 如何降低显存占用

• 使用量化精度更低的 GGUF、AWQ、GPTQ 模型。
• 减小最大上下文长度，控制 KV Cache 占用。
• 使用 Flash Attention 等融合注意力实现，减少临时矩阵的显存开销。
• 减少批量大小或关闭多余缓存。

这些措施可以间接决定显存里还能放多少数据，也会影响矩阵分块的形状。并不是一定降低显存就更好，有时显存降低但分块过碎，矩阵乘法的效率也会下降。

9. 常见问题与排查方法

这类问题在大模型本地部署中很常见。重点不是背命令，而是形成一套排查路径：先确认硬件支持，再确认精度类型，再确认算子库是否被调用，最后通过 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_tensors="pt", padding=True).to("cuda") outputs = model.generate(**inputs, max_new_tokens=128)

实际部署中，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 高得多的吞吐。

建议你做三件事：

1. 用torch.cuda.get_device_capability确认自己的 GPU 架构。
2. 用 FP16 和 FP32 跑同一个模型任务，对比推理耗时，建立 Tensor Core 是否生效的直接感知。
3. 无论用 Ollama、vLLM 还是 llama.cpp，多看启动日志和 profiling 工具，确认算子真正走了 Tensor Core 路径，而不是只听“GPU 利用率 100%”。

下一步可以从 Flash Attention 或 CUTLASS 入手，进一步看注意力算子的分块设计。矩阵分块只是入口，里面的数据布局、Bank Conflict、双缓冲机制，才是真正拉开性能差距的细节。

这篇内容是原理和部署排查向的入门路线参考，实际性能受显卡型号、驱动版本、CUDA 版本、模型架构、量化方式影响，应以本机测试为准。