Kestrel_VLLM - 18 · 投机解码位级对拍收口：batch 与 M=1 一致，注意力路径与 CPU 逐字节对齐

18 · 投机解码位级对拍收口：batch 与 M=1 一致，注意力路径与 CPU 逐字节对齐

一句话：把「M>1 batch verify 与 M=1 decode」的对拍从 top-1 下沉到完整 logits 逐字节，纠正一处 64-token 长序列的误读，并把 CPU↔GPU 的分歧一路消到 layer 0 的 MoE FFN 为止——层 0 整条注意力路径（RMSNorm → QKV GEMM → QK-norm → RoPE → KV 量化 → attention → O 投影+残差）已与 CPU 逐位全等。 前置：建议先读 第 16 篇（batch verify 的 top-1 位级一致）。 环境：x86-64 + NVIDIA RTX 5090（CUDA 后端）· Qwen3-30B-A3B q4local · 开关 VLLM_CUDA_ATTN_BITEXACT=1 / VLLM_CUDA_L2EXACT=1（位级档，opt-in）。

一、问题与结论

状态：注意力路径 ✅ 位级全等（opt-in 档）；MoE FFN 在本篇仍 ❌ 未对齐（已登记，下一篇收口）。

二、背景：为什么 top-1 对拍不够

第 16 篇 只对拍了 top-1（[SPECTOP1]/[ISOTOP1]）。top-1 相同不构成位级证明：仓库方法论的「位级门禁」原则说得直白——两个 logits 可能在 top-1 之外有 ULP 差，只在边界 token 处翻 top-1。于是本篇加 [ISODIFF] 探针：batch 列 j 的完整 logits 与「同 token 顺序 M=1 decode」逐 uint32 字节比，报第一个非零 ULP 差的位置。

[SPECTOP1] pos=33 col=0 top1=99692 [ISOTOP1] col=0 99692 [ISODIFF] col=0 ndiff=0
[SPECTOP1] pos=33 col=1 top1=99371 [ISOTOP1] col=1 99371 [ISODIFF] col=1 ndiff=0
…（pos 33–38，col 0–4 全 ndiff=0）

→ batch 列与顺序 M=1 的完整 logits 逐字节相同，位级一致成立。

三、核心机制

3.1 误读纠正：token 58 不是翻转

实测：

「参考 81167」是 token21 的值。token58 两边一致，不存在翻转。

3.2 唯一真实分歧：t=27，源在 CUDA attention

实测：

隔离对照（MoE/lm_head 全留 CPU，实测）：

→ 分歧源是 CUDA attention，不是 MoE GEMM，也不是 batch verify。

对拍口径上的判定：

• batch verify vs M=1（同为 CUDA）：ndiff=0，位级一致——投机正确性红线不破。
• CPU↔GPU：本篇之前是 greedy 级（跨硬件 ULP 差）。本篇把它提升到位级（opt-in）。

3.3 消除 t=27：逆向 UCRT 的 expf

第一步：算子顺序对齐（不够）。 既有 k_attn_decode_f32q 是两遍 softmax（全局 max → exp → 并行归约求和），归约树求和序与 CPU 的 online softmax 不同。新增 k_attn_decode_online，逐字复刻 CPU 的 flash_attn_single_q_q8_neon：int8 query 量化 + int32 QK 点积 + 逐 token online softmax + 顺序 PV，每 head 一 block、thread 0 串行扫整条 KV。

用 device expf 时反而更差（实测，t=24 就翻）：

探针：量化 expf 差异（写对拍探针，400002 组 [-100,0] 含随机/密集近 0，实测）：

结论：CUDA 的 expf 与 UCRT 差很多；(float)exp(double)（= 正确舍入）能把差异从 28.3% 压到 0.007%（换成它后 64 token 一致数从 26 提到 56，分歧从 t=24 推到 t=57）。但 UCRT 的 expf 本身不是正确舍入（28 例差 1 ULP），且「到舍入边界的距离」在 bad/ok 之间不可分离——它不是「double 结果加噪声」，而是独立的 float 例程。

