EdgeTransformerEngine
在一张 4 GB 的 NVIDIA T1000 上，把显存占用和模型大小、上下文长度解耦
不调 cuBLAS、不调 CUTLASS、不用 TensorRT，用手写 CUDA kernel 把 Llama-style Transformer 的每一步算完：
RMSNorm、GQA、RoPE、online softmax、SwiGLU 全部自己写，编译成一个 extern "C" 的 DLL，
Python 侧只负责分词和采样循环。
在此基础上做了三件让显存"够用"的事——prefill 路径裁剪、layer streaming、KV cache offload，
最终在一张 4 GB 的 T1000（sm_75） 上，让 9.6B 的 Qwen3.5 只用约 0.5 GB 显存跑出正确结果。
Qwen3.5-9B（int8）
（卡只有 4 GB）
10.2s → 5.2s（200 token prompt）
engine.cu + engine_qwen35.cu
两个大词表矩阵（embed_tokens / lm_head，各 130560×1536）留在主机内存、多线程 fp32 在 CPU 上算；
24 个 transformer block 和 final norm 在 GPU 上。这是 2.16 GB 的 bf16 模型能塞进 4 GB 卡的第一层原因。
token id │ ▼ CPU embedding lookup（大词表不上卡） │ H2D ▼ GPU Layer 0 … Layer N-1 │ RMSNorm → Q/K/V → (Q/K norm) → RoPE │ → KV cache append → causal GQA + online softmax │ → Wo → residual → RMSNorm → SwiGLU → residual ▼ GPU final RMSNorm │ D2H ▼ CPU lm_head GEMV（8 线程） logits[vocab]
| 模式 | GPU 常驻 | 速度 | 适用 |
|---|---|---|---|
| 全常驻（默认） | ~1296 MB 权重 + 完整 KV | ~55 ms/tok（~18 tok/s） | 放得下就用它 |
--resident-layers 2 | ~108 MB 权重 | ~139 ms/tok（~7 tok/s） | 权重放不下 |
--kv-offload | KV 恒定 ~0.5 MB tile | 受主机 KV 传输限制 | 长上下文放不下 |
| 两者叠加 | ~150 MB 总计 | 能跑优先 | 全模型 + 长上下文 |
按"这条知识改变了我怎么写代码"排序，不是按教科书章节排序。每条都附一句我自己的收敛结论。
教科书讲 online softmax 时强调的是减最大值防溢出。真正让我改变写法的是它的另一半性质：
softmax 的状态可以压缩成三个量 (m, l, acc)——当前最大分数、分母累计、加权 V 的分子累计——
并且这三个量可以增量合并。
newm = max(m, score) ecorr = exp(m - newm) // 把旧账按新基准折算 p = exp(score - newm) l = l * ecorr + p acc = acc * ecorr + p * V
一旦承认这一点，"完整 KV cache 必须同时在显存里"这个前提就没了： 可以一次只搬 512 个位置上卡，算完更新状态，再搬下一块。
这是把上面那条数学落成工程的关键一步，也是整个项目我最满意的设计。
单个 kernel 里的寄存器和 shared memory 在 kernel 结束时全部蒸发，
所以要跨多个 tile 累积 softmax 状态，就必须让 m / l / acc 活在 global memory 里，
拆成三个 kernel：
attn_init_kernel m = -1e30, l = 0, acc = 0
↓
for tile in KV: // 每块 512 个位置
cudaMemcpyAsync(host_kv[tile] → d_ktile/d_vtile)
attn_tile_kernel // 原地 rescale m/l/acc
↓
attn_finalize_kernel out = acc / l
结果是 GPU 上的 KV 占用变成常数，和上下文长度彻底无关；增长的只是主机内存。
Layer streaming 让 GPU 只保留 N 个权重槽位（slot = layer % N），
copy_stream 预取第 i+1 层，compute_stream 算第 i 层。
难点不在"开两条 stream"，而在于两个方向的依赖必须都表达出来，少一个就出错：
cudaStreamWaitEvent(compute_stream, copy_done[slot]); // 数据 ready 才能读：防止算到还没拷完的权重 cudaStreamWaitEvent(copy_stream, compute_done[next_slot]); // buffer free 才能写：防止覆盖还在被读的槽位
本质上是在描述一个 buffer 的状态机：空闲 → 正在写 → 可计算 → 正在读 → 可覆盖。
写对了，PCIe 传输就藏进了计算的影子里；每层耗时从
T_copy + T_compute 降到 max(T_copy, T_compute)。
T1000 是 sm_75，没有高效的原生 bf16 运算。所以策略是bf16 只当存储格式：
权重和 KV cache 存 bf16（省一半），kernel 读进来立刻 __bfloat162float，
所有点积和 softmax 用 fp32 累加。
bf16↔fp32 的转换之所以是一次移位而不是查表，是因为 bf16 就是 fp32 砍掉低 16 位尾数——指数位宽完全一样：
FP32: [sign][8-bit exponent][23-bit mantissa] BF16: [sign][8-bit exponent][ 7-bit mantissa] uint32_t x = (uint32_t)h << 16; // 就这一行
这也是 bf16 在训练里赢过 fp16 的原因：动态范围和 fp32 相同，代价只是精度， 而长点积用 fp32 累加就把精度补回来了。
batch=1 的自回归解码，每个线性层都是 GEMV 而非 GEMM： 每个权重元素只被用一次，算术强度接近 1。这意味着堆 FLOPS 没有用，瓶颈全在搬数据。
这条认知直接解释了本项目里三个看似无关的实测现象：
- ·
--resident-layers 2和3测出同样的速度——因为路径是 PCIe 传输 bound，多留一层权重不改变每 token 要搬的总字节数。所以 N=2 是显存甜点。 - ·int8 量化带来 约 2 倍加速——不是因为 int8 算得快，而是每 token 搬运量减半。
- ·CPU lm_head 每个 token 要扫过整个 130560×1536 的 bf16 矩阵（~400 MB）。按典型双通道内存带宽估算，这一步单独就是几十 ms——和实测的 55 ms/tok 同一量级。全常驻模式下，时间主要花在 CPU lm_head，不在 GPU。
Qwen3.5-9B 的 bf16 导出是 17.9 GB，装不进 15 GB 内存，只能从硬盘 mmap 流式——很慢。 做了 int8 对称 per-row 量化后降到 8.4 GB，整模型常驻内存，硬盘彻底退出链路。
关键不在于"量化到 int8"，而在于哪些张量不能量化：
大矩阵 + embed + lm_head 走 int8；而
A_log / dt_bias / conv1d / 各 norm / in_proj_a·b
这些 SSM 敏感小张量保留 f32。它们参数量占比极小，但直接进指数/门控，误差会沿时间维累积放大。
验证方式是硬指标而非观感：贪心解码的 token id 序列与 bf16 逐字节一致。
移植 Qwen3.5 的混合 Gated-DeltaNet 线性注意力时，没有直接写 CUDA， 而是搭了一条每一环都能单独证伪的链：
HuggingFace 真实前向 ── 生成 golden（tiny 随机模型 + logits + 权重）
↓ 对比
纯 numpy fp32 复现 ── max|numpy − HF| ≈ 7.6e-7 → 架构确实读懂了
↓ 对比
独立 CUDA 前向 ── max|CUDA − HF| ≈ 7.9e-7 → 算子写对了
↓ 对比
生产 DLL 增量推理 ── f32 9.5e-7；int8 argmax 全中 → 状态缓存/驻留模式都没错
每一环出问题，都能立刻定位到是"理解错了架构"还是"CUDA 写错了"还是"工程链路错了"。
贯穿整个项目的一条主线：GPU 显存从来不是被"节省"了，只是被换成了别的东西——
| 手段 | 省下 | 付出 |
|---|---|---|
| 大词表放 CPU | ~800 MB 显存 | 每 token 一次 CPU 内存全扫 |
| Layer streaming | 权重显存 ∝ 层数 → ∝ N | 每 token 走一遍 PCIe |
| KV offload | KV 显存 ∝ 上下文 → 常数 | 每层每 token 重传历史 KV |
| int8 量化 | 体积减半 | 精度（此处实测无损） |
注意最后一行的不同：前三行都是拿时间换空间，只有量化是拿精度换空间—— 而当路径是带宽 bound 时，量化同时换来了时间。这就是为什么 int8 是这里唯一"两头赚"的手段， 也是下一步该做 int4 的理由。
每条都对应仓库里一个可复跑的脚本，判据是可证伪的硬指标（token id 逐位一致 / 数值误差上界），不是"输出看着挺像"。
server.py），支持 SSE 流式，任意 OpenAI 客户端可直接调用。
lm_head GEMV 是单 token 成本的大头。
于是 engine_prefill 跑完 24 层填 KV cache，但跳过 final RMSNorm 和 lm_head。
200 token 的 prompt 上首 token 从 ~10.2 s 降到 ~5.2 s。
--resident-layers 2 时 GPU 权重从 ~1296 MB 降到 ~108 MB，输出与全常驻模式逐位相同。
N=3（162 MB）实测与 N=2 同速——印证了路径是 PCIe bound，因此 N=2 是显存甜点。
[242, 242, 178, 137, 244, 178]，GPU KV 占用保持 ~0.5 MB。
--resident-layers 2 --kv-offload 叠加后，完整模型加长上下文一共只占用约 150 MB 显存。
这是项目的主结果：显存占用同时和模型深度、上下文长度解耦。
visual.* / mtp.* 只导文本子图；
int8 后 8.4 GB 整个常驻内存。显存账：双缓冲权重 2 槽 ~419 MB + 24 个线性层的定长 state cache 50 MB
+ 8 个全注意力层的 KV（32 KB/token）= ~0.5 GB。
实测 "The capital of France is" → " Paris."，
token id 与 bf16 逐字节一致；纯流式 + 双缓冲 ~1.41 tok/s。
--resident-layers 5 时横幅显示 weights on GPU=1247 MB (5 resident + 27 streamed)，输出仍正确。
7.6e-7 → 独立 CUDA 前向 vs HF 7.9e-7 → 生产 DLL vs HF 9.5e-7；
三种驻留模式（全常驻 -1 / 纯流式 0 / 混合 N）输出逐位一致。
这是一个教学导向的引擎，不是 vLLM 的替代品。下面这些不是"还没来得及做"，是我读完自己代码后能明确指出的边界。
当前局限
- batch=1 且 prefill 逐 token。成熟引擎的 prefill 是 GEMM + FlashAttention 一次吃整段；这里对每个 prompt token 单独跑一遍全网络，长 prompt 代价大。
- GEMV 没有向量化 load / warp shuffle / Tensor Core，性能明显低于 cuBLAS。
- KV offload 没有 tile 双缓冲：拷贝和
attn_tile_kernel都在同一条 compute stream、只有一套 tile buffer，所以是"拷一块算一块"，没有重叠。而 layer streaming 的双缓冲代码就在旁边，可以直接复用。 head_dim有未校验的隐含约束：attention 和 qk_norm 用了固定__shared__ float s[256]和折半归约，要求 head_dim ≤ 256 且为 2 的幂，但加载路径没有检查——喂错模型会静默算错。- 部分 capability 读了但没接线：
sliding_window、attn_logit_softcap等字段从权重文件读入却未参与计算。喂 Gemma 权重不会报错，会静默给出错误结果——这比崩溃危险。 - CUDA 错误检查不完整：kernel launch、
cudaMemcpyAsync、event/stream 创建大多没查返回值。 - 只有贪心采样，无 temperature / top-p；
-arch固定 sm_75。
下一步（按性价比排序）
- 先量再改：上 Nsight Systems 看两条 stream 的真实重叠率，确认 compute 到底在不在
copy_done上等——现在的流水线收益是推算的，不是测出来的。 - int4 量化：路径是带宽 bound，每 token 搬运量再减半是当前最大的单点杠杆。
- KV tile 双缓冲：直接复用 layer streaming 的 event 模式，成本低、收益确定。
- CPU lm_head 提速：持久线程池（现在每 token 建销线程）+ SIMD 点积 + lm_head 量化。按第 05 条的估算，这是全常驻模式下最大的一块时间。
- 批量 prefill：把
[1, hidden]的 GEMV 升级成[seq_len, hidden]的 GEMM，配 causal mask 和 FlashAttention 式 prefill。这是从教学引擎走向可用引擎的分水岭。 - 加载期 capability 门禁：遇到未实现的 capability 直接拒绝加载，把"静默算错"变成"明确报错"。
我日常做的是 Edge 里的 agent memory——偏系统与架构。做这个项目是想把 "模型跑起来"这件事从推理框架的黑箱里拆出来，亲手验证一遍： 数学公式怎么变成 kernel、KV cache 为什么吃显存、长上下文的代价具体贵在哪、 权重放不下时 stream 到底怎么调度才不出错。
最后留下的不是这两千行 CUDA，而是一个可复用的判断顺序： 先算算术强度定瓶颈 → 再想状态住在哪层存储 → 最后用可证伪的参照链守住正确性。 现在再看任何一篇推理优化的论文或框架，我知道该先去看它的哪一部分。