SYSTEMS · INFERENCE DATAFLOW

一次大模型推理不只是在算，更是在搬数据

沿一次请求看清三类东西：长期驻留的权重（weights）、算子之间流动的激活（activations），以及逐层增长的键值缓存（Key–Value Cache, KV Cache）。它们在冷启动、预填充（Prefill）与逐 token 解码（Decode）三个时间尺度里走不同路线，也因此形成不同瓶颈。

COLD START权重加载一次checkpoint → device memory PER REQUESTPrefill 建历史批量计算并逐层写 KV PER TOKEN STEPDecode 反复读权重 + 历史 KV +（多 GPU 时）通信

发光的数据块从远端存储流向 GPU 大容量显存，再进入更靠近计算单元的片上工作区 — **直觉总览：**存储像总仓，GPU device memory 像大仓库，片上 cache、shared memory、register 或 TMEM 像工作台。图表达“把工作集搬近并复用”的方向，不代表每个字节必须按同一固定阶梯经过。

00 · 数据流全景

从存储到计算单元，一次推理经过哪些地方

在普通 x86 + 独立 GPU 服务器上，模型加载通常经过存储与 CPU 内存，再由直接内存访问（Direct Memory Access, DMA）通过 PCI Express（PCIe）进入 GPU；计算时，数据在 device memory、cache 与流式多处理器（Streaming Multiprocessor, SM）的片上工作区之间流动和复用。满足条件的系统也可用 GPUDirect Storage（GDS）避开 CPU bounce buffer。

大模型推理数据流：存储、CPU 内存、GPU device memory、L2、SM 内 L1 shared memory、寄存器与 Tensor 或 CUDA 计算单元，并标出 GPUDirect Storage 旁路 — 这是普通 x86 独立 GPU 的逻辑路线；cache 命中与架构相关工作区不是每个字节都必须逐级经过的固定物理阶梯。

性能观察

理解推理性能，需要同时关注数据放在哪里、每轮逻辑上要访问多少、物理上从哪一级命中、能复用几次，以及通信能否与计算重叠。

别混两类路径

PCIe / NVLink 决定设备之间怎样传；PagedAttention 一类机制组织 KV Cache 的显存 block。前者改变通信链路，后者主要改变分配、共享与可承载并发，不会凭空消除 dense attention 对历史 KV 的访问。

01 · 模型加载

权重先“驻场”，请求来了通常不再读 SSD

推理服务启动时，框架读取 safetensors 等 checkpoint shards（检查点分片），完成数据类型（data type, dtype）、分片、量化格式或 kernel 布局转换，再把权重放到目标设备。加载完成后，常用权重通常长期驻留 device memory；CPU offload（卸载）、按层流式加载和稀疏专家分级存储，是容量不足时用额外传输换空间的设计。

模型加载三条路径：普通路径由 SSD 或网络存储经过 CPU DRAM 和 DMA 进入 GPU HBM；GPUDirect Storage 在条件满足时直接进入 HBM；offload 在 CPU 或 NVMe 与 GPU HBM 间换入换出 — 蓝线是常见启动加载，橙线是有条件的 GDS 旁路，绿线是容量不足时付延迟换空间的运行时 offload。

权重容量的第一笔账

dense BF16 / FP16

7\times 10^9\ \text{parameters}\times 2\ \text{bytes}\approx 14\ \text{GB}\approx 13.0\ \text{GiB}

Brain Floating Point 16（BF16，脑浮点 16 位）与 Half-Precision Floating-Point 16（FP16，半精度浮点）在这里都按每参数 2 bytes 粗算。“约 14GB”是十进制结果，换成二进制约 13.0GiB。真实显存还要加量化 scale/metadata、对齐、临时 workspace、CUDA graph、激活和 KV Cache；若 embedding 与 LM head 共享权重，也不能重复计数。