第二步：反汇编 ucrtbase.dll，位级复刻。 用 objdump + dumpbin 定位 ucrtbase.dll 的 expf（RVA 0xA84F0，SSE2 路径），读出算法与常量：

t = x*K; n = rint(t); j = n&63; m = n>>6;
r = fma(-n, C2, x); z = r*r;
p = fma(fma(r,C0,C1), z, r);
res = 2^m * fma(T[j], p, T[j]); return (float)res;
K = 64/ln2, C2 = ln2/64, C0 = 1/6, C1 = 1/2, T[j] = 2^(j/64)（64 项表）

即「3 项泰勒 + 64 项表」，内层全 double + 显式 FMA——所以它的相对误差 ~r⁴/24 ≈ 3.6e-11，这也解释了它为何「几乎正确舍入但差 28 例」。复刻要点（device）：__double2int_rn（对应 vcvtpd2dq）、__fma_rn（对应 vfmadd213sd/vfnmadd231sd；-fmad=false 只禁自动融合，显式 FMA 不受影响）、__longlong_as_double 造 2^m，64 项表按位型内嵌。特殊值仅在 x<-104/x>88.72 退化。

仓库里落地后的 device 复刻（src/npu/cuda/vllm_cuda_kernels.cu，64 项 EXPF_T 表略）：

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu */
/* UCRT expf 的位级复刻。逆向自 C:\\Windows\\System32\\ucrtbase.dll 的 expf
* （RVA 0xA84F0，SSE2 路径）。host 复刻对 400002 组 [-100,0] 样本与 UCRT
* expf 逐位一致（0 差异）。 */
__device__ __forceinline__ float expf_ucrt(float xf) {
if (xf > 88.7228394f) return (float)exp((double)xf); /* 上溢 -> +inf */
if (xf < –104.0f) return 0.0f; /* 深度下溢 */
const double K = 9.23324826168936567683e+01; /* 64/ln2 */
const double C2 = 1.08304246962491450973e-02; /* ln2/64 */
const double C0 = 1.66666666666666657415e-01; /* 1/6 */
const double C1 = 5.00000000000000000000e-01; /* 1/2 */
double x = (double)xf;
int n = __double2int_rn(x * K); /* 对应 vcvtpd2dq（就近偶数） */
double dn = (double)n;
double r = __fma_rn(–dn, C2, x); /* 对应 vfnmadd231sd */
int j = n & 63;
int m = n >> 6;
double z = r * r;
double t1 = __fma_rn(r, C0, C1);
double p = __fma_rn(t1, z, r);
double Tj = __longlong_as_double((long long)EXPF_T[j]);
double q = __fma_rn(Tj, p, Tj);
double two_m = __longlong_as_double((long long)(unsigned long long)(m + 0x3ff) << 52);
return (float)(two_m * q);
}

结果（实测）：

解码文本也与 CPU 完全相同。代价：ctx=1024 下 89 → 200 ms/tok（2.25×，double exp + 每 head 串行扫描）。按仓库方法论的默认化门禁（需「位级等价 且 性能领先」），位级路径性能不领先，故不默认化，保留为 opt-in：默认仍走原两遍 softmax（greedy 级，与改动前逐 token 一致）。

3.4 L2 完整模型对拍 harness

前面几节把 attention 做到位级，但那不是 L2。按仓库方法论的 L2 口径，L2 = 逐层 hidden + 完整 logits + KV 行哈希全等。

harness（两侧同 key、同常数）：VLLM_L2_DUMP=<path>（默认关）：对缓冲求 FNV-1a 64，按 <tag> <hash16> 追加写；CPU 与 CUDA 两侧使用同一批 key，diff 文本即得「第一个差异落在哪层哪步」。CPU 侧在 st_qwen_model_forward 直接算；GPU 侧在 vc_decode_body 用单线程 FNV 内核写到设备暂存、token 末一次性 D2H（避免逐层同步）。

