惊呆了！25GB 内存上运行 GLM 5.2 744B 模型
数
数据STUDIO约 40 分钟
数据STUDIO约 40 分钟GLM 5.2 是一个 7440 亿参数的 Mixture of Experts（MoE）模型。按常规思路，它应该需要一整排数据中心 GPU；但 MoE 的关键不在于“所有参数是否同时参与计算”，而在于每个 token 到底会激活多少参数。
我发现有个大佬用纯 C 写了一个推理引擎：把 GLM 5.2 的路由专家转成 int4，按需从磁盘读取，再配合压缩 KV cache、专家路由和分层缓存，在约 20GB RAM 的 CPU-only 配置下运行。是真的夯。长文警告，可收藏后阅读。
实现路径可以压缩成几件事：
先检查模型和机器，再从一个 C 文件构建引擎； 用约 400 行的更小 streaming engine，先和 PyTorch oracle 逐 token 对齐； 把 FP8 权重解码为 float32，再按行量化并打包成 int4； 通过 pread、safetensors 索引和posix_fadvise，只读取当前需要的专家；用 Multi-head Latent Attention（MLA）把 KV cache 压缩到 latent； 用 AVX2 int8/int4 kernel、GPU、LRU、page cache 和磁盘组成多级驻留； 用 MTP、n-gram 和 grammar 做 speculative decoding，再用一次 batch forward 验证； 最后用一个 OpenAI-compatible HTTP server 对外提供服务。
代码仓库是 glm-5.2-in-c，模型来源是 GLM-5.2-FP8。
仓库结构如下，目录职责和后面的复现步骤一一对应：
glm-5.2-in-c/
├── engine/ # the C engine
│ ├── glm.c # the whole engine in one file
│ ├── st.h, tier.h # streaming loader and expert tiering
│ ├── olmoe.c # the 400 line stepping stone we build first
│ ├── backend_cuda.cu # optional CUDA kernels
│ └── Makefile # one command builds it
├── tools/ # offline Python
│ ├── convert_fp8_to_int4.py # the FP8 to int4 requantizer
│ └── make_glm_oracle.py # the tiny PyTorch reference we validate against
├── cli/ # entry point and server
│ └── openai_server.py # the OpenAI compatible API
├── docs/ # the longer theory write ups
├── results/ # every log and number quoted in this post
├── deploy/ # provisioning and run scripts
└── bench/ # the capture scripts behind the results
01为什么 744B 可以放进小内存
问题不是“买多少张 GPU”，而是不要假设 7440 亿参数必须同时驻留。
GLM 5.2 有 78 层，前 3 层是 dense，之后 75 层每层有 256 个路由专家。每个 token 在每层只取 top-8，因此大部分专家在当前 token 上都处于“睡眠”状态。
按这套实现的估算，始终参与计算的部分约 170 亿参数，按 4 bit 计算约 9.9GB，可以放在 RAM 中；其余约 7270 亿个路由专家约 362GB，留在磁盘上。每个 token 真正触碰的是每层的 8 个专家，总体约 11GB 的专家数据，而且路由有明显的热点，很多专家会反复被访问。
一个专家包含 gate、up、down 三个矩阵。若矩阵形状为 O x I，int4 每字节装两个值，再加上每行一个 scale，则单个专家大约是 19MB。19MB 既是一次冷读取的工作单元，也是后面缓存、预取和并行化都要围绕的粒度。
所以“744B 跑在 20GB”并不意味着把完整模型塞进 20GB，而是把常驻部分、当前工作集、KV 状态和睡眠专家的磁盘存储拆开管理。
02模型与硬件
同一个 glm 二进制，从笔记本到测试机
最低形态是：
The floor: a machine you already own
CPU any modern x86-64 (AVX2 helps) or Apple Silicon
RAM about 16 to 26 GB
disk an NVMe SSD with room for the int4 model
GPU none required
公开测试记录显示，Framework 13 笔记本约 0.37 token/s，桌面机约 0.1 到 0.3 token/s。这里的意义是说明 GPU 不是运行的必要条件；速度则取决于 CPU、磁盘、RAM 和缓存状态。
开发和测量使用的是另一台更大的机器：4 张 NVIDIA L40、AMD EPYC 7763、228GB RAM 和 3.2TB NVMe。它用于获得更快、更可重复的测量，也用于展示增加 RAM 或 VRAM 后瓶颈如何移动，并不是 20GB 模式的最低硬件要求。
Model name: AMD EPYC 7763 64-Core Processor
CPU(s): 124
Thread(s) per core: 1
Core(s) per socket: 62
Socket(s): 2
NUMA node(s): 2
L3 cache: 1.9 GiB
CPU AVX flags (note: AVX2 present, NO avx512/vnni)
avx avx2 fma sse4_1 sse4_2
total used free shared buff/cache available
Mem: 228Gi 136Gi 2.9Gi 51Mi 90Gi 91Gi
Swap: 0B 0B 0B
CPU 有 AVX2 和 FMA，但没有 AVX-512 或 VNNI。后面的整数 kernel 因此以 AVX2 为主。
NVIDIA-SMI 570.195.03 Driver Version: 570.195.03 CUDA Version: 12.8
GPU Name Memory-Usage GPU-Util
0 NVIDIA L40 44543MiB / 49140MiB 0%
1 NVIDIA L40 44543MiB / 49140MiB 0%
2 NVIDIA L40 44545MiB / 49140MiB 0%
3 NVIDIA L40 44543MiB / 49140MiB 0%
软件依赖只有 C 编译器、可选的 CUDA 编译器和离线转换用的 Python：
gcc (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0
nvcc: Cuda compilation tools, release 12.8, V12.8.93
Python 3.12.3
模型配置中的关键字段如下：
{
"architectures": ["GlmMoeDsaForCausalLM"],
"model_type": "glm_moe_dsa",
"hidden_size": 6144,
"num_hidden_layers": 78,
"first_k_dense_replace": 3,
"num_attention_heads": 64,
"num_key_value_heads": 64,
"n_routed_experts": 256,
"num_experts_per_tok": 8,
"n_shared_experts": 1,
"moe_intermediate_size": 2048,
"q_lora_rank": 2048,
"kv_lora_rank": 512,
"qk_nope_head_dim": 192,
"qk_rope_head_dim": 64,
"v_head_dim": 256,
"index_head_dim": 128,
"index_n_heads": 32,
"index_topk": 2048,
"num_nextn_predict_layers": 1,
"vocab_size": 154880,
"scoring_func": "sigmoid",
"topk_method": "noaux_tc",
"routed_scaling_factor": 2.5,
"rope_parameters": { "rope_theta": 8000000 },
"quantization_config": {
"quant_method": "fp8", "fmt": "e4m3", "weight_block_size": [128, 128]
}
}
状态工具输出：
model /nvme/glm52_i4
arch hidden 6144 · 78 layer · 256 expert/layer · top-8
shards 144 files · 384 GB on disk
RAM 239 GB total · 232.6 GB available
disk 3029 GB free
engine ready
这里的 glm52_i4 是把供应商 FP8 checkpoint 转换后的目录。144 个分片在磁盘上约 384GB，但运行时只需要让 dense 部分和当前专家工作集进入内存。
03一个 C 文件组成的推理引擎
引擎的核心是约 3900 行的 glm.c，配合少量 header-only helper，不使用 BLAS 或推理框架。
GPU 构建：
nvcc -O3 -std=c++17 -arch=sm_89 -c backend_cuda.cu -o backend_cuda.o
gcc -O3 -march=native -fopenmp -DGLM_CUDA glm.c backend_cuda.o -o glm -lm -fopenmp \
-L/usr/local/cuda/lib64 -lcudart -lstdc++
CPU 构建：
gcc -O3 -march=native -fopenmp -Wall -Wextra glm.c -o glm -lm -fopenmp
编译产物只有约 377KB：
-rwxrwxr-x 1 ubuntu ubuntu 376648 Jul 14 06:37 glm
测试输出：
----- test_json ----- json tests: ok
----- test_st ----- safetensors primitive tests: ok
----- test_tier ----- tier tests: ok
----- test_grammar ----- test_grammar: ok
----- test_decode_batch ----- decode batch helper tests: ok
----- test_idot ----- idot kernel exactness (avx2): ok
----- test_i4_acc512 ----- test_i4_acc512: skipped (no AVX-512 on this build)
test_i4_acc512 跳过是因为测试 CPU 没有 AVX-512；idot kernel exactness (avx2): ok 则说明 AVX2 整数点积和普通 C reference 在测试范围内逐位一致。
另一个容易被忽略的点是 OpenMP。专家矩阵乘法很小且连续，默认的 passive wait 会让线程反复睡眠、唤醒，唤醒成本超过计算本身。引擎启动时会调整为 active spin，并重新执行一次：
[OMP] hot-thread tuning: re-exec once (GLM_NO_OMP_TUNE=1 to skip)
测试记录中，这个调整把 matmul 时间从 66.9 秒降到了 20.9 秒，输出没有变化。
04手写的基础数学
没有 BLAS，归一化、激活和 RoPE 都是 C 函数。
/* RMS norm: divide by the root-mean-square of the vector, then scale by w. */
static void rmsnorm(float *out, const float *x, const float *w, int D, float eps) {
double ms = 0; for (int i = 0; i < D; i++) ms += (double)x[i] * x[i];
float r = 1.f / sqrtf((float)(ms / D) + eps);
for (int i = 0; i < D; i++) out[i] = x[i] * r * w[i];
}
/* Softmax, shifted by the max for numerical stability. */
static void softmax(float *x, int n) {
float m = -1e30f; for (int i = 0; i < n; i++) if (x[i] > m) m = x[i];
float s = 0; for (int i = 0; i < n; i++) { x[i] = expf(x[i] - m); s += x[i]; }
for (int i = 0; i < n; i++) x[i] /= s;
}
/* SiLU, the activation inside every SwiGLU expert. */
static inline float siluf(float x) { return x / (1.f + expf(-x)); }
RoPE 通过旋转成对的 query/key 元素，让 attention 感知位置：
/* Rotary position embedding, interleaved. Each pair (2j, 2j+1) is rotated by an
* angle that grows with the position, so attention becomes position-aware. */
static void rope_interleave(float *v, int pos, const Cfg *c) {
int half = c->qk_rope / 2; float in[256]; memcpy(in, v, c->qk_rope * sizeof(float));
for (int j = 0; j < half; j++) {
float inv = powf(c->theta, -2.0f * j / c->qk_rope);
float ang = pos * inv, cs = cosf(ang), sn = sinf(ang);
float a = in[2*j], b = in[2*j + 1];
v[j] = a * cs - b * sn;
v[half + j] = b * cs + a * sn;
}
}
一个 transformer layer 的顺序是：输入归一化、attention、残差、第二次归一化、MoE 或 dense feed-forward、再次残差。
/* One layer: in_norm -> attention -> residual -> post_norm -> MoE/dense -> residual. */
for (int s = 0; s < S; s++) rmsnorm(nrm + (int64_t)s*D, x + (int64_t)s*D, l->in_ln, D, c->eps);
attention_rows(m, l, li, nrm, S, pos_base, kvs, positions, tmp);
for (int64_t j = 0; j < (int64_t)S*D; j++) x[j] += tmp[j];
for (int s = 0; s < S; s++) rmsnorm(nrm + (int64_t)s*D, x + (int64_t)s*D, l->post_ln, D, c->eps);
if (l->sparse) moe(m, l, li, nrm, S, tmp);
else dense_mlp(l, nrm, S, D, c->dense_inter, tmp);
for (int64_t j = 0; j < (int64_t)S*D; j++) x[j] += tmp[j];
05先用 400 行验证流式思路
完整引擎之前，先实现了一个更小的 olmoe.c。目标很单纯：常驻 dense 权重，把专家从磁盘读入小缓存，并与 PyTorch reference 逐 token 对齐。
专家缓存槽位保存量化后的 gate、up、down，以及每行的 scale：
/* One expert's weights, held quantized. Each matrix [out,in] is int8 per row
* plus one float scale per row. This is what takes the RAM cost from
* 4 bytes/param (f32) down to 1 byte/param. We dequantize on use in the matmul. */
typedef struct { int eid; int8_t *g, *u, *d; float *gs, *us, *ds; uint64_t used; } Slot;
typedef struct { Slot *slots; int n, cap; } LCache;
量化采用逐行 absmax：
/* Quantize a weight [O,I] to int8 q[O,I] plus a per-row scale, symmetric. */
static void quantize_rows(const float *w, int8_t *q, float *scale, int O, int I, int bits) {
int qmax = (1 << (bits - 1)) - 1;
#pragma omp parallel for schedule(static)
for (int o = 0; o < O; o++) {
const float *wr = w + (int64_t)o * I;
float amax = 0.f;
for (int i = 0; i < I; i++) { float a = fabsf(wr[i]); if (a > amax) amax = a; }
float s = amax / qmax; if (s < 1e-8f) s = 1e-8f;
scale[o] = s;
int8_t *qr = q + (int64_t)o * I;
for (int i = 0; i < I; i++) {
int v = (int)lrintf(wr[i] / s);
if (v > qmax) v = qmax;
if (v < -qmax - 1) v = -qmax - 1;
qr[i] = (int8_t)v;
}
}
}
缓存命中就提升 used；未命中则淘汰 LRU 槽位，读取三个矩阵：
/* Return the quantized weights of one expert, from cache or from disk. */
static void expert_get(Model *m, int layer, int eid, Slot **out) {
LCache *lc = &m->cache[layer];
for (int i = 0; i < lc->n; i++) if (lc->slots[i].eid == eid) {
m->hits++; lc->slots[i].used = ++m->clock; *out = &lc->slots[i]; return;
}
m->miss++;
Slot *s;
if (lc->n < lc->cap) {
s = &lc->slots[lc->n++];
s->g = malloc(ng); s->u = malloc(ng); s->d = malloc(nd);
s->gs = falloc(c->inter); s->us = falloc(c->inter); s->ds = falloc(c->hidden);
} else {
int lru = 0;
for (int i = 1; i < lc->n; i++) if (lc->slots[i].used < lc->slots[lru].used) lru = i;
s = &lc->slots[lru];
}
load_expert_w(m, gate_name, s->g, s->gs, c->inter, c->hidden, tmp);
load_expert_w(m, up_name, s->u, s->us, c->inter, c->hidden, tmp);
load_expert_w(m, down_name, s->d, s->ds, c->hidden, c->inter, tmp);
s->eid = eid; s->used = ++m->clock;
*out = s;
}
MoE 的形状就是路由、取 top-k、取缓存或磁盘中的专家、运行 SwiGLU，再将加权结果相加：
/* The MoE step for tokens x[S,hidden] -> out[S,hidden]. */
static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out) {
Cfg *c = &m->c; int D = c->hidden, E = c->n_experts, K = c->topk, I = c->inter;
float *logits = falloc((int64_t)S * E);
matmul(logits, x, l->gate, S, D, E);
memset(out, 0, (int64_t)S * D * sizeof(float));
for (int s = 0; s < S; s++) {
float *pr = logits + (int64_t)s * E;
softmax_row(pr, E);
int idx[64]; float val[64];
for (int kk = 0; kk < K; kk++) {
int best = -1; float bv = -1e30f;
for (int e = 0; e < E; e++) {
int taken = 0;
for (int j = 0; j < kk; j++) if (idx[j] == e) { taken = 1; break; }
if (!taken && pr[e] > bv) { bv = pr[e]; best = e; }
}
idx[kk] = best; val[kk] = bv;
}
const float *xs = x + (int64_t)s * D;
for (int kk = 0; kk < K; kk++) {
Slot *e; expert_get(m, layer, idx[kk], &e);
matmul_q(g, xs, e->g, e->gs, D, I);
matmul_q(u, xs, e->u, e->us, D, I);
for (int i = 0; i < I; i++) { float gv = g[i]; g[i] = (gv / (1.f + expf(-gv))) * u[i]; }
matmul_q(hh, g, e->d, e->ds, I, D);
float w = val[kk];
float *os = out + (int64_t)s * D;
for (int d = 0; d < D; d++) os[d] += w * hh[d];
}
}
}
小模型的结果：
== Streaming C engine, cache = 16 experts/layer, experts @ 8-bit ==
resident weights loaded in ... | RSS after load: ... GB
Reference: 207 187 119 103 103 103 103 103 119 34 ...
C engine : 207 187 119 103 103 103 103 103 119 34 ...
Matching tokens: 20/20
Expert cache hit rate: ... (hit=... miss=...)
20/20 token 对齐不是完整模型质量 benchmark，而是验证流式加载、专家缓存和解码路径的实现一致性。
06从 FP8 到 int4
供应商 checkpoint 使用 FP8 e4m3，按这套实现的估算约 756GB。CPU 不能直接高效地对这种 FP8 格式做完整计算，因此先离线解码为 float32，再转换成 C 引擎需要的 int4 容器。
每行先找绝对值最大项，得到 scale，再将值舍入到 [-8, 7]：
两个 int4 nibble 装进一个字节。为避免有符号存储，先加 8 映射到 0..15，第一个值放低四位，第二个值左移 4 位放高四位。
Python 转换器：
def quant_int4(w, bits):
O, I = w.shape
qmax = (1 << (bits - 1)) - 1
amax = np.abs(w).max(axis=1, keepdims=True)
s = np.maximum(amax / qmax, 1e-8)
q = np.clip(np.rint(w / s), -8, qmax).astype(np.int32)
rb = (I + 1) // 2
out = np.zeros((O, rb), np.uint8)
v0 = (q[:, 0::2] + 8).astype(np.uint8)
out[:, :v0.shape[1]] = v0
if I > 1:
v1 = (q[:, 1::2] + 8).astype(np.uint8)
out[:, :v1.shape[1]] |= (v1 << 4)
return out.reshape(-1), s[:, 0].astype(np.float32)
C 端的 packer：
/* Pack w[O,I] f32 -> int4 (2 per byte) + a per-row scale. */
static void pack_int4(const float *w, uint8_t *q4, float *scale, int O, int I, int bits) {
int qmax = (1 << (bits - 1)) - 1, rb = (I + 1) / 2;
#pragma omp parallel for schedule(static)
for (int o = 0; o < O; o++) {
const float *wr = w + (int64_t)o * I; float amax = 0;
for (int i = 0; i < I; i++) { float a = fabsf(wr[i]); if (a > amax) amax = a; }
float s = amax / qmax; if (s < 1e-8f) s = 1e-8f; scale[o] = s;
uint8_t *qr = q4 + (int64_t)o * rb;
for (int i = 0; i < I; i += 2) {
int v0 = (int)lrintf(wr[i] / s);
if (v0 > qmax) v0 = qmax; if (v0 < -8) v0 = -8;
int v1 = 0;
if (i + 1 < I) {
v1 = (int)lrintf(wr[i+1] / s);
if (v1 > qmax) v1 = qmax; if (v1 < -8) v1 = -8;
}
qr[i >> 1] = (uint8_t)((v0 + 8) | ((v1 + 8) << 4));
}
}
}
FP8 block scale 解码：
def dequant(f, name):
import torch
sl = f.get_slice(name); dt = sl.get_dtype()
if dt in ("F8_E4M3", "float8_e4m3fn"):
w = f.get_tensor(name).to(torch.float32)
sc = f.get_tensor(name + "_scale_inv").to(torch.float32)
O, I = w.shape
sc = sc.repeat_interleave(128, 0).repeat_interleave(128, 1)[:O, :I]
return (w * sc).numpy()
return f.get_tensor(name).to(torch.float32).numpy()
不同 tensor 不使用同一精度策略：
def classify(name, n_layers, keep_mtp=False, keep_idx=False):
if name.endswith("_scale_inv"): return "consumed"
if name.endswith("e_score_correction_bias"): return "f32"
if name.endswith("mlp.gate.weight"): return "f32"
if name.endswith("norm.weight") or name == "model.norm.weight": return "f32"
if name in ("model.embed_tokens.weight", "lm_head.weight"): return "io"
if ".mlp.experts." in name and name.endswith(".weight"): return "x"
if name.endswith(".weight"): return "q"
return "f32"
路由专家占模型绝大多数，转成 int4；embedding、输出头和 MTP 等边界部分保留更高精度；router、norm 和 bias 保持 float32。转换器按 shard 工作，避免同时保留完整 FP8 checkpoint 和 int4 输出：
python tools/convert_fp8_to_int4.py --repo zai-org/GLM-5.2-FP8 --outdir /nvme/glm52_i4 \
--ebits 4 --io-bits 8
一个转换后的专家 gate projection 被解码为：
name: model.layers.0.mlp.gate_proj.weight
dtype: U8 shape(packed bytes): [37748736] nbytes: 37748736
rows (from .qs): 12288 -> implied cols = nbytes*2/rows = 6144 => int4 (2 vals/byte) CONFIRMED
first 8 packed bytes: 87 78 77 89 b5 b8 77 77
decoded int4 values (nibble-8): [-1, 0, -1, 1, -3, 0, -1, -1] (low nibbles)
FP8 block dequant 自测：
[selftest fp8 block-dequant] mean relative error = 0.0226 (OK)
模型统计图显示：约 362GB 是 int4 routed experts，约 12GB 是 int8 的输入输出边界和 MTP，router、bias、norm 不到 1GB。
07int4 的代价
int4 并不是免费压缩。测试在一个小型、同架构模型上用 PyTorch oracle 做 teacher forcing，比较不同 bit-width 下有多少位置完全一致。
== ./glm 64 16 16 == experts@16-bit dense@16-bit -> 32/32 positions
== ./glm 64 8 8 == experts@8-bit dense@8-bit -> 30/32 positions
== ./glm 64 4 8 == experts@4-bit dense@8-bit -> 28/32 positions
== ./glm 64 4 4 == experts@4-bit dense@4-bit -> 9/32 positions
== ./glm 64 8 4 == experts@8-bit dense@4-bit -> 6/32 positions
== ./glm 64 2 2 == experts@2-bit dense@2-bit -> 1/32 positions
这个结果告诉我们：routed experts 对 int4 更耐受，而 dense path、attention 和 shared experts 对低比特更敏感。专家 int4、dense int8 得到 28/32；专家 int8、dense int4 只有 6/32。
这里的 9/32 是小模型的实现敏感度实验，不是完整 GLM 5.2 的质量 benchmark，也不能单独推出“完整模型质量下降了多少”。
Python converter 和 C runtime 的打包一致性也进行了对照：
A: oracle full weights, C engine runtime-quantizes to int4 (pack_int4):
PREFILL (teacher-forcing) C vs oracle: 9/32 positions
B: python-converter int4 container, C engine reads it verbatim:
PREFILL (teacher-forcing) C vs oracle: 9/32 positions
(identical X/32 => python np.rint quantizer == C lrintf pack_int4)
两条路径得到相同结果，说明 np.rint 和 lrintf 产生了相同的量化字节，离线 Python 转换和 C 运行时读取可以对齐。
08按需读取权重
int4 容器有了，下一步是让权重读取不把 384GB 全部变成进程常驻页。
GLM 5.2 的约 12 万个 tensor 分布在 144 个 shard 中。启动时先建立 hash index：
typedef struct {
char *name; int fd; int64_t off; int64_t nbytes; int dtype; int64_t numel;
} st_tensor;
static st_tensor *st_find(shards *S, const char *name) {
uint64_t h = st_hash(name) & (S->hcap - 1);
while (S->hidx[h] >= 0) {
st_tensor *t = &S->t[S->hidx[h]];
if (!strcmp(t->name, name)) return t;
h = (h + 1) & (S->hcap - 1);
}
return NULL;
}
运行时只读取指定 tensor 的字节，并在需要时提示内核丢弃文件页：
static int64_t st_read_f32(shards *S, const char *name, float *out, int drop) {
st_tensor *t = st_find(S, name);
if (!t) { fprintf(stderr, "missing tensor: %s\n", name); exit(1); }
void *raw = malloc(t->nbytes);
if (pread(t->fd, raw, t->nbytes, t->off) != t->nbytes) { perror("pread data"); exit(1); }
if (t->dtype == 2) memcpy(out, raw, t->nbytes);
else if (t->dtype == 0) {
uint16_t *p = raw;
for (int64_t i = 0; i < t->numel; i++) out[i] = bf16_to_f32(p[i]);
} else {
uint16_t *p = raw;
for (int64_t i = 0; i < t->numel; i++) out[i] = f16_to_f32(p[i]);
}
free(raw);
if (drop) posix_fadvise(t->fd, t->off, t->nbytes, POSIX_FADV_DONTNEED);
return t->numel;
}
实现还保留了一个 O_DIRECT 文件描述符，用于绕过 page cache。19MB 专家随机读的报告是：
buffered x8 threads: 64 reads x 19MB = 1.3 GB in 0.40s -> 3.22 GB/s (6.2 ms/block)
O_DIRECT x8 threads: 64 reads x 19MB = 1.3 GB in 0.31s -> 4.13 GB/s (4.8 ms/block)
这不是跨设备基准，而是该测试环境的冷读取结果。8 个专家跨 75 个 MoE 层大约需要 600 次读取，因此 tiering 的目标不是让磁盘无限快，而是尽量不触发冷读。
09MLA：把 KV cache 压到 576 个 float
GLM 5.2 使用 Multi-head Latent Attention。每个 token 不保存 64 个 head 的完整 K/V，而只保存一个 512 维 latent 和 64 维 rotary 部分。
完整 K/V 需要 64 * (256 + 256) = 32768 个 float；MLA 的缓存只需 512 + 64 = 576 个 float，约小 57 倍。
/* The MLA compressed KV cache: per token we keep only the normalized latent
* Lc [kv_lora] and the rotary key Rc [qk_rope]. */
typedef struct {
float **Lc, **Rc, **Ic;
int *kv_start, max_t;
int disk_nrec;
char disk_path[2048];
} KVState;
decode 时通过 weight absorption，把 key 的上投影吸收到 query 中，直接对 latent 打分，而不是对每个历史 token 重新重建 K/V。
/* Weight absorption for decode (small S). */
#pragma omp parallel for collapse(2) schedule(static)
for (int s = 0; s < S; s++) for (int h = 0; h < H; h++) {
const float *qp = Q + (int64_t)s*H*qh + (int64_t)h*qh;
const float *qr = qp + c->qk_nope;
float qabs[512]; memset(qabs, 0, kvl * sizeof(float));
for (int d = 0; d < c->qk_nope; d++)
qt_addrow(&l->kv_b, rbase + d, qp[d], qabs);
float *sc = sc_all + (int64_t)omp_get_thread_num() * sc_cap;
for (int jj = 0; jj < nt; jj++) {
const float *Lt = kv_row(ks->Lc[layer], st0 + jj, kvl);
const float *kr = kv_row(ks->Rc[layer], st0 + jj, c->qk_rope);
float a = 0; for (int i = 0; i < kvl; i++) a += qabs[i] * Lt[i];
for (int d = 0; d < c->qk_rope; d++) a += qr[d] * kr[d];
sc[jj] = a * c->attn_scale;
}
softmax(sc, nt);
float clat[512]; memset(clat, 0, kvl * sizeof(float));
for (int jj = 0; jj < nt; jj++) {
const float *Lt = kv_row(ks->Lc[layer], st0 + jj, kvl);
float a = sc[jj]; for (int i = 0; i < kvl; i++) clat[i] += a * Lt[i];
}
qt_matvec_rows(&l->kv_b, rbase + r0v, vh, clat, ctx + ((int64_t)s*H + h)*vh);
}
实现用显式 K/V 重建和 absorption 两条路径比较：
== ABSORB=1 (absorbed low-rank MLA) ==
GLM C engine : 207 187 119 103 103 103 103 103 119 34 ...
== ABSORB=0 (explicit K/V reconstruction) ==
GLM C engine : 207 187 119 103 103 103 103 103 119 34 ...
缓存还可以持久化：
file: /nvme/glm52_i4/.glm_kv
-rw-rw-r-- 1 ubuntu ubuntu 45108756 ... .glm_kv
跨进程恢复的日志：
===== RUN 1: fresh prompt, generate, persist KV =====
One Two Three Four Five Six Seven Eight Nine Ten Eleven Twelve Thirteen Fourteen Fifteen Sixteen Se
===== RUN 2: NEW process reopens .glm_kv, continues, no re-prefill =====
[KV] resumed conversation from disk: 57 tokens in 0.0s (no re-prefill)
Seventeen Eighteen Nineteen Twenty
10稀疏 attention indexer
长上下文时，模型使用 sparse attention indexer 先对历史位置打分，只保留 index_topk=2048；上下文短于 2048 时，保留的就是全部历史位置，因此该路径应与 dense attention 等价。
/* For each query, score every past key with a ReLU'd per-head dot product,
* weight and sum across heads, then keep the top index_topk positions. */
float wsc = 1.f / sqrtf((float)nh), rs = 1.f / sqrtf((float)hd);
float *isc = falloc(nk);
for (int t = 0; t < nk; t++) {
const float *kt = kv_row(ks->Ic[layer], t, hd);
float a = 0;
for (int h = 0; h < nh; h++) {
const float *qhp = qi + (int64_t)h*hd;
float d0 = 0; for (int i = 0; i < hd; i++) d0 += qhp[i] * kt[i];
d0 *= rs; if (d0 > 0) a += w32[h] * d0;
}
isc[t] = a * wsc;
}
float *tmp = falloc(nk); memcpy(tmp, isc, nk * sizeof(float));
qsort(tmp, nk, sizeof(float), cmp_fdesc);
float thr = tmp[keep - 1];
int *dst = m->dsa_sel + (int64_t)s * dtopk, nd = 0;
for (int t = 0; t < nk && nd < keep; t++) if (isc[t] > thr) dst[nd++] = t;
开关对照：
=== DSA_FORCE=1 vs default === generated text: identical (only the timing banner differs)
=== DSA=0 vs default === generated text: identical (only the timing banner differs)
11MoE 路由与专家执行
router 对 256 个专家打分并选 top-8。这里有一个实现细节：sigmoid(logit) + correction bias 用于选择，但最终加权使用的是未加 bias 的 sigmoid 值。
/* sigmoid(logit)+bias picks top-K, weight is plain sigmoid(logit). */
for (int s = 0; s < S; s++) {
const float *xs = x + (int64_t)s * D;
matmul(logit, xs, l->router, 1, D, E);
for (int e = 0; e < E; e++) {
logit[e] = sigmoidf(logit[e]);
choice[e] = logit[e] + l->router_bias[e];
}
int *idx = idxs + (int64_t)s*K;
float *w = ws + (int64_t)s*K;
for (int kk = 0; kk < K; kk++) {
int best = -1; float bv = -1e30f;
for (int e = 0; e < E; e++) {
int tk = 0;
for (int j = 0; j < kk; j++) if (idx[j] == e) { tk = 1; break; }
if (!tk && choice[e] > bv) { bv = choice[e]; best = e; }
}
idx[kk] = best; w[kk] = logit[best];
}
for (int kk = 0; kk < K; kk++) w[kk] *= c->routed_scale;
}
批量 prefill 或 speculative verification 时，会先对 batch 中出现的专家做 unique union，每个专家只读一次：
int *uniq = malloc((size_t)E * sizeof(int)); int nu = 0;
unsigned char seen[E]; memset(seen, 0, (size_t)E);
for (int s = 0; s < S; s++) for (int kk = 0; kk < keff[s]; kk++) {
int e = idxs[(int64_t)s*K + kk];
if (!seen[e]) { seen[e] = 1; uniq[nu++] = e; }
}
每个专家是 SwiGLU：
expert_gate_up(gg, uu, xg, &e->g, &e->u, nr);
for (int64_t z = 0; z < (int64_t)nr*I; z++) gg[z] = siluf(gg[z]) * uu[z];
matmul_qt(hh, gg, &e->d, nr);
for (int r = 0; r < nr; r++) {
float *os = out + (int64_t)rows[r]*D, wgt = rw[r];
float *hr = hh + (int64_t)r*D;
for (int d = 0; d < D; d++) os[d] += wgt * hr[d];
}
路由统计显示专家使用有明显偏斜，测试记录中最热的专家在短时间内被触发 776 次。热点分布是缓存有效的前提。
12AVX2 整数 kernel
矩阵乘法是主要 CPU 成本。引擎同时保留 float、int8、int4 和 int2 路径；整数路径会先将 activation 量化为 int8，再把权重和 activation 的整数点积乘以 scale。
static inline float qrow_i8(const float *x, int8_t *q, int I) {
float amax = 0;
for (int i = 0; i < I; i++) { float a = fabsf(x[i]); if (a > amax) amax = a; }
float s = amax / 127.f; if (s < 1e-12f) s = 1e-12f;
float inv = 1.f / s;
for (int i = 0; i < I; i++) q[i] = (int8_t)lrintf(x[i] * inv);
return s;
}
AVX2 没有直接的 signed-byte 乘加指令，实现使用符号折叠，让权重绝对值变成 unsigned、把权重符号并入 activation：
static inline int32_t dot_i8i8(const int8_t *w, const int8_t *x, int I) {
int32_t sum = 0; int i = 0;
#if defined(__AVX2__)
__m256i acc = _mm256_setzero_si256();
const __m256i ones = _mm256_set1_epi16(1);
for (; i + 32 <= I; i += 32) {
__m256i wv = _mm256_loadu_si256((const __m256i*)(w + i));
__m256i xv = _mm256_loadu_si256((const __m256i*)(x + i));
__m256i p = _mm256_maddubs_epi16(_mm256_sign_epi8(wv, wv),
_mm256_sign_epi8(xv, wv));
acc = _mm256_add_epi32(acc, _mm256_madd_epi16(p, ones));
}
sum = hsum256_i32(acc);
#endif
for (; i < I; i++) sum += (int32_t)w[i] * x[i];
return sum;
}
统一 dispatcher 先尝试 GPU，再尝试整数 kernel，最后回退到 float：
static void matmul_qt(float *y, const float *x, QT *w, int S) {
#ifdef GLM_CUDA
if (g_cuda_enabled && w->cuda_eligible && !w->cuda_failed && !omp_in_parallel()) {
if (glm_cuda_matmul(&w->cuda, y, x, weights, w->s, w->fmt,
S, w->I, w->O, w->cuda_device)) return;
}
#endif
if (w->fmt == 0) { matmul(y, x, w->qf, S, w->I, w->O); return; }
if (g_idot && (w->fmt == 1 || (w->fmt == 2 && S >= g_i4s))) {
int I = w->I; int8_t *xq; float *sx;
quant_scratch((size_t)S*I, (size_t)S, &xq, &sx);
for (int s = 0; s < S; s++)
sx[s] = qrow_i8(x + (int64_t)s*I, xq + (int64_t)s*I, I);
if (w->fmt == 1) matmul_q_idot(y, xq, sx, w->q8, w->s, S, I, w->O);
else matmul_i4_idot(y, xq, sx, w->q4, w->s, S, I, w->O);
return;
}
if (w->fmt == 1) matmul_q(y, x, w->q8, w->s, S, w->I, w->O);
else if (w->fmt == 3) matmul_i2(y, x, w->q4, w->s, S, w->I, w->O);
else matmul_i4(y, x, w->q4, w->s, S, w->I, w->O);
}
测试：
----- test_idot ----- idot kernel exactness (avx2): ok
但 int8 activation path 会引入约 0.3% RMS 的量化噪声，偶尔让 token 与 float path 不同：
=== IDOT=1 (int8 activation) === A mutex (mutual exclusion) is
=== IDOT=0 (f32 reference) === A mutex (mutual exclusion object)
这不是“一个输出正确、另一个错误”，而是速度和逐 token 可复现性之间的取舍。需要严格复现时可以关闭整数 activation path。
13五级专家驻留
专家的层级是：VRAM、pinned RAM、按层 LRU cache、操作系统 page cache、磁盘。
替换策略对近期热度设置了 25% 加 4 的 margin：
static int tier_pick_swap(const uint32_t *heat, int nexpert, const int *pinned, int npin,
int *slot, int *eid, long *gain) {
if (!heat || !pinned || npin < 1 || nexpert < 1) return 0;
int cold = 0;
for (int z = 1; z < npin; z++)
if (heat[pinned[z]] < heat[pinned[cold]]) cold = z;
int hot = -1; uint32_t fh = 0;
for (int e = 0; e < nexpert; e++) {
int resident = 0;
for (int z = 0; z < npin; z++) if (pinned[z] == e) { resident = 1; break; }
if (!resident && heat[e] > fh) { fh = heat[e]; hot = e; }
}
if (hot < 0) return 0;
uint32_t fc = heat[pinned[cold]];
if (fh <= fc + (fc >> 2) + 4) return 0;
*slot = cold; *eid = hot; *gain = (long)fh - (long)fc;
return 1;
}
频率优先，近期性只在接近时打破平局：
static uint64_t tier_lfru_score(uint32_t heat, uint32_t last, uint32_t clock) {
uint32_t age = clock - last, recent = age < 255 ? 255 - age : 0;
return ((uint64_t)heat << 8) | recent;
}
预量化容器中一个专家的三个矩阵连续存储时，可以用一次约 19MB 的 pread 读入 slab，再让三个 QT 指向同一块内存：
int contig = tw[ord[0]].off + tw[ord[0]].nbytes == tw[ord[1]].off
&& tw[ord[1]].off + tw[ord[1]].nbytes == tw[ord[2]].off;
if (contig) {
if (pread(tw[ord[0]].fd, s->slab, wtot, off0) != wtot) {
perror("pread expert"); exit(1);
}
pos[ord[0]] = 0;
pos[ord[1]] = tw[ord[0]].nbytes;
pos[ord[2]] = tw[ord[0]].nbytes + tw[ord[1]].nbytes;
}
QT *qt[3] = {&s->g, &s->u, &s->d};
for (int k = 0; k < 3; k++) {
qt[k]->q8 = (int8_t*)(s->slab + pos[k]);
qt[k]->q4 = s->slab + pos[k];
qt[k]->s = fp[k];
}
内存预算不是“尽量多塞缓存”，而是显式扣除 dense、KV、工作集、page cache reserve 和 activation：
double ws_b = 64.0 * (double)eb;
double kv_b = kv_pool_bytes(m, max_ctx);
double kvb_b = (double)max_ctx * c->n_heads * (c->qk_nope + c->v_head) * 4.0;
double pc_b = 2.5e9;
double slack = 1.2e9 + pc_b + ws_b + kv_b + kvb_b;
double avail = ram_gb * 1e9 - (double)m->resident_bytes - slack;
int capmax = (avail > 0 && nsp > 0)
? (int)(avail / ((double)nsp * eb)) : 0;
if (capmax < 1) capmax = 1;
20GB 和 200GB 预算的日志：
[RAM_GB=20.0] resident 12.1 GB + reserve 6.1 GB (ws 1.2, KV 1x4096 0.7, kvb 0.5),
experts 18.9 MB x 77 layers -> cap lowered 8->1 (projected peak 19.7 GB)
[RAM_GB=200.0] cap raised 8->45: budget allows it (projected peak 199.3 GB)
四张 L40 的 placement 计划：
policy quality · quality-preserving yes
model 144 shards · 383.7 GB
disk 4.9 GB cold experts · 3029.2 GB free
RAM 200.0 GB budget · 10.9 GB dense · 6.1 GB runtime · 183.0 GB warm experts · cap 125/layer
VRAM 185.0 GB hot tier · ~9780 experts · 4x NVIDIA L40
limit disk expert misses
启动前的 doctor：
[ ok] model.path model directory is readable
[ ok] model.config config.json is valid
[ ok] model.tokenizer tokenizer.json found
[ ok] engine.binary engine executable is ready
[ ok] accelerator.cuda CUDA engine and devices are available
[ ok] model.shards safetensors headers are valid
[ ok] memory.ram RAM budget is viable
[warn] placement.plan cold expert misses may reach disk; normal decode speed depends on hit rate
result warning
GPU 热专家分布：
[CUDA] hot expert tier: 9780/15860 experts, VRAM 184.99 GB (total budget 185.0 GB)
[CUDA] device 0: 2445 experts, 46.25 GB
[CUDA] device 1: 2445 experts, 46.25 GB
[CUDA] device 2: 2445 experts, 46.25 GB
[CUDA] device 3: 2445 experts, 46.25 GB
[RAM_GB=200.0] cap raised 8->45: budget allows it (projected peak 199.3 GB)
OpenAI-compatible API listening on http://0.0.0.0:8000/v1
VRAM residency (per GPU): 44543 / 44543 / 44545 / 44543 MiB of 49140 MiB
RAM: 137 GB used
page cache holds streamed experts: Cached 93066072 kB
系统还把历史访问统计写到 <model>/.glm_usage，启动时据此预先 pin 热专家：
double conf = (double)hist / 200000.0;
if (conf > 1) conf = 1;
double pin_gb = expert_avail(m, ram_env, ebits, est_ctx) * 0.5 * conf / 1e9;
if (pin_gb >= 0.5) pin_load(m, g_usage_path, pin_gb);
RAM 增加后，测试记录显示命中率上升，但 CPU-only 吞吐并不跟随上升；瓶颈会从磁盘转移到 CPU matmul。
增加 VRAM 则能继续降低专家 matmul 时间：
14处理磁盘延迟
page cache 是最直接的缓存层：
COLD (drop page cache): 24 tokens in 76.20s (0.31 tok/s), expert-disk 43.71s
WARM (immediate repeat): 24 tokens in 50.15s (0.48 tok/s), expert-disk 21.92s
DROP=1 (fadvise DONTNEED evicts): 24 tokens in 184.81s (0.13 tok/s), expert-disk 149.54s
warm 运行明显快于 cold，而强制 DONTNEED 会让每次专家访问回到磁盘。
引擎也尝试了预取、I/O worker 和跨层 router lookahead：
但在这套测试使用的虚拟化文件系统和 NVMe 上，并行读取反而变慢：
PIPE=0 (serial load then matmul): 32 tokens in 83.27s (0.38 tok/s)
PIPE=1 (overlap disk load with matmul): 32 tokens in 197.65s (0.16 tok/s)
lookahead 虽然把命中率从 58% 提到 62%，却因预取造成驱逐压力而降低吞吐。最终不同 I/O 模式的报告是：
buffered: 0.17 tok/s drop: 0.20 tok/s direct: 0.23 tok/s mmap: 0.19 tok/s
这组结果的边界很重要：预取是否有用取决于设备是否能承受并发读，不能把某一台机器上的 pipeline 结果外推到所有 NVMe。
15三种 speculative decoding
逐 token 解码每次都要付出一轮权重访问。speculative decoding 先生成多个草稿 token，再用一次 batched forward 验证，接受与 greedy 结果匹配的最长前缀。
/* Lossless self-speculation: draft, verify in one batched forward. */
while (emitted < n_new && !done) {
int next = pick_tok(logit, V, carry_ban); free(logit);
if ((eos >= 0 && next == eos) || is_stop(next)) break;
emit(next, ud); all[kv] = next; emitted++;
int g = 0, gsrc = 0;
if (g_gr_on) { g = grammar_draft(draft, g_gr_max); if (g > 0) gsrc = 1; }
if (!g && g_draft > 0) {
if (m->has_mtp) g = mtp_draft(m, next, kv, g_draft, draft);
else g = ngram_draft(all, kv + 1, g_draft, draft);
}
int S = 1 + g; int batch[64]; batch[0] = next;
memcpy(batch + 1, draft, g * sizeof(int));
float *lo = step_all(m, batch, S, kv);
int k = 0;
while (k < g && emitted < n_new) {
int accept = (argmax_v(lo + (int64_t)k*V, V) == draft[k]);
if (!accept) break;
emit(draft[k], ud); all[kv + 1 + k] = draft[k]; emitted++; k++;
}
kv += 1 + k;
logit = falloc(V);
memcpy(logit, lo + (int64_t)k*V, V * sizeof(float));
free(lo);
}
草稿来源有三种：模型自带的 multi-token prediction head、n-gram，以及 grammar。MTP 必须保留 int8，否则接受率接近于零。
MTP ON (draft=3): 16 tokens in 60.03s (0.27 tok/s) | speculation 2.67 tokens/forward | acceptance 50%
MTP OFF (draft=0): 16 tokens in 64.84s (0.25 tok/s) | speculation 1.07 tokens/forward | acceptance 0%
Grammar 示例：
[GRAMMAR] person.gbnf: 10 rules, forced span capped at 24 tokens/forward
{"name": "Maxwell Vance", "age": 34, "city": "Seattle"}
grammar: 50% acceptance (3/6 forced drafts)
采样参数也需要收紧：
temp=0 greedy: "Euphoric" (stable across repeats)
temp=0.7 top_p=0.9: "Elatated" / "Euphoric" / "Elysian"
temp=1.0 top_p=0.95: "Elate" / "Elated" / "Euphoric"
temp=1.3 top_p=0.98: "Ebullient" / "Ecstatic" ... then degrades into gibberish
这里可以归因于 int4 后分布尾部的量化噪声：温度和 nucleus 过高时，更容易采到噪声。
16OpenAI-compatible 服务
引擎外层是只使用 Python 标准库的 HTTP server。server 负责调度、队列、SSE streaming 和 backpressure；常驻的 C engine 通过管道上的字节协议接收任务。
请求以 SUBMIT 行开始，随后读取可能包含换行的 prompt payload。C 端会验证长度、slot、token 数、temperature 和 top_p：
static inline int submit_parse(const char *line, Submit *s) {
char tail;
if (!line || !s ||
sscanf(line, "SUBMIT %llu %d %llu %d %f %f %c", &s->id, &s->slot,
&s->bytes, &s->max_tokens, &s->temperature, &s->top_p, &tail) != 6)
return 0;
return s->id > 0 && s->bytes <= (16u << 20) && s->slot >= 0 && s->max_tokens >= 1 &&
isfinite(s->temperature) && isfinite(s->top_p) &&
s->temperature >= 0 && s->temperature <= 2 &&
s->top_p > 0 && s->top_p <= 1;
}
每个 slot 只有一个可变 KV context，因此 server 使用固定容量和有界队列；满载时返回 HTTP 429。
printf("DONE %llu STAT %d %.2f %.1f %.2f %d %d\n", r->id, r->emitted,
r->emitted / dt, (dh + dm) > 0 ? 100.0 * dh / (dh + dm) : 0.0, rss_gb(),
r->prompt_tokens, r->length_limited);
prefill 和 API 日志：
[prefill] layer 1/78 · 13 token
[prefill] layer 5/78 · 13 token
...
[prefill] layer 77/78 · 13 token
[prefill] layer 78/78 · 13 token
[api] "POST /v1/chat/completions HTTP/1.1" 200
一次请求的流式统计：
{
"ttft_s_incl_prefill": 5.6807,
"content_deltas": 200,
"last_content_at_s": 104.299,
"mean_inter_token_s": 0.4956,
"decode_tok_per_s_excl_prefill": 2.02,
"usage": { "prompt_tokens": 28, "completion_tokens": 200, "total_tokens": 228 }
}
首次 token 包含 28 token prompt 的 prefill，用时 5.6807 秒；随后平均每 token 约 0.4956 秒，decode 速度约 2.02 token/s。
并发也不会线性扩展：
单请求约 2.08 token/s，四个并发请求的总吞吐约 2.9 token/s，约为 1.39 倍；单个请求延迟从 72 秒增加到约 206 秒。20 个请求同时到达时，5 个被接纳，其他 15 个以 queue timeout 或 queue full 得到 HTTP 429。
17实际对话与性能边界
公开测试记录中有几组实际服务请求。例如 Python 素数函数请求耗时 119.17 秒，在 200 token 上限前停止；火车速度题正确完成单位换算；澳大利亚首都问题遵守了三句限制；haiku 请求得到 5-7-5 结构。
prompt: Explain in three sentences why a mixture-of-experts model activates only a few experts per token.
wall time: 84.68s
loaded in 11.33s | resident dense: 9912.75 MB | layers=78 experts=256 | MTP ACTIVE (draft=3)
24 tokens in 84.68s (0.28 tok/s) | expert hit rate 3.5% | RSS 16.09 GB
持久化 KV 也通过跨轮对话做了检查：模型能回忆前一轮的 hummingbird，并补充相关事实。
[smoke] 3 questions acc 66.7% acc_norm 66.7% [RSS 142.67 GB | hit 74%]
这只是 3 个问题的 smoke test，不能当作完整质量 benchmark。
同一台大机器上比较了三种放置方式：
CPU 20 GB: 24 tokens in 84.68s (0.28 tok/s) | hit 3.5% | RSS 16.09 GB
PROFILE: expert-disk 46.5s | expert-matmul 27.0s | attention 5.0s
CPU 200 GB: 24 tokens in 85.56s (0.28 tok/s) | hit 68.6% | RSS 181.92 GB
PROFILE: expert-disk 36.0s | expert-matmul 36.3s | attention 5.9s
两个 CPU 配置都是 0.28 token/s，但瓶颈不同：20GB 模式命中率只有 3.5%，主要受专家磁盘读取限制；200GB 模式命中率达到 68.6%，磁盘和 matmul 更接近均衡。
线程数继续增加也不会线性加速：
1 thread: 48 tokens in 887.83s (0.05 tok/s) | IPC 2.11
124 threads: 48 tokens in 227.46s (0.21 tok/s) | IPC 0.23
这组结果可以归纳为 memory wall：32 线程左右达到峰值，继续增加线程会因为内存带宽竞争而下降。
[CUDA] expert groups timing: H2D 36.6 ms | kernel 289.7 ms | D2H 58.2 ms
GPU 配置中，kernel 时间明显高于 H2D/D2H，瓶颈转向 GPU compute。
另有不同设备上的公开报告：Framework 13 约 0.37 token/s，Ryzen 9950X 桌面机约 0.10 到 0.28 token/s，Intel i5-12600K 原生 Windows 约 0.08 token/s，Apple M5 Max 通过 Metal backend 约 2.06 token/s。这些是不同硬件、磁盘和缓存条件下的分散结果，不是受控横向 benchmark。
18如何确认实现正确
验证时先在 PyTorch 中搭建同架构 tiny model：MLA、sparse indexer、sigmoid router、shared expert 都保留，只缩小权重，使 reference 可以快速运行。
关键 tensor 名称：
model.layers.3.self_attn.q_a_proj.weight
model.layers.3.self_attn.kv_a_proj_with_mqa.weight
model.layers.3.self_attn.kv_b_proj.weight
model.layers.3.self_attn.indexer.wq_b.weight
model.layers.3.mlp.experts.gate_up_proj
model.layers.3.mlp.gate.e_score_correction_bias
model.layers.3.mlp.shared_experts.gate_proj.weight
prefill：
PREFILL (teacher-forcing) C vs oracle: 32/32 positions | 723.8 pos/s
greedy decode：
Reference (oracle): 207 187 119 103 103 103 103 103 119 34 ...
GLM C engine : 207 187 119 103 103 103 103 103 119 34 ...
Matching tokens: 20/20 | Expert cache hit rate: 88.1% | 227.5 tok/s
这证明的是 C 实现和 tiny reference 在指定路径上保持一致，不是完整模型的通用能力评测。还需要注意：运行日志中的加载时间 banner 可能变化，而 token 本身保持一致；开启 int8 activation path 时，也可能因为量化噪声和浮点累加顺序导致单 token 差异。
19从 checkpoint 到服务
先编译并跑 tiny oracle：
make glm
SNAP=./glm_tiny TF=1 ./glm 64 16 16
PREFILL (teacher-forcing) C vs oracle: 32/32 positions | 723.8 pos/s
再分 shard 把 FP8 转成 int4，并单独转换 MTP：
python tools/convert_fp8_to_int4.py --repo zai-org/GLM-5.2-FP8 --outdir /nvme/glm52_i4 \
--ebits 4 --io-bits 8
python tools/convert_fp8_to_int4.py --repo zai-org/GLM-5.2-FP8 --outdir /nvme/glm52_i4 --mtp
规划资源并检查机器：
python resource_plan.py --model /nvme/glm52_i4
python doctor.py --model /nvme/glm52_i4 --gpu 0,1,2,3
20GB CPU-only 模式：
PROMPT="Explain in three sentences why a mixture-of-experts model activates only a few experts per token." \
NGEN=24 RAM_GB=20 SNAP=/nvme/glm52_i4 ./glm 64
loaded in 11.33s | resident dense: 9912.75 MB | layers=78 experts=256 | MTP ACTIVE (draft=3)
24 tokens in 84.68s (0.28 tok/s) | expert hit rate 3.5% | RSS 16.09 GB
四张 L40 的 GPU tier：
GLM_CUDA=1 GLM_GPUS=0,1,2,3 CUDA_EXPERT_GB=185 RAM_GB=200 \
PROMPT="Summarize the transformer architecture in about 150 words." \
NGEN=200 SNAP=/nvme/glm52_i4 ./glm 64
OpenAI-compatible server：
GLM_MODEL=/nvme/glm52_i4 GLM_API_KEY=local-secret \
python openai_server.py --host 127.0.0.1 --port 8000 --kv-slots 4
[CUDA] hot expert tier: 9780/15860 experts, VRAM 184.99 GB (total budget 185.0 GB)
OpenAI-compatible API listening on http://0.0.0.0:8000/v1
这条路径说明的不是“把 744B 直接塞进 20GB”，而是一套内存层级工程：int4 降低专家体积，磁盘保存睡眠专家，pread和缓存控制工作集，MLA 压缩上下文状态，RAM 和 VRAM 分别承载热专家，GPU 或 CPU 再承担实际矩阵计算。
在无 GPU、RAM 较小的机器上，主要瓶颈是磁盘冷读；RAM 增大后，瓶颈会转到 CPU matmul；VRAM 增加后，专家 compute 会更多地转移到 GPU。20GB 模式可以运行，但它依赖具体实现、int4 转换、流式读取、缓存预算和对低吞吐的接受度。四张 L40 是开发与测量设备，不是最低硬件要求。
相关仓库与模型：
C implementation — GitHub Model checkpoint — Hugging Face