加载路线	适用场景	运行特点
Storage → CPU buffer → GPU	通用框架、普通文件系统与独立 GPU 服务器	由 DMA 与 runtime 协调传输，通用性最好
Storage → GPU memory	满足 GPUDirect Storage 条件的系统	减少 CPU memory 中转；仍依赖存储、文件系统、驱动、拓扑与对齐配置
CPU / NVMe offload	权重或 KV 放不进 GPU，需要用延迟换容量	扩大可用容量，同时增加 PCIe 传输；实现可能使用显式 copy / prefetch，也可能触发 page migration 或 fault，抖动风险不同

02 · 输入与 Embedding

文字先变 ID，ID 再变隐藏向量

用户文本经过 Chat Template 和 Tokenizer 变成 Token IDs，通过主机到设备传输进入 GPU，再从 Embedding Matrix 查表得到 Hidden States — Token ID 只是索引；真正进入 Transformer 的是 embedding 查表后形成的三维 hidden-state tensor。

CPU 侧常见工作

Tokenizer 把文本切成 token 并映射为整数 ID，还会拼 chat template、special tokens 与 position。多数服务在 CPU 完成；GPU tokenizer 或专用预处理流水线也可以承担这部分工作。

GPU 侧常见工作

token IDs 很小，传入 GPU 后对 embedding matrix 做 gather，得到形如 [batch, sequence, hidden_size] 的 hidden states。4096 hidden size 表示每 token 有 4096 个数，不代表它们分别拥有固定的人类语义。

输入小，状态大

1000 个 32-bit Integer（INT32，32 位整数）token IDs 只有约 4KB；若 hidden size=4096、BF16，则 1000 个 token 的一层 hidden states 约 1000×4096×2 ≈ 8.2MB（未计 batch）。真正的大头通常是权重、层间激活、attention 工作区和逐层 KV Cache。

03 · Transformer 层

一层 Transformer 由 Attention 与多层感知机（Multi-Layer Perceptron, MLP）两条分支组成

现代 decoder-only Large Language Model（LLM，大语言模型）常采用前置归一化（Pre-Norm）：attention 前进行一次 Root Mean Square Layer Normalization（RMSNorm，均方根层归一化）或 Layer Normalization（LayerNorm，层归一化），MLP 前再进行一次。数据沿 residual stream（残差主干）穿过整层，两条分支分别完成 token mixing（位置间混合）与 channel mixing（通道间混合）；具体模型也可能使用并行分支、混合注意力或不同 Norm。

X_\ell\xrightarrow{\mathrm{Norm\,+\,Attention}}A_\ell=X_\ell+\mathrm{Attn}(\mathrm{Norm}(X_\ell))\xrightarrow{\mathrm{Norm\,+\,MLP}}X_{\ell+1}=A_\ell+\mathrm{MLP}(\mathrm{Norm}(A_\ell))

现代 decoder-only Transformer 单层数据流：Residual Stream 进入 RMSNorm、QKV Projection、Attention、输出投影和残差加法，再进入第二次 RMSNorm、MLP 或 MoE、第二次残差加法，并与 KV Cache 读写相连 — 蓝色上半段混合 token，绿色下半段混合 channel；两条 residual bypass 才是贯穿层间的主干。

①

Attention 分支

mix tokens

归一化后的 X 与 Wq/Wk/Wv 做投影得到 Q/K/V；位置编码常作用于 Q/K。Prefill 会为 prompt 各位置生成 K/V，Decode 则为新 token 生成一小段 K/V 并追加缓存。attention 输出经输出投影后与 residual stream 相加。

②

MLP / MoE 分支

mix channels

第二次归一化后进入大矩阵乘法与门控激活。Dense MLP 通常是权重与 Floating-Point Operations（FLOPs，浮点运算次数）大户；Mixture of Experts（MoE，混合专家）只执行被路由专家，但所有专家的存储、跨卡 token dispatch 和 shared expert 仍要计入系统账。

③

逐层向前

residual stream