两侧同常数的哈希实现（src/model/vllm_safetensors.c 与 src/npu/cuda/vllm_cuda_kernels.cu）：

/* 文件名：src/model/vllm_safetensors.c */
static uint64_t l2_hash64(const void *p, size_t len) {
const uint8_t *b = (const uint8_t *)p;
uint64_t h = 1469598103934665603ull;
for (size_t i = 0; i < len; i++) { h ^= b[i]; h *= 1099511628211ull; }
return h;
}
void st_l2_dump(const char *tag, const void *p, size_t bytes) {
const char *path = getenv("VLLM_L2_DUMP");
if (!path || !path[0] || !p || !bytes) return;
FILE *f = fopen(path, "a");
if (!f) return;
fprintf(f, "%s %016llx\\n", tag, (unsigned long long)l2_hash64(p, bytes));
fclose(f);
}

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu，设备侧同常数 FNV-1a 64 */
#define L2_FNV_OFF 1469598103934665603ull
#define L2_FNV_PRM 1099511628211ull
__global__ void k_l2_hash(const unsigned char *__restrict p, int nbytes,
unsigned long long *__restrict out) {
if (threadIdx.x != 0 || blockIdx.x != 0) return;
unsigned long long h = L2_FNV_OFF;
for (int i = 0; i < nbytes; i++) { h ^= (unsigned long long)p[i]; h *= L2_FNV_PRM; }
*out = h;
}

每个 decode token 的 key 集合（48 层 → 每 token 530 key）：

3.5 前提：对拍配置本身必须干净（VLLM_CUDA_FORCE_CPU=1）

这是本轮踩过的最大的坑。 全 GPU decode 需要 VLLM_CUDA_INFER=1（把权重预载驻留），但该开关同时也会让 prefill 的 GEMM 走 GPU offload——于是 prefill 写入 KV cache 的 K/V 行与 scale 就已经和纯 CPU 参考差 ~1 ULP。此时无论 decode 侧多干净，L2 都必然不等，且首个差异会伪装成 t32_l0_ks（看起来像 RoPE/attention 的问题），把定位带偏。

正解：加 VLLM_CUDA_FORCE_CPU=1——设备仍初始化、权重仍预载，但所有 host 侧 offload 回退 CPU（prefill 因此纯 CPU）；而 decode 走的是设备内的 vc_decode_body，不受该开关影响，仍全 GPU。位级对拍的正确命令：

# 参考：纯 CPU
VLLM_L2_DUMP=ref.txt vllm_kestrel_x64 –model <m> –stream-test –stream-n 4
# 被测：全 GPU decode + 位级档（prefill 留 CPU）
VLLM_L2_DUMP=gpu.txt VLLM_CUDA_DECODE=1 VLLM_CUDA_INFER=1 \\
VLLM_CUDA_FORCE_CPU=1 VLLM_CUDA_L2EXACT=1 \\
vllm_kestrel_x64 –cuda –model <m> –stream-test –stream-n 4

3.6 逐项消差（VLLM_CUDA_L2EXACT=1 档）

按 harness 钉出的顺序一个个改，每改一次复跑、看首个差异是否后移：

