说实话如果你打算把一个量化好的三元量化模型直接丢给 vLLM 加载第一眼看到的一定是报错key 对不上、dtype 不认、模型注册表里根本找不到对应的网络结构。这很正常因为三元量化ternary quantization和 GPTQ、AWQ 那种“压缩权重再用近似浮点去算”的思路完全不是一回事。权重变成 {-1, 0, 1} 之后线性层里的乘法基本可以改成符号选择加加减减算子语义都变了vLLM 现有的量化 kernel 覆盖不到这里。这篇文章是我把一个小 7B 级开源底座的三元量化版本真正塞进 vLLM 的完整记录核心就是三件事先从 safetensors 二进制里逆向出三元权重的打包规则再做数值对账确认解包和 scale 的配合没有半点偏差最后手写 CUDA kernel 把这条链路焊进 vLLM 的自定义模型推理流程里。整个过程踩了不少坑但走通之后收益非常明确——权重显存直接降到 FP16 的八分之一decode 单 token 延迟也有肉眼可见的下降。如果你正在折腾低比特量化落地、想在 vLLM 里跑非官方权重或者纯粹对自定义 kernel 集成推理框架感兴趣这篇应该能帮你省掉至少两三天试错时间。1. 三元量化不是普通量化vLLM 官方支持矩阵为什么没有它1.1 BitNet 带起的 1.58-bit 热潮与落地现实三元量化真正进入大众视野是 BitNet 那套 1.58-bit 工作带起来的。它把线性层权重约束到三个符号值-1、0、1。理论上的信息熵是 log2(3) ≈ 1.58 bit工程上为了对齐字节通常用 2 bit 存一个权重一个字节能 pack 四个权重。相比 FP16 的每权重 2 字节直接压缩 8 倍相比 INT8 的每权重 1 字节也有 4 倍优势。这个收益在 7B 级别模型上非常诱人。7B 模型的线性层attention 的 q/k/v/o 加上 MLP 的 gate/up/down大概能占到总参数的 90% 以上粗算 6.5B 参数FP16 下就是 13GB 显存。如果全部换成三元打包权重本体只需要 1.6GB 左右再算上 per-channel 的 scale整体也就 1.8GB 上下。对大模型推理来说这省出来的可不是一点点。但现实很骨感。模型权重开源出来了能把它跑起来的推理引擎却少得可怜。vLLM 作为目前部署大模型最主流的框架官方量化支持已经覆盖了 GPTQ、AWQ、FP8 这些路线但对三元量化一直没有原生方案。我拿到权重后查了一圈 issue发现社区也有不少人在问不过答案基本都是“自行实现”。这就是这个项目最开始的动机既然官方不做那就自己动手。1.2 vLLM 现有量化方案的底层假设要理解为什么 vLLM 没有原生支持三元量化得先看它现有量化方案共同依赖的一个假设。我整理了一张表方案权重存储形式核心假设计算方式GPTQ / Marlin4-bit 权重 group scale权重仍可反量化为近似 fp16反量化后做常规 GEMMAWQ4-bit per-channel scale同上但 scale 来自激活分布反量化后做常规 GEMMFP88-bit 浮点硬件本身支持 fp8 精度计算FP8 tensor coreInt8 / SmoothQuant8-bit 整数激活与权重用同一量纲缩放INT8 tensor core三元量化本项目2-bit 符号 per-channel scale权重是语义符号不是近似数值符号选择 加/减前四行的共同点是量化后权重依然是一个“近似数值”拿到之后可以反量化回接近浮点数的值然后扔给常规 GEMM 去算。整套 vLLM 的 kernel、模型加载、量化和反量化逻辑都是围绕这个假设设计的。三元量化打破了这个假设。权重是 -1/0/1它不是某个浮点数的低精度近似而是一个纯粹的符号语义。你当然可以把它反量化成 fp16 的 -1.0、0.0、1.0然后再跑 FP16 GEMM但这样做等于放弃了三元量化最大的红利乘法变加减法的计算优势。真正该做的是让 GEMM kernel 直接消费打包后的符号权重把“乘以 1 或 -1”变成“加上当前值或减去当前值”把“乘以 0”变成跳过从而省掉一大堆浮点乘法指令。也正因为如此三元量化不能简单通过“新增一个量化方法”接进 vLLM它需要从模型定义、权重加载到 kernel 执行一整条链路全部重写。1.3 我拿到的模型长什么样目标又是什么具体的项目目标是这样我手上有一个社区朋友基于开源底座训练导出的三元量化版本底座是 Qwen2.5-7B 的架构。发布方提供的权重结构大概是所有 transformer 层的 q/k/v/o、gate/up/down 线性层权重是 int8 dtypeshape 是[out_features, in_features / 4]每 4 个三元符号 pack 进一个字节。每个线性层带一个独立的 fp16 scale 张量长度等于out_features做 per-channel 缩放。embedding 和 lm_head 保持 fp16没有做三元量化。这部分对精度太敏感词表映射搞坏了整个模型基本就废了。我的目标非常明确把这个权重通过 vLLM 起一个 OpenAI 兼容的推理服务同时尽可能保留三元量化带来的显存收益。也就是说不能简单加载时解包成 fp16 然后假装是普通模型——虽然这个方案最简单但那样显存又回到 13GB做这个项目的意义就没了。2. 逆向 checkpoint从 int8 字节码还原三元打包规则2.1 第一眼看到的权重张量先说清楚这里的“逆向”不是破解什么东西而是对一个公开可加载的 checkpoint 做二进制格式分析。safetensors 本身是开放的权重文件里的 key、shape、dtype 都是明文我们要做的就是从字节排列里反推出打包规则。步骤完全合规性质上跟分析一个未知格式的录音文件差不多。拿到权重后我第一件事是写了一段脚本看张量的原始信息from safetensors import safe_open with safe_open(model-00009-of-00012.safetensors, frameworkpt, devicecpu) as f: w f.get_tensor(model.layers.10.self_attn.q_proj.weight) scale f.get_tensor(model.layers.10.self_attn.q_proj.weight_scale) print(w.shape, w.dtype) # torch.Size([4096, 1024]) torch.int8 print(scale.shape, scale.dtype) # torch.Size([4096]) torch.float16第一眼最奇怪的就是 shape。q_proj 的输入维度明明是 4096为什么权重张量第二维只有 1024答案其实很容易猜4096 / 1024 4也就是每个字节里 pack 了 4 个三元权重。这正好和 BitNet 系列常见的 2-bit packing 对上了。接下来就是验证这个猜测同时确认符号映射和位域顺序。2.2 2-bit 符号解包映射表与边界检查如果确实是每字节 4 个权重每个权重占 2 bit那么一个字节从低位到高位可以拆成 4 个 code00→ 符号 001→ 符号 110→ 符号 -111→ 非法/未使用我在验证这个映射前先做了一个简单的统计把所有权重字节拿出来按 2-bit 位域扫描看看有没有出现11这个组合。结果跑了整个 checkpoint11组合出现次数为 0。这是一个很强的信号说明合法三元值确实只有 {0, 1, -1} 三个第三个组合被刻意留空。解包的代码这样写import torch def unpack_ternary(packed: torch.Tensor) - torch.Tensor: 把 pack 后的 int8 张量解包成三元符号张量 [-1, 0, 1] b packed.to(torch.uint8) shifts torch.arange(0, 8, 2, devicepacked.device) # [0, 2, 4, 6] codes (b.unsqueeze(-1) shifts) 0b11 # [..., 4] signs (codes 1).to(torch.int8) - (codes 2).to(torch.int8) return signs.reshape(packed.shape[0], packed.shape[1] * 4)(codes 1)得到 1 的位置(codes 2)得到 -1 的位置两者相减正好把00、01、10映射到 0、1、-1。如果出现11它在两个布尔张量里都是 False相减等于 0但实际并不会发生。我在解包时还加了一个 assert统计codes 3的数量一旦非零说明我猜的打包规则是错的需要立刻停下来重新分析。这里还有一个小坑容易踩位域顺序。同样的字节最低位在前和最高位在前解出来的符号序列是反的。我当时用第一个字节的二进制手工推演了一遍确认是低位在前little-endian bit order。假设第一个字节是0b00011010从低到高的 2-bit 组就是10、10、01、01对应的符号是 -1、-1、1、1。这个顺序直接写进了解包函数后续对账部分会自动验证它对不对。2.3 两个容易翻车的细节padding 与未量化层逆向上来就翻车的主要是两个点值得单独写一下。第一个是 padding。打包的前提是输入维度能被 4 整除因为 4 个符号才凑一个字节。大部分主流模型的 hidden size 都是 4 的倍数7B 模型里常见的 3584、4096 也都满足。但如果遇到 hidden size 不是 4 的倍数或者因为 head_dim 切分导致某条路径的输入维度被 pad 过pack 的字节数就会和理论值对不上。这个情况在分析权重 shape 的时候就要留意如果in_features // 4不是整数通常说明原模型在导出时做过一次 padding解包后必须把 padding 位裁掉不然符号序列会整体错位。第二个是未量化层。不是所有 Linear 都被打包成三元这次拿到的模型里embedding 和 lm_head 保持 FP16。我一开始想当然地写了一个全量加载逻辑结果报错说lm_head.weight根本不是 int8。后来改成按 key 动态判断遇到q_proj/k_proj/v_proj/o_proj/gate_proj/up_proj/down_proj走三元解包路径遇到embed_tokens/lm_head走普通 FP16 路径。这里也提醒大家不要假设同一个 checkpoint 里所有层的数据格式是统一的导出脚本的取舍各不相同。3. 数值对账如何证明我解出来的权重是对的3.1 对账口径要跟谁比、误差落在哪逆向解包写完接下来不是急着写 kernel而是先回答一个问题我解出来的权重到底对不对这就要做数值对账。但“对不对”需要一个准确的口径。这里最忌讳的是拿解包结果去和原始 FP16 权重比——因为三元量化本身是有损压缩误差本来就是天量你比不出任何有效结论。真正要对着比的是量化作者在导出模型时使用的参考反量化实现。也就是说假设作者把权重从 FP16 三元化后才打包成 int8那他必然有一套“解包 乘 scale”的逻辑来恢复计算用的张量我要让我的解包结果和他那套逻辑的输出对齐。如果作者公开了反量化代码直接拿来对比就行。如果没有就得靠统计特征交叉验证。我当时验证了这么几条解包后的符号分布是否正常1 / -1 比例接近 1:10 的比例与发布方描述一致per-channel scale 是否全为正数解包后的符号张量乘 scale 之后与原始 FP16 权重在同位置上的正负号相关性是否足够高这几条都过了再进入逐层数值对比。3.2 参考实现与批量对账脚本我这里比较幸运发布方给了一段很短的反量化参考代码逻辑和我猜的基本一致只是 scale 的维度处理略有不同。我写了一个批量对账脚本遍历 checkpoint 里所有线性层def compare_layer(key, packed, scale, ref_dequant_fn): mine unpack_ternary(packed).float() * scale.view(-1, 1).float() ref ref_dequant_fn(packed, scale).float() diff (mine - ref).abs() return diff.max().item(), diff.mean().item()注意scale.view(-1, 1)这个细节权重张量的 shape 是[out, in]scale 对应的是每一个输出通道所以必须把 scale 变成列向量去做广播也就是按行乘。如果写成scale.view(1, -1)scale 会错位到输入通道那一维结果必然对不上。我跑完整个 checkpoint 后绝大多数层的max_diff在 1e-5 到 1e-4 之间。这个量级的误差完全来自 FP16 的舍入说明解包规则和参考实现完全一致可以放心进入下一步。对账脚本也顺手固化到了仓库里后面换权重版本时第一时间就能发现打包规则有没有变。3.3 一次真实的失败scale 的维度顺序这里想分享一次实际的对账失败因为排查过程对大家应该很有参考价值。第一次跑对账脚本时绝大部分层都通过了但有一层的max_diff高达 0.5而且出错位置非常规律每隔一段就会出现。我当时第一反应是解包逻辑有问题但奇怪的是同一层的 k_proj 和 v_proj 又全部通过。这就把范围缩小到了 q_proj 这一条单独路径上。对比之后发现问题出在 scale 的读取顺序。那一层的 q_proj 权重在导出时被打过维度重排scale 张量的排列方式和我的假设不一样。我原以为是“第 i 个输出通道对应第 i 个 scale 值”实际却是“输出通道已经被 split 成若干组scale 是按组内顺序排列的”。说白了就是 scale 的维度顺序和权重张量的实际布局不一致。最后解决方式也很简单加载那一层时把 scale 先做一次 reshape/permute对齐权重张量的维度顺序再乘进解包结果里。这个失败给了一个非常深刻的教训遇到对账误差先别急着怀疑 kernel 和浮点精度第一优先级要检查数据布局。数据的维度顺序错了后面所有优化都白搭。4. 手写 kernel从能跑的朴素版到值得上线的 tiled 版4.1 三元 GEMM 的计算特征对账确认后终于进入核心部分写三元 GEMM 的 CUDA kernel。常规 GEMM 是 Y[M, N] X[M, K] × W[N, K]^T每一组乘加都要做一次浮点乘法和一次浮点加法。三元 GEMM 的不同在于W 的每个元素只有 -1/0/1 三种取值所以乘法的结果只有三种可能取 X 当前值、取 X 当前值的相反数、取 0。这意味着每个乘法都可以变成code 1acc xvcode -1acc - xvcode 0什么都不做省掉一次浮点乘法代价是引入一次符号判断。如果判断没有处理好kernel 照样慢得没法看。所以全部功夫都花在怎么让符号判断尽量廉价甚至完全变成算术操作。另外一个重要特征是计算和访存的平衡会在不同阶段切换。decode 阶段 M1只有一个 token 参与计算瓶颈是访存要把 W 从 HBM 读进来把 X 的一整行读进来。这时候省权重内存就是省时间。prefill 阶段 M 很大计算量变成主导此时符号运算到底能省多少取决于 tile 设计好不好不能简单下结论。4.2 朴素 kernel先跑对写 kernel 的原则永远是先跑对再跑快。我第一版写了一个最直接的实现每个线程负责输出矩阵中的一个元素即一对(row, col)然后串行遍历 K 维度逐字节解包并累加。__global__ void ternary_gemm_naive( const half* __restrict__ X, // [M, K] const int8_t* __restrict__ W, // [N, K / 4] packed const half* __restrict__ scale, // [N] half* __restrict__ Y, // [M, N] int M, int N, int K) { int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; if (row M || col N) return; float acc 0.f; const half* x_row X row * K; const int8_t* w_col W col * (K / 4); for (int k4 0; k4 K / 4; k4) { unsigned char byte (unsigned char)w_col[k4]; #pragma unroll for (int j 0; j 4; j) { int code (byte (j * 2)) 0b11; float xv __half2float(x_row[k4 * 4 j]); if (code 1) acc xv; else if (code 2) acc - xv; } } Y[row * N col] __float2half(acc * __half2float(scale[col])); }这个版本能跑但性能很惨。实测在 4096 × 4096 的 GEMM 上比 cuBLAS 的 FP16 GEMM 还要慢 30% 到 40%。原因很简单每个线程每次都单独访问 X 行和 W 列的全局内存没有向量化加载也没有 shared memory 复用符号判断虽然不会有真正的 warp divergence因为同一个字节解出的四个 code 对相邻线程是类似的判断路径不会大规模分叉但指令开销并不小。这个朴素版本不是拿来用的是拿来当基准和正确性参照物的。我在这个版本的基础上加了和参考实现的对账确认它能算出正确结果才开始做性能优化。4.3 tiled 版本shared memory 复用 位运算符号选择tiled 版本的思路和普通 GEMM 优化完全一样让一个 block 处理一小块输出 tile把对应的 X tile 和 W tile 先加载到 shared memory然后反复复用减少全局内存访问。不同的地方主要在符号选择的处理上。我的做法是 block tile 取BM64, BN64, BK32X tile 用float4向量化加载到 shared memoryW tile 保持 int8 packed 形式放进去scale tile 单独放一份。内层循环里每次拿到一个 code不再做分支判断累加而是用位运算把符号选择透明化float xv __half2float(xs[k4 * 4 j]); // code: 0 - 0, 1 - xv, 2 - -xv float contrib (code 1) ? xv : 0.f; if (code 2) contrib -contrib; acc contrib;这里的code 1判断是否非零code 2判断是否要取负两条逻辑都是极廉价的位运算加一次条件赋值比浮点乘法便宜得多也比if/else分支更容易被编译器优化。这套设计跑下来prefill 场景的提升非常明显。M1024、N4096、K4096 的 GEMM 实测比 cuBLAS FP16 快了约 30%。但 decode 场景反而不理想原因很简单M1 时一个 block 只需要处理一行输出shared memory tile 的优势发挥不出来反而多了一堆 tile 加载和同步开销。所以我又单独写了一版 decode 专用 kernel。decode 专用 kernel 的思路回到朴素版的“一个线程算一个输出元素”但做了两个改进一是不再拆块加载X 的整行在 L2 里可以被大量线程复用实测冗余读并没有想象中严重二是去掉 shared memory直接用__ldg走只读缓存路径。最终 decode 单 token 延迟从 FP16 cuBLAS 的 85us 左右降到 58us 左右降幅约 30%。4.4 实测数据decode 提速多少prefill 为什么变化不大我在 A100 80G 上跑了一组对比结果如下场景K4096, N4096FP16 cuBLAS三元朴素三元 tileddecode 专用M1decode~85us~120us~90us~58usM1024prefill~70us~110us~45us不适用这里要解释一下为什么 prefill 提升明显而 decode 只提升 30% 左右。decode 是一个天然 memory-bound 的场景。一个 token 的 hidden state 只有 8KB权重按三元 pack 后也只有原来的八分之一但读取 X 的那一次 8KB 无法继续压缩它在延迟里的占比被抬高了。换句话说省下来的主要是 W 的读取时间X 的读取时间还是那个固定开销。所以 decode 的加速比不可能像权重压缩比那样到 8 倍能省 30% 已经很不错。prefill 是典型的 compute-bound 场景。M 足够大计算时间主导一切而三元 GEMM 每 k 步省掉的浮点乘法在这种场景下才能充分兑现。再加上 tiled 设计把访存重叠得很好45us 对 70us 的差距就是这么来的。顺带算一笔显存账7B 模型的线性层权重如果按 FP16 存大约是 13GB三元 pack 后权重本体只有 1.6GB 左右加上 per-channel scale 和少量 meta 信息总共 1.8GB 上下。这个数字才是整个项目里最确定的、不可辩驳的收益。5. 焊进 vLLM模型注册、权重加载与服务化验证5.1 两条路线假 FP16 还是真三元kernel 写好了接下来要把它接进 vLLM。这里有一个路线选择而且在做之前很容易低估两条路线的差距。路线一加载时解包把三元权重反量化成 FP16然后当普通模型跑。优点是对 vLLM 完全透明你只需要在权重加载层做一个 hook把 int8 packed 权重解包成 fp16 替换进去。缺点是显存收益清零解包后的权重还是 13GB。如果只是为了业务快速上线这条路线完全 OK但我不甘心。路线二保持 int8 packed 权重常驻显存自定义一个 QLinear 模块forward 里调用三元 kernel。真正的显存收益和计算收益都来自这条路线但工作量也大得多——需要把模型类里所有 Linear 替换掉还要保证 vLLM 的 model executor 能识别这个自定义结构。我最后选了路线二。5.2 注册自定义模型类的关键代码vLLM 的模型注册机制相对清晰register_model可以挂一个新的模型类。我的做法是继承同底座的 Qwen 模型类只重写产生 Linear 的地方from vllm.model_executor.models.registry import register_model from torch import nn import torch class TernaryLinear(nn.Module): 替换标准 Linear权重保持 int8 packedforward 调用三元 kernel def __init__(self, input_size, output_size): super().__init__() self.input_size input_size self.output_size output_size # input_size // 4因为每 4 个三元符号 pack 成一个 int8 self.weight nn.Parameter( torch.empty(output_size, input_size // 4, dtypetorch.int8)) self.scale nn.Parameter(torch.empty(output_size, dtypetorch.float16)) # 初始化略 def forward(self, x): # x: [num_tokens, input_size] fp16 return ternary_gemm_forward(x, self.weight, self.scale) register_model(QwenTernaryForCausalLM) class QwenTernaryForCausalLM(Qwen2ForCausalLM): # 继承原结构在 __init__ 中把 nn.Linear 替换为 TernaryLinear ...权重加载时有一个关键点不要在线加载时做任何 unpack。safetensors 里的 int8 权重直接copy_进self.weight参数scale 直接进self.scale只有 forward 时才会被 kernel 消费。这样权重常驻显存的形态始终是打包后的 1.8GB不会出现中间膨胀。还有一个容易被忽视的坑vLLM 在计算模型显存占用时会根据model.get_model().get_memory_footprint()来估算权重占了多少显存进而决定 KV cache 能预留多少。如果你不覆写这个方法vLLM 会按默认 dtype比如 fp16去估算 7B 模型的权重得出一个 13GB 级别的错误数字导致 KV cache reserve 严重偏小。我这里覆写了它返回 int8 packed 权重 scale 的真实字节数。5.3 scheduler 一个字母都不用改但别高兴太早有个热词叫 “vLLM scheduler 逻辑”很多人在讨论低比特模型时都会问是不是要改 scheduler 才能适配。这里可以给出一个明确结论不用改一个字母都不用改。vLLM 的 scheduler 只负责序列级别的调度、KV cache 分配、preemption 决策它完全不感知模型内部权重是 FP16 还是三元打包。只要你的自定义模型类返回标准的 hidden_states 和 logitscontinuous batching、paged attention 这些机制会照常工作。省显存这件事发生在权重存储层面和 KV cache 管理是两个完全不相干的维度。但这不意味着没有坑。最大的坑是显存预算计算就是上面说的get_memory_footprint()。另一个坑是框架版本兼容。vLLM 的模型注册和 Linear 抽象层在不同版本之间变过好几次我用的那版是自定义模型能注册进去的但升级版本时这些接口很可能又变了。我的应对方式是把 kernel 和自定义模型封装在独立仓库里vLLM 只作为依赖引用每次升级前先跑一遍对账脚本和端到端测试再决定。5.4 端到端验证与显存核算集成完成后我从头到尾验证了一遍用本地权重目录启动 vLLM--dtype float16--max-model-len 4096。先跑一个极短的序列把模型的最后一层 logits dump 出来和 HF 参考实现的输出对比top-1 一致率必须达到 100%。用 OpenAI 兼容接口发起多轮请求确认 chat template、采样参数、停止符这些都正常。用nvidia-smi观察显存占用。实际结果中7B 模型启动后权重部分只占 1.8GB 左右加上 KV cache 和激活总显存占用大约 5GB而同样配置下 FP16 原版要接近 16GB。对于一个 7B 级别模型来说这个差距直接决定了一台卡上能塞多少个并发模型实例。权重形式线性层权重占用启动后总显存约FP16 原版~13GB~16GB三元 packed本项目~1.8GB~5GB这里的数字依赖--max-model-len和 KV cache 配置不同参数下会有浮动但量级差距不会变。6. 复盘这个项目里真正值得复用的东西6.1 收益账显存、延迟、吞吐分开算整个项目下来收益最确定的是显存。线性层权重从 13GB 级降到 1.8GB 级对部署方来说意味着同样的 GPU 显存可以跑更多并发实例或者给每个实例分配更大的 KV cache这些都会直接换算成成本下降。延迟方面decode 单 token 时间在 A100 上降了约 30%。这个加速比说不上夸张但在推理服务里30% 的延迟改善对用户体验的影响是实实在在的。prefill 场景提升更明显tiled kernel 相比 FP16 cuBLAS 能快 30% 到 40%。但是要泼一盆冷水吞吐和延迟不能画等号。一旦 batch 变大访存竞争加剧三元 kernel 的收益会被稀释。我自己的实测是batch size 超过 16 之后端到端 QPS 的提升就明显收窄了。所以如果你想靠三元量化来解决高并发吞吐问题预期要放低一些如果你的瓶颈是显存不够、并发实例数上不去那这个方案的收益会非常直接。6.2 代价账维护成本与框架升级风险收益之外也要算代价。最主要的是维护成本vLLM 每升级一个版本模型注册接口、Linear 抽象、甚至get_memory_footprint的调用路径都可能变化需要同步适配。我见过一些团队因为这个原因干脆把 vLLM 版本锁死在一个老版本上但那样又牺牲了新功能比如更好的调度策略或新的 attention 实现。其次是 kernel 的 shape 特殊性。我为 4096 维度的 GEMM 做了不少针对性调优换一个 hidden size 或者 head_dim 不同的模型性能特征很可能完全不一样需要重新调 tile 大小和向量化宽度。这不是一个“跑遍天下模型都通用”的 kernel。最后是模型权重本身的稳定性。三元量化模型目前没有像 GPTQ 那样形成事实标准不同团队导出的权重packing 规则、scale 维度、甚至是否做了 padding 都可能不同。对账脚本必须作为门禁留下来每次拿到新权重先跑一遍确认格式没漂移再继续下一步。6.3 给想抄作业的人的建议最后按照我自己的体会给想复现这条路的朋友三个建议。第一先对账再写 kernel。我在这件事上的排序是数据布局问题必须在数据层解决只有对账完全通过kernel 才能安心调优。否则后面 kernel 出了任何 bug你都要花大量时间去区分到底是 kernel 错了还是权重本身就错了。第二从 1B 左右的模型起步。第一次做逆向和 kernel 集成用 7B 模型会让每次调试迭代都慢到怀疑人生。先用 0.5B 到 1B 级别的小模型把整条链路打通、把显存估算和模型注册跑顺再切到大模型效率高很多。第三区分“先跑通业务”和“真做优化”。如果只是想让一个三元量化模型尽快上线提供服务路线一加载时解包成 FP16完全足够你甚至不需要写 kernel。只有当显存收益成为刚需、或者你想真正研究低比特推理时才值得投入路线二。业务优先级永远高于炫技这一点是我折腾完这个项目后最深的体会。