当前层输出成为下一层输入。高效实现会复用 buffer、原地执行部分算子并进行 kernel fusion，让更多中间结果留在片上或直接被下一个算子消费，从而减少 HBM 往返。

04 · GPU 芯片内

矩阵太大装不进片上，就切成 tile 流水计算

对 Y=XW，kernel 会把工作拆成 tile（数据块），让同一块 X/W 尽量服务多次乘加。共同逻辑是取数 → 片上暂存/复用 → 矩阵乘加 → 写回；但操作数和累加器放在 register、shared memory 还是 Tensor Memory（TMEM），取决于 GPU 代际、指令与 kernel，不能画成一条跨架构不变的硬件管线。

GPU 矩阵乘 tile 的共同数据流与架构差异：global memory 请求由 cache 命中或 device memory 服务，kernel 可把 tile 暂存在 shared memory；Ampere、Hopper 与 Blackwell 分别使用不同的寄存器、shared memory、TMA 与 TMEM 放置方式，完成结果再写回 global memory — **共同逻辑不等于固定放置：**Hopper 的 Tensor Memory Accelerator（TMA）可在 global 与 shared memory 间搬 tensor 而不借寄存器中转；Blackwell 的 TMEM 是另一块专用于 Tensor Core 数据、尤其累加器的片上存储，TMA 与 TMEM 不是同一件事。

层级	谁能看到 / 谁管理	推理里常放什么
Device memory（HBM/GDDR）	所有 SM 可访问；容量最大、离计算最远	权重、激活、KV Cache、工作区
L2 cache	芯片级、硬件管理	跨 SM 或连续 kernel 可能复用的数据
L1 / Shared memory	每个 SM；现代 NVIDIA 架构常共享可配置的物理资源，但 L1 是硬件 cache、shared 是显式可编程空间	GEMM/attention tile、归约中间量；Hopper TMA 可异步搬入/搬出 shared
Registers	每线程私有的逻辑寄存器；不足时会 spill，而 CUDA local memory 的物理存储仍在 device memory	地址、标量与部分指令族的 operand / accumulator fragments
TMEM（Blackwell）	Blackwell SM100 新增的 Tensor Memory；由 `tcgen05` Tensor Core 指令直接访问	累加器固定在 TMEM，operand A 也可来自 TMEM；operand B 来自 shared memory
Tensor / CUDA cores	计算执行单元	矩阵乘加、向量与标量运算

真实执行过程：global address 的 load 可能命中 L2/L1，也可能访问 High Bandwidth Memory（高带宽显存，HBM）或 Graphics Double Data Rate memory（图形双倍数据率显存，GDDR）。kernel 可用异步预取、double buffer（双缓冲）和 warp specialization（warp 专职分工）重叠搬运与计算，也可能因 register 压力 spill。要确认物理字节从哪一级来，应看 Nsight Compute 等 profiler 的 memory transactions，而不是只看源码中的逻辑 load。

术语范围：本节用 NVIDIA CUDA / Tensor Core 命名讲解；AMD、Intel 及专用加速器也有片上缓存、可编程工作区与矩阵单元，但层级名称、可见性、指令和 profiler 指标并不相同，不能直接把 TMA/TMEM 套过去。

05 · Prefill / Decode

同一份权重，为什么“读题”和“答题”瓶颈不同

Prefill 与 Decode 逐层数据流：两阶段都在每一层的 QKV 投影中产生 K 和 V 并写入该层 KV Cache；只有完成全部 Transformer 层后才经 LM Head 采样，采样 token 只作为下一轮 Decode 的输入 — Prefill 的复用来自“一个权重 tile 服务多个 prompt positions”；Decode 的循环依赖让每轮新增位置很少。注意：KV 在逐层 QKV 分支写入，不是采样后才生成。

Prefill · 常更偏 compute-bound

整段 prompt 的多个 token 可组成较大的通用矩阵乘（General Matrix Multiplication, GEMM），同一块权重能服务很多 token，算术强度通常更高；同时生成 prompt 的各层 KV。实际瓶颈会随上下文长度、batch、模型形状、kernel 效率和主机输入流水线变化。