仓库里这两个「同序」内核（src/npu/cuda/vllm_cuda_kernels.cu）：

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu，k_rms_norm_cpu / k_qk_norm_cpu */
__global__ void k_rms_norm_cpu(float *__restrict out, const float *__restrict x,
const float *__restrict w, int d, float eps) {
__shared__ float rd[1];
const int i = threadIdx.x;
if (i == 0) {
float ss = 0.0f;
const int d8 = d & ~7;
for (int b = 0; b < d8; b += 8) { /* CPU 的 8 元素分块 hsum8 树 */
float lo[4], hi[4];
for (int k = 0; k < 4; k++) { lo[k] = x[b+k]*x[b+k]; hi[k] = x[b+4+k]*x[b+4+k]; }
for (int k = 0; k < 4; k++) lo[k] += hi[k];
const float s01 = lo[0] + lo[1], s23 = lo[2] + lo[3];
ss += s01 + s23;
}
for (int j = d8; j < d; j++) ss += x[j] * x[j];
rd[0] = 1.0f / sqrtf(ss / (float)d + eps); /* 与 UCRT 逐位一致；GPU rsqrtf 差 34.9% */
}
__syncthreads();
const float rms = rd[0];
for (int j = i; j < d; j += DC_BLK) out[j] = x[j] * rms * w[j];
}
__global__ void k_qk_norm_cpu(float *__restrict x, const float *__restrict w,
int hd, int nheads) {
const int h = blockIdx.x;
if (h >= nheads) return;
float *xh = x + (size_t)h * hd;
__shared__ float rd[1];
const int i = threadIdx.x;
if (i == 0) {
float ss = 0.0f;
for (int j = 0; j < hd; j++) ss += xh[j] * xh[j]; /* 逐元素升序 ss += x[i]^2 */
rd[0] = 1.0f / sqrtf(ss / (float)hd + 1e-6f);
}
__syncthreads();
const float rms = rd[0];
for (int j = i; j < hd; j += DC_BLK) xh[j] = xh[j] * rms * w[j];
}

结果：层 0 内每加一个探针就把范围缩小一次——_nrm ok 后从 _kraw 继续，_kraw ok 后再看 _qrope，_qrope ok 后再看 _attn，最后只剩 _x。

3.6.1 UCRT cosf/sinf 的反汇编复刻

与 3.3 的 expf 同一套方法。GetProcAddress(ucrtbase) 定 RVA：cosf = 0xa7d20、sinf = 0xac110（expf = 0xa84f0，与 3.3 一致，验证方法自洽）。读出的算法：

| x | <= pi/4 : 直接双精度多项式
|x| < 2^-13 : cosf -> 1.0f ；sinf -> x（恒等）
|x| < 2^-7 : cosf -> (float)fma(-(x*0.5), x, 1.0) ；
sinf -> (float)fma(-(x*x*x), 1/6, x)
否则 : z=x*x；cos 多项式 p = A3 + z(A5 + z(A7 + z*A9))，结果 fma(p, z*z, 1 – 0.5*z)
sin 多项式 p = B3 + z(B5 + z(B7 + z*B9))，结果 fma(p, x*z, x)
| x | > pi/4 : Cody-Waite 归约
n = (int)trunc( fma(|x|, 2/pi, 0.5) ) ; r = fma(-n, PIO2_HI, |x|) – n*PIO2_LO
cosf: n 奇 -> sin 多项式(参数 r)，偶 -> cos 多项式；符号位 ^= ((n&3)+1)>>1
sinf: n 奇 -> cos 多项式(参数 r)，偶 -> sin 多项式；符号位 ^= (n&3>=2) ^ sign(x)
最终一次 (float) 舍入

常量（.rdata 按位型读出）：2/pi = 0x3fe45f306dc9c883、PIO2_HI = 0x3ff921fb54400000、PIO2_LO = 0x3dd0b4611a626331、A3..A9 = 1/24, -1/720, 1/40320, -1/3628800、B3..B9 = -1/6, 1/120, -1/5040, 1/362880。

复刻要点：多项式系数必须按「213 形式」写——vfmadd213sd xmm1, xmm3, mem 是 xmm1 = xmm3*xmm1 + mem，不是 xmm3*mem + xmm1；写反会让整个多项式量级错（本次第一版就把 cosf 做成 98% 不一致，定位到这条即修好）。|x| ≥ 0x4f490fdb（cosf）/ 0x4b800456（sinf）走 Payne-Hanek 慢路，RoPE 不可达（angle ≤ max_seq），退化为 device 库函数。

仓库里落地后的 device 复刻（src/npu/cuda/vllm_cuda_kernels.cu，sinf_ucrt 同法）：

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu，cosf_ucrt() */
/* host 侧对照真实 UCRT 对拍 600002 组样本 → cosf/sinf 各 0 差异。 */
static __device__ __forceinline__ double vc_ucd(unsigned long long b) {
return __longlong_as_double((long long)b);
}
__device__ __forceinline__ float cosf_ucrt(float xf) {
const double INV_PIO2 = vc_ucd(0x3fe45f306dc9c883ull);
const double PIO2_HI = vc_ucd(0x3ff921fb54400000ull);
const double PIO2_LO = vc_ucd(0x3dd0b4611a626331ull);
const double A3 = vc_ucd(0x3fa5555555555555ull); /* 1/24 */
const double A5 = vc_ucd(0xbf56c16c16c16c16ull);
const double A7 = vc_ucd(0x3efa01a01a01a019ull);
const double A9 = vc_ucd(0xbe927e4fb7789f5cull);
...
const unsigned au = __float_as_uint(xf) & 0x7fffffffu;
if (au >= 0x7f800000u) return cosf(xf);
const double xd = (double)xf;
if (au <= 0x3f490fdbu) { /* |x| <= pi/4 */
if (au < 0x39000000u) return 1.0f; /* |x| < 2^-13 */
...
const double z = xd * xd;
double p = __fma_rn(z, A9, A7);
p = __fma_rn(z, p, A5);
p = __fma_rn(z, p, A3);
return (float)__fma_rn(p, z * z, 1.0 – 0.5 * z);
}
...
}

验证（实测）：host 复刻对 UCRT 对拍 600002 组样本（[-10000,10000] 均匀 + [-35,35] 密集 + 随机）→ cosf/sinf 各 0 差异；随后移植为 device cosf_ucrt/sinf_ucrt。RoPE 内核里按档位选择（src/npu/cuda/vllm_cuda_kernels.cu）：

/* 文件名：src/npu/cuda/vllm_cuda_kernels.cu，k_rope(…, bitexact) */
const float ca = bitexact ? cosf_ucrt(angle) : cosf(angle);
const float sa = bitexact ? sinf_ucrt(angle) : sinf(angle);

四、实测数据

实测（–stream-n 4，48 层 → 每 token 530 key，4 token 共 2120）：

层 0 逐字段（数据流序）：

ok _nrm ok _kraw ok _knorm ok _k ok _ks ok _v ok _vs
ok _qrope ok _attn ok _ores
DIFF _x <– 唯一残项：MoE FFN 输出

结论：全 GPU decode 在 opt-in 档（VLLM_CUDA_L2EXACT=1）下，层 0 的整条注意力路径（RMSNorm → QKV GEMM → QK-norm → RoPE → KV 量化 → attention → O 投影+残差）已与纯 CPU 逐字节一致；沿数据流继续推进，第一个差异点就是 MoE FFN。

五、边界与已知限制

已完成（位级，opt-in）

未完成 / 已知墙

CPU 对照（迁移前基线）

• CPU 参考：kestrel-llm/src/model/vllm_safetensors.c:f32x8_hsum / hsum_neon4 —— CPU 归约序（8 元素树 / 4 元素树）；超越函数直接用 UCRT 库 expf/cosf/sinf。
• 迁移要点：跨硬件位级的拦路虎是超越函数——反汇编 ucrtbase.dll 位级复刻 expf_ucrt / cosf_ucrt / sinf_ucrt；RMSNorm/QK-norm 改串行同序（k_rms_norm_cpu / k_qk_norm_cpu）；QKV GEMM 加 pairwise+nsplit=1；attention 改 online softmax 镜像。
• 真机验证：未独立复现：需 30B + VLLM_L2_DUMP（530 key/token，2120 key 一次跑完）与位级档 A/B。

六、小结

相关篇目：第 16 篇、第 17 篇、第 22 篇 源码与配套资源：本仓库 https://gitee.com/pei-xiaoguang/kestrel-llm-cuda.git； CPU 推理源码 https://gitee.com/pei-xiaoguang/kestrel-llm