Decode · 小 batch 常更偏 memory-bound

每轮只新增一个或少量 token，却仍要穿过所有层并逻辑上访问大量权重与历史 KV。连续批处理把多个请求合成更大的矩阵后，可以摊薄权重流量，瓶颈也可能转向计算、KV 或通信；“Prefill 算力受限、Decode 带宽受限”是常见区域，不是阶段定义。

算

7B dense BF16 的理想“搬权重”下界

bandwidth estimate

\text{ideal lower bound}\approx\frac{14\ \text{GB}}{3.35\ \text{TB/s}}\approx 4.2\ \text{ms}

用 H100 SXM 的官方 3.35TB/s 峰值粗算，若 batch=1、7B dense BF16 的一次 decode step 从 HBM 近似读取 14GB 权重，仅这笔流量的理想下界约 4.2ms，对应约 239 step/s 上限。真实执行还要付 KV、激活、kernel launch、同步和带宽未达峰值等成本；若权重命中 cache、采用量化、被张量并行分片或换成 MoE，分子与链路都会改变。这是教学 Roofline 下界，不是 H100 实测吞吐。

batch=1

14GB 权重主要服务 1 个 token position，复用低，容易“搬得多、算得少”。

batch=8 的直觉

同一轮大致仍读取一份权重，却服务 8 个 token positions；粗看权重流量摊成约 1.75GB/token。实际收益取决于 batch 是否把 GEMM 做大、KV 长度差异和调度开销。

06 · Attention 与 KV Cache

KV Cache 省了重算，却新增了持续增长的状态

自回归 decode 不必重算历史 token 的 K/V，但每层 dense attention 仍要让当前 Q 与允许访问的历史 K/V 交互，再把新 K/V 追加进去。序列越长、并发越高，这部分容量与逻辑访问量越重要；Multi-Query Attention（多查询注意力，MQA）、Grouped-Query Attention（分组查询注意力，GQA）、Multi-head Latent Attention（多头潜在注意力，MLA）、滑窗、稀疏注意力与 KV 量化会改变这笔账。

\text{KV bytes / sequence}=L\times S\times H_{kv}\times D\times 2_{K,V}\times b

Decode 第 t 步的完整 Attention 链：当前 hidden state 投影为 Q(t)、K(t)、V(t)，Q(t) 与 K[1:t] 计算 scores，经因果 mask 和 softmax 得到权重，再与 V[1:t] 做加权和得到当前输出；K(t) 与 V(t) 同时追加到缓存尾部 — 绿色实线读取 K[1:t] / V[1:t]，绿色虚线写入当前 K(t) / V(t)；压 KV、少读 block、改善 block 分配是三类不同优化。

32K GQA 手算

32 层、32K tokens、8 KV heads、head dimension 128、BF16：32×32768×8×128×2×2 = 4GiB / sequence。这是未计 batch、block 尾部、allocator metadata、跨 rank 分片与低精度 scale 的逻辑容量；dense decode 每步逻辑上访问允许范围内的大部分历史 KV，物理 HBM 字节还受 cache、tiling 与复用影响。

KV Cache 的时间复杂度

它把“每步重算全部历史 K/V”变成“复用历史 K/V”；在 dense attention 中，对历史缓存的读取和点积仍会随序列长度近似线性增长。更详细推导见 KV Cache 专题。

技术	主要优化对象	机制特点
FlashAttention	dense attention 的中间矩阵 HBM 读写，长序列训练/prefill 尤其明显	通过 tiling + online softmax 减少中间结果往返，同时保持同一 dense attention 定义；decode 常配合 KV-aware kernel
PagedAttention	连续预留、外部碎片、重复前缀与动态增长	以固定大小 block 管理 KV，并可做块级共享/copy-on-write；最后一个未填满 block 仍有尾部浪费，block table 也有间接寻址成本
KV 量化	每个 K/V 元素的存储字节与读取带宽	通过降低数值位宽缩小缓存，并配合量化 kernel 完成计算

07 · 多 GPU

模型一拆开，数据流里就多了“通信税”

多 GPU 推理通信图：流水线并行在 stage 边界传 activation；张量并行中所有 rank 都向 All-Reduce 提供 partial result 并各自接收完整结果；上下文并行交换 K/V blocks；跨节点数据依次经过 GPU A、NIC A、InfiniBand 或 RoCE fabric、NIC B 和 GPU B — Pipeline Parallel（流水线并行，PP）搬 stage activation；Tensor Parallel（张量并行，TP）的 ranks 参与 collective；Context Parallel（上下文并行，CP）交换 K/V 或 attention 中间量；跨机数据仍经过两端网络接口卡（Network Interface Controller, NIC）与 network fabric。

并行方式	什么被切开	推理时常见通信
Pipeline Parallel（PP）	不同 GPU 放不同层	stage 边界传 activation；单请求串行延迟会叠加，batch/micro-batch 才能填流水线
Tensor Parallel（TP）	同一层的矩阵/heads 分片	每层根据切分方式使用 all-reduce、all-gather 或 reduce-scatter
Expert Parallel（EP）	Mixture-of-Experts（混合专家，MoE）的 experts 分布到多卡	router 后 token dispatch/combine 常形成 all-to-all
Context / Sequence Parallel（CP / SP）	长序列或激活沿序列维切分	交换 K/V、部分 attention 结果或归一化统计，取决于算法，CP 与 SP 也不是同义词

单机

GPU↔GPU 可能走 NVLink/NVSwitch，也可能只走 PCIe；拓扑决定带宽和延迟。CPU↔GPU 在常规 x86 服务器多走 PCIe，Grace Hopper/Blackwell Superchip 才提供 memory-coherent（内存一致性）NVLink Chip-to-Chip（NVLink-C2C）。

跨机

数据经 NIC 与 InfiniBand 或 RDMA over Converged Ethernet（融合以太网上的 RDMA，RoCE）。满足 GPUDirect Remote Direct Memory Access（远程直接内存访问，RDMA）条件时，NIC 可 DMA 到 GPU memory，避免 CPU bounce buffer；仍需正确拓扑、driver、memory registration 与通信库。

08 · 优化地图

每种优化，其实在改不同的一笔“数据账”

推理优化与数据成本的四条独立映射：量化减少权重与 KV 每元素字节数；算子融合和 FlashAttention 减少 HBM 往返；PagedAttention 改善 KV block 分配并提升有效 batch 容量；连续批处理和投机解码提高每次权重读取服务的 token 数 — 先定位最贵的数据链，再选择减少字节、减少落盘、减少碎片或提高复用的工具。

方法	主要动作	最直接影响
INT8 / INT4 / FP8 / FP4	降低权重、激活或 KV 的 bytes/element；前提是 kernel 真以低比特搬运并高效计算	容量 + 带宽，有时也提升 Tensor Core 吞吐
Kernel fusion	把 Norm、bias、activation、residual 等连续算子合并，减少中间量落回 HBM与 launch	HBM traffic + launch overhead
FlashAttention	attention tiling、online softmax、用重计算换输入/输出（Input/Output, I/O）流量	attention 中间矩阵 HBM traffic
PagedAttention	以块管理 KV，并支持 copy-on-write / prefix sharing	显存利用率、可承载 batch 与共享
Continuous batching	动态把不同请求的 decode token 合并到更大的 step	权重复用与系统吞吐；未必降低单请求尾延迟
Speculative decoding	一次目标模型前传验证多个候选 token	把一次权重读取摊到多个被接受 token
Prefill / Decode disaggregation	把两阶段放到不同 engine，并把 Prefill 产出的 KV 传给 Decode	可分别调 TTFT 与 ITL、隔离尾延迟，但新增 KV 传输与路由；不自动提高吞吐，vLLM 稳定版文档在本页核查日仍明确其实现不提升 throughput

最后只记五问

逻辑上访问什么？物理上从哪一级来？要搬多少？能复用几次？搬运能否和计算重叠？ 这五问比“某 GPU 有多少 Tera Floating-Point Operations per Second（TFLOPS，每秒万亿次浮点运算）”更能解释真实推理性能；最终仍要用 profiler、Roofline、端到端 Time to First Token（首 token 延迟，TTFT）、Inter-Token Latency（相邻 token 延迟，ITL）/Time per Output Token（每输出 token 时间，TPOT）、throughput 与目标 workload 验证。

来源与核查

把硬件路径、算法机制与 Serving 行为分开取证

本页于 2026-07-16 重新核对 CUDA 13.3、Hopper/Blackwell 指令路径、当前 KV Cache 系统与 vLLM 稳定版 P/D 分离文档。硬件峰值只用于明确假设下的教学下界；cache 命中、实际 HBM bytes、通信重叠与端到端延迟必须在目标模型、GPU、并发和软件版本上重新 profile。

主题	一手来源	本页采用的边界
CUDA memory model	CUDA Programming Guide 13.3；CUDA Best Practices Guide	host/device address space、global/shared/register/local memory、统一 L1/Shared 资源与 L2 的语义；逻辑 load 不等于固定物理阶梯
Ampere / Hopper / Blackwell tile path	Ampere Tuning Guide 13.3；CUTLASS Warp-Level MMA guide；Hopper Tuning Guide 13.3；CUTLASS tcgen05 guide	Ampere 的代表性 operand/accumulator register fragments、Hopper TMA 的 global↔shared 搬运与 Blackwell TMEM 累加器分别取证，不拼成一条固定阶梯
Pre-Norm decoder layer	LLaMA paper；RMSNorm paper	用代表性 decoder-only 架构说明两次子层前归一化与残差主干；公式是教学抽象，不外推到并行分支或全部模型
模型加载与 GDS	Safetensors docs；GPUDirect Storage Overview	checkpoint 分片与部分 tensor 读取；GDS 是有硬件、文件系统、驱动和拓扑前提的 direct DMA path，不是默认路径
CPU↔GPU 与 H100 数量级	Grace Performance Guide；H100 官方规格	普通 x86 PCIe 与 Grace NVLink-C2C 分账；H100 SXM 3.35TB/s 只用于 14GB 权重的峰值带宽下界
Attention I/O	FlashAttention paper	dense exact attention 通过 tiling 与 online softmax 减少 HBM↔片上静态随机存取存储器（Static Random-Access Memory, SRAM）往返，不等于把 token-pair 关系变少
Paged KV 与复用	PagedAttention paper；TensorRT-LLM KV Cache System	block allocation、尾块浪费、prefix sharing、copy-on-write、offload 与当前 pool/eviction 限制；分页主要改内存管理，不直接删掉 dense KV 访问
Prefill / Decode 与 P/D 分离	DistServe paper；vLLM stable · Disaggregated Prefilling	两阶段可分别匹配资源与服务等级目标（Service Level Objective, SLO），但要传 KV；vLLM 稳定版页面在 2026-07-16 仍标为 experimental，并明确其实现不提高 throughput
多 GPU / 跨节点通信	NVIDIA Collective Communications Library（NCCL）2.30.7 Collectives；GPUDirect Remote Direct Memory Access（RDMA，远程直接内存访问）	collective 语义与所有 rank 的参与合同；RDMA 可避开 CPU bounce buffer，但不会消除 Network Interface Card（NIC，网卡）、fabric、registration 与拓扑成本

计算说明：文中的 14GB、4.2ms、239 step/s 与 4GiB 均为明确假设下的教学手算，用来建立数量级直觉；GB/TB 使用十进制，GiB 使用二进制。任何厂商峰值、论文 throughput 或框架行为都不能脱离 SKU、dtype、模型、长度分布、batch、SLO 与版本外推。