amd-infer

一个从零写的 LLM（大语言模型）解码引擎，专门针对 AMD RDNA4 显卡（gfx1201，例如 RX 9070 XT）。 它用 C 语言写成。它的索引、量化和内存规划逻辑来自一个经过形式化验证的核心（core/ie_core.h）。

名词

• GGUF：llama.cpp 使用的模型文件格式。本引擎读 GGUF v3。
• Q4_0：4 位权重量化。每 32 个权重为一块，一块有一个 f16 缩放因子。
• Q8：8 位激活量化。每 32 个激活为一块，d = amax/127，q = round(x/d)。
• GEMV：矩阵乘向量。解码时，几乎全部时间都花在 GEMV 上。
• wave：RDNA 上 32 个线程一起执行（wave32）。
• 工作组（workgroup）：256 个线程，也就是 8 个 wave。
• LDS：工作组内共享的片上内存。
• arena：一块连续内存。一步解码的所有激活缓冲区都放在里面。
• 验证核心：core/ie_core.h。它的定律（laws/ie_laws.cpp）在 Lean 4 中被证明，对所有输入成立。
• hsaco：AMD GPU 的代码对象文件（ELF）。

文件

文件	内容
core/ie_core.h	验证核心（C 与 C++ 共用一个头文件）。本项目不修改它。
laws/ie_laws.cpp、proofs/	定律与 Lean 证明（另一位开发者维护）。
src/gguf.c/.h	mmap 并解析 GGUF v3：所有类型的元数据（含数组）、张量信息、对齐、数据偏移。
src/model.c/.h	Llama 与 Qwen2 模型；加载时把 Q4_0 重排成 GPU 布局；构造一步解码的算子列表。
src/plan.c	激活 arena 的静态规划（贪心），然后用 ie_plan_check 检查。
src/cpu.c	CPU 后端。算法和工作划分与 GPU 内核相同。
src/hip.c	GPU 后端。运行时用 dlopen 加载 libamdhip64.so。
src/main.c	命令行工具 ie-run。
src/gen.c/.h	采样、预填充、MTP 推测解码（ie-run 和 ie-serve 共用）。
src/serve.c、src/tok.c、src/json.c	HTTP 服务 ie-serve（OpenAI 兼容接口）、分词器、JSON。见 SERVE.md。
kernels/ie_kernels.c	GPU 内核，用 C 写（不是 HIP C++），由 clang 编译到 amdgcn。
kernels/ie_mk.c	大核函数（默认开启，IE_MK=0 关闭）：解码一步或验证一步的核函数在一次常驻启动里按顺序执行。
tools/mk_gen.py	生成大核函数的核函数表（参数布局），设备端和主机端共用。
prof/sm_profiler.h、prof/sm_profiler.c	核函数内的逐 wave 事件时间线（参照 leepoly/sm-profiler 改写给 AMD）：导出 Chrome 追踪 JSON，用 ui.perfetto.dev 打开。
tools/mk_bench.c	大核函数探针：32 个矩阵向量乘，一次启动对比逐个启动。
tools/make_tiny_gguf.py	生成小的随机模型（GGUF v3），两种架构都有。
tools/ref_forward.py	独立的 numpy float64 参考前向计算。
tools/check_objdump.sh	检查 GEMV 内核的机器码。
tools/gdn_chunk.py	在 numpy 中检查分块并行的线性注意力等式（与逐 token 递推比较）。
tests/kernel_test.c	CPU GEMV 与逐元素点积对比；重排后的半字节与 GGUF 对比。
tests/laws_test.cpp	在随机输入上运行所有定律（g++）。

构建

需要：gcc（或 clang）、clang 18 以上（带 amdgcn 目标）、ld.lld、llvm-objdump、python3 + numpy + gguf（pip）。 构建时不需要 ROCm。

make          # build/ie-run（CPU 引擎 + GPU 主机代码）和 build/ie_kernels.hsaco
make test     # 全部测试（见下文）
make proofs   # Lean 证明：proofs/check.sh（需要环境变量 CPP_LEAN、EDG_DIR，以及 lean）

内核编译过程：clang -x c -std=c11 --target=amdgcn-amd-amdhsa -mcpu=gfx1201 -nogpulib -O3 先生成 LLVM IR， 再编译成目标文件，最后 ld.lld -shared 链接成 hsaco。 clang 18 在 C 语言里不接受 amdgpu_flat_work_group_size 属性， 所以 Makefile 在 IR 中加入 "amdgpu-flat-work-group-size"="256,256"。 所有内核都用 256 线程的工作组。

运行

build/ie-run MODEL.gguf [--backend cpu|gpu] [--tokens 1,2,3] [--n 32] [--ctx N] \
             [--dump-logits FILE] [--hsaco build/ie_kernels.hsaco] [--info] \
             [--tokens-file FILE] [--ppl [--ppl-first N]] [--stop] \
             [--draft D [--mtp FILE]] [--chunk N] [--temp T [--top-k K] [--top-p P] [--seed S]]

• --tokens：提示词的 token id 列表。这一版没有 BPE 编码，只接受 id。
• --n：生成的 token 数。
• --stop：生成到结束符（tokenizer.ggml.eos_token_id）就停止。
• --chunk N：GPU 上提示按每块最多 N 个 token 处理（1 到 512，默认 256）。一块只读一次权重。
• --dflash 文件（与 --draft 7 一起用）：用 DFlash 2 起草器（另一个 GGUF 文件）起草，见下文“DFlash 2 起草器”。
• --draft D：用模型的 MTP 层做推测解码，每步最多 D 个草稿（1 到 7，只用于 GPU）。--mtp 文件：MTP 层在另一个文件里。 --draft-pmin P：MTP 层对某个草稿的概率低于 P 时，这一步不再继续猜（默认 0.6）。推荐 --draft 7。
• --temp T：采样温度（默认 0 = 贪心）。--top-k K：只在最大的 K 个里采样（0 = 不限制）。 --top-p P：只保留累计概率达到 P 的最大的那些 token（K = 0 且 P < 1 时，先取最大的 1000 个）。--seed S：随机数种子。
• --dump-logits：每一步写一行 位置 token logit0 logit1 ...。
• 输出：所有 token id，以及解码后的文本（gpt2 分词器用字节级解码，例如 Qwen2；llama 分词器处理 ▁ 和 <0xXX>）。
• 结束时打印每 token 时间、tokens/s，以及有效带宽（每个 token 读取的权重字节数 ÷ 时间）。

在 gfx1201 机器上运行（GPU）

1. 安装 ROCm 运行时，确保可以加载 libamdhip64.so（例如 export LD_LIBRARY_PATH=/opt/rocm/lib）。
2. 运行 make。GPU 主机上没有 clang 时，在另一台机器上运行 make build/ie_kernels.hsaco，再把这个文件复制过去。
3. 运行 build/ie-run model.gguf --backend gpu --tokens 1,2,3 --n 64。

如果找不到库，程序会退出并给出明确的提示。 GPU 后端只在 gfx1201 上编译。核函数需要 v_dot4_i32_iu8 指令和 wave32。

环境变量：

变量	作用
IE_PROFILE=1	每个核函数后记录一个 GPU 事件，结束时打印按核函数和按层的时间表
IE_PROFILE_CSV=文件	与 IE_PROFILE=1 一起用：把每个算子的时间写进 CSV 文件
IE_TRACE=文件	按启动顺序把每次核函数启动写进 CSV：运行序号、阶段（decode、prefill、verify、mtp、sample）、位置、算子、层、核函数、工作组数、权重字节数。不加 GPU 事件，不改变时间
IE_ATTN=old	用旧的注意力核函数 ie_attn（每个头一个工作组）
IE_ATTN=check	每层同时运行新旧两个注意力核函数，结束时打印两者输出的最大相对差
IE_ATTN=split	不用 ie_attn_kv，只用 ie_attn_split 和 ie_attn_split_t（结果逐位相同，用来对比）
IE_ATTN_KV_MIN=N	位置数大于 N 时用 ie_attn_kv（默认 768）
IE_ATTN_SPW=N	ie_attn_kv 每个工作组做 N 段（默认自动：约 128 个工作组）
IE_SPEC_TIME=1	MTP：打印每步草稿和验证的时间
IE_LOOKUP=0	MTP：关闭从上下文起草（默认开）
IE_DFLASH=0	加了 --dflash 也只用 MTP 层起草
IE_DFLASH_AUTO=0	每一步都用 DFlash 起草（默认：按测得的速度在 DFlash 和 MTP 之间自动选择）
IE_DFLASH_N=N / IE_DFLASH_PMIN=P	DFlash：每块最多 N 个草稿（默认 15）；草稿概率低于 P 时截断（默认 0.3）
IE_DFLASH_SEL=0	DFlash：不用选择器，每个位置取概率最高的候选（对比用）
IE_LOOKUP_AGREE=N	从上下文起草：MTP 草稿前 N 个和上下文一致时才用（默认 2）
IE_LOOKUP_LONG=N	从上下文起草：匹配不少于 N 个 token 时，MTP 链只跑 IE_LOOKUP_AGREE 步（默认 16）
IE_LOOKUP_MIN=N / IE_LOOKUP_MAX=N	从上下文起草：最短匹配（默认 2）；每步最多草稿数（默认 15，最大 15）
IE_ATTN_FA=0	预填充不用 ie_attn_fa，改用 ie_attn_pfg
IE_ATTN_FA_MIN=N	位置数大于 N 时预填充用 ie_attn_fa（默认 1024）
IE_ATTN_FA2=0	预填充不用 ie_attn_fa2（S 只算一次的版本），改用 ie_attn_fa
IE_GEMM_F32_S=0	预填充的小 f32 矩阵不用 ie_gemm_f32_s
IE_Q4K_SB=1	预填充的 Q4_K 矩阵乘法先把激活值重新量化成每 256 个元素一个缩放系数，再用整数累加（ie_gemm_q4kr_sb，快约 2%，默认关）
IE_GEMM_H=1	预填充的 Q4_K 矩阵乘法用 ie_gemm_q4k_h（激活值转 fp16，权重在 LDS 里解压成 fp16，fp16 矩阵指令）。默认关：比 ie_gemm_q4kr 慢 13% 到 36%（tools/tilelang/README.md）
IE_PROFILE=2	按预填充块统计每个算子的时间（配合 IE_STREAM=0）
--tokens-file 文件、--ppl、--ppl-first N（命令行选项）	从文件读 token；计算从第 N 个位置开始的困惑度
IE_GRAPH_BENCH=N	调试：把一个解码步骤录成 HIP 图，重放 N 次，和普通流比较每步时间，然后退出（27B Q4_K_M：30.87 对 30.42 毫秒）
IE_NORM_FUSE=1	把 RMSNorm 合进前一个带残差的矩阵向量乘（最后完成的工作组做）。在 R9700 上更慢，所以默认关闭（27B：48.2 对 46.0 毫秒/token）
IE_STREAM=0	关闭流式加载（先在主机上加载全部权重，再上传）
IE_TR_MIN=N	多 token 矩阵向量乘：token 数 ≥ N 时用每个 wave 算 4 行的核函数 ie_gemv_q8q8_tr（默认 2）
IE_GEMM_MIN=N	token 数 ≥ N 时用矩阵核心（WMMA）的矩阵乘 ie_gemm_q8 和预填充注意力 ie_attn_pf（默认 17）
IE_GEMM=old	预填充矩阵乘用旧的 ie_gemm_q8（权重经过 LDS），不用 ie_gemm_q8r
IE_GEMM=h	用 fp16 WMMA 的矩阵乘 ie_gemm_h（实验：比 int8 版慢）
IE_GEMM_CHECK=1	调试：每个 WMMA 矩阵乘也用矩阵向量乘算一遍，打印最大相对差（遇到带偏置或残差的矩阵时停止）
IE_SPEC_TRACE=1	推测解码时打印每一步的草稿、模型结果和接受数
IE_DRAFT_HEAD=q8	草稿用原来的 Q8_0 输出矩阵（默认：加载时做一份 Q4_0，只给草稿用）
IE_DRAFT_VOCAB=N	草稿头只含 token 0 到 N−1 和最后的特殊 token（默认 131072；0 = 全部 token）。验证不变，输出不变
IE_GEMV_NOBATCH=1	调试：多 token 运行时，矩阵向量乘每个 token 单独启动
IE_GDN_CHUNK=0	提示块不用分块并行的线性注意力（ie_gdn_prep/wy/seq），改用逐 token 的 ie_gdn
IE_GDN_CHECK=1	调试：每个提示块的线性注意力层先用 ie_gdn 算一遍，再用分块并行算法算，打印输出和状态的最大差

硬件追踪：用 rocprofv3 --kernel-trace 运行引擎，再用 tools/rocprof_ops.py 把追踪结果对应到算子（见脚本开头的说明）。 按核函数的时间：在 rocprofv3 下运行一次 ie-run（带 IE_TRACE），再用 tools/rocprof_kernels.py 日志.csv 追踪.csv [阶段] 汇总。 线程追踪：rocprofv3 --att（一次 kernel 调用）之后，用 tools/att_summary.py 按指令类别汇总延迟和停顿（用法见脚本开头）。 注意：硬件计数器（rocprofv3 --pmc）要先把显卡的性能等级设成 profile_standard： echo profile_standard | sudo tee /sys/class/drm/card1/device/power_dpm_force_performance_level，用完改回 auto。 在默认的 auto 下，SQ 的指令、等待、LDS 冲突和 L2 计数器读出来都是 0（只有 GRBM_GUI_ACTIVE、SQ_BUSY_CYCLES、SQ_WAVES 有值），线程追踪不受影响。 硬件计数器（2026-10-10，4K 提示，每块 512 个 token；周期的比例对 SQ_WAVE_CYCLES）：

kernel	向量指令周期	等待指令	等待（全部）	LDS 存储体冲突	L2 命中
ie_gemm_q4kr	10.2%	48.7%	42.9%	10.6%	98.8%
ie_gemm_q6kr	13.2%	39.7%	48.2%	33.8%	98.5%
ie_gemm_q8r	16.4%	32.3%	56.1%	9.2%	95.9%
ie_gdn_seq	13.6%	24.8%	59.0%	1.1%	89.1%
ie_attn_fa2	11.5%	8.5%	84.0%	28.2%	92.9%
ie_swiglu	2.9%	2.0%	94.2%	0	83.5%

核函数时间轴：tools/viz.sh 模型.gguf [步数] [输出.html] 在 rocprofv3 下运行一次 ie-run（带 IE_TRACE）， 再用 tools/ie_viz.py 生成一个独立的 HTML 页面和一个 Perfetto 追踪文件（.json，在 ui.perfetto.dev 打开）。 页面按运行（一个解码步骤、一个提示块……）显示每个核函数的时间段和核函数之间的空隙，泳道可以按类别、按层或单条； 颜色可以按类别或按带宽（权重字节 ÷ 时长，占标称 640 GB/s 的比例）。页面还给出每个运行的空隙时间、启动次数、 权重读取的下限时间，以及按核函数、按类别、按层的平均值。VIZ_ARGS="--draft 7" 可以看 MTP 的草稿和验证步骤。 矩阵向量乘基准测试：tools/gemv_bench.c 对不同形状单独计时（见文件开头的说明；KIND=q4k、q6k：K 系列格式）。 tools/gemv_eq.c 检查多 token 核函数与单 token 核函数逐位相同（随机 Q4_K、Q6_K 或 Q8_0 数据），并给两者计时。 tools/gemm_bench.c 每次启动换一份权重（共约 512 MB），权重从显存读，不从 64 MB 缓存读。 tools/gemv_rows_proto.c 是一个原型（每个 wave 算 R 行），测试结果：R = 2 和 R = 1 一样快，R = 4 和 8 更慢，所以引擎没有用它。

支持的模型：Llama 和 Qwen2 架构。Q4_0 和 Q8_0 矩阵走整数路径（dot4 指令）。F32、F16 和 BF16 矩阵在加载时转成 f32（BF16 精确：f32 的高 16 位）。 Q4_K 矩阵（Q4_K_M 文件的主体）保持 4 位（MAT_Q4K）：每 32 个权重一块，半字节按 Q4_0 的顺序重排（值 0..15）， 每块存 6 位缩放和 6 位最小值，每 256 个权重存 f16 的 d 和 dmin。块点积：dq = Σ q·a（整数，精确）， 然后 acc += da · ((d·sc)·dq − (dmin·mn)·asum)，和 ggml 的反量化公式相同。核函数：单 token ie_gemv_q4kq8， 2 到 16 个 token ie_gemv_q4kq8_tr（与单 token 逐位相同），17 个以上 ie_gemm_q4kr（WMMA）。 Q6_K 矩阵保持 6 位（MAT_Q6K）：低 4 位按 Q4_0 布局存放，高 2 位每 32 个权重另存 2 个字， 每 16 个权重一个 int8 缩放，每 256 个权重一个 f16 的 d。块点积：q − 32 由字节运算得到有符号 int8， s0·dq0 + s1·dq1（两个 16 权重的半块），浮点写成同样的显式 fma。核函数：ie_gemv_q6kq8、ie_gemv_q6kq8_tr、 ie_gemv_q6kq8_tr8、ie_gemm_q6kr（WMMA，两个半块各一次 WMMA）。设置 IE_Q6K_AS_Q8=1 时仍按 Q8_0 加载。 Q5_K 和 Q4_K 的词嵌入在加载时再量化成 Q8_0（近似）。Q4_K 和 Q6_K 的整数部分（地址、半字节和高位重排、子块点积、预解包的点积） 由 Lean 证明（11 条定律，见“什么被验证”）。qwen35 的 Q4_0 矩阵（Q4_K_M 文件里的 MTP 层）按 Q8_0 加载（精确：int8 = q − 8，同一缩放）。 已确认：make test 的 tiny-llama-k（Q4_K、Q5_K、Q6_K）与 numpy 参考（gguf 包自己的反量化）相对差 3e-7。

Qwen3.8-27B-Coder390-EfficientThink Q6_K-MTP（nerkyor 的微调版，23.2 GB，2026-10-10）：结构与 Qwen3.8-27B 相同； 矩阵几乎全是 Q6_K（原生读取），输出层和词嵌入是 Q8_0，ssm_alpha / ssm_beta 是 BF16（加载时精确转成 f32），MTP 层 Q6_K。

项目	Coder390 Q6_K	官方 Q4_K_M
每个 token 读的权重	21.4 GB	16.5 GB
困惑度（wikitext-2，4 段）	5.084	5.173
解码，不用草稿（个/秒）	24.0（514 GB/s）	33.0
写代码，7 个草稿	108.5	112.7
续写，7 个草稿	63.2	89.8
聊天，7 个草稿	71.0	68.3
中文，7 个草稿	41.5	48.9
预填充 512 个 token	1222	1439

四个提示、3 和 7 个草稿，输出都与不用草稿时逐个 token 相同。困惑度的差别同时来自微调和量化（Q6_K 比 Q4_K_M 精度高），无法分开。 同一仓库的 Q8_0-MTP 也能加载（全是 Q8_0 加 BF16）；Q4LynnStyle 含 Q5_K（转成 Q8_0，近似）；Q3LynnStyle（Q3_K）和 Q2LynnStyle（IQ 系列）不支持。 注意：测量前 GPU 曾处于低功耗状态（power_dpm_force_performance_level 不是 auto），所有核函数慢约 3.5 倍。

Qwen3.8-27B Q4_K_M（bartowski，17.4 GB；主体 Q4_K，另有 Q6_K、Q8_0、Q5_K，自带 MTP 层 blk.64）。2026-10-08，同一台 R9700：

项目	本引擎 Q8_0	本引擎 Q4_K_M	llama.cpp Q4_K_M
困惑度（wikitext-2，4 段 × 512，后半计分）	5.1852	5.173	5.1697
解码，不用草稿（个/秒）	21.9	33.0	30.4
写代码，7 个草稿	93.6	112.7	-
续写（prose 提示），7 个草稿	56.4	89.8	-
聊天，3 个草稿	51.2	69.8	-
预填充 512 个 token（每块 512）	1785	1439	1134（pp512）

三个提示、3 和 7 个草稿，输出都与不用草稿时逐个 token 相同。MTP 层用文件自带的（--mtp 指向同一个文件）。 改动：ssm_alpha / ssm_beta 是 F32 矩阵，多 token 时以前每个 token 启动一次（预填充只有 336 个/秒）， 现在 ie_gemv_f32_t 一次做完，F32 读取一次发出 8 个；验证（≤ 8 个 token）用每个 wave 8 行的 ie_gemv_q4kq8_tr8 （激活值的读取减半：8 个 token 时 435 → 约 500 GB/s）。Q4_K 的浮点运算在所有核函数里都写成同样的显式 fma， 所以单 token 与多 token 仍逐位相同。 Q6_K 原生读取（以前转成 Q8_0）：每个 token 少读约 1.3 GB，解码 30.1 → 32.3 个/秒（533 GB/s，标称值的 83%）， 困惑度 5.1745 → 5.173（不再有再量化的误差）。 预填充 1433 → 1451，两处改动：

1. ie_gemm_q4kr 47 → 51 T MAC/s。最小值项 −Σ (dmin·mn)·(da·asum) 是一个 K = 块数的小矩阵乘： 每 4 块用一次 fp16 WMMA（v_wmma_f32_16x16x16_f16）直接加进 f32 累加器，每块每个输出元素少一条 fma。 fp16 的相对误差 2^-11，远小于 4 位权重和 8 位激活值本身的舍入误差。
2. ie_gemm_q6kr 36 → 40 T MAC/s：两个半块写成 acc += da·(s0·d0 + s1·d1)，全部是可以双发射（VOPD）的浮点指令。 试过的另外两种写法更慢：整数合并（v_mad_i32_i24 不能双发射，30 T MAC/s）；把魔数减法预先算好（40，但多了相消误差）。 MTP 的速度变化（写代码 107.3 → 106.2，聊天 65.6 → 70.0）来自接受率：生成的文本变了；每步时间没有变。

解码和验证的改动（2026-10-08；两个模型、三个提示、3 和 7 个草稿，输出仍与不用草稿时逐个 token 相同）：

改动	效果
ie_rmsnorm：平方和一次发出 8 个读取（求和顺序不变）	每个 token 1.36 → 1.10 毫秒
ie_gemv_f32 每个工作组一行（以前每个 wave 一行：48 行只有 6 个工作组）；每层 ssm_beta 和 ssm_alpha 合成一个矩阵	每个 token 1.17 → 0.54 毫秒；Q4_K_M 解码 32.3 → 33.0 个/秒
验证（6 到 8 个 token）：ie_gemv_q4kq8_ts2、ie_gemv_q6kq8_ts2，8 个 token 分给 2 个 wave（每个 wave 8 行 × 4 个 token），工作组数翻倍	验证的矩阵向量乘快 5% 到 24%（行数少的矩阵快得多）
验证的注意力一次启动（ie_attn_split_t，网格 y = token；每个 token 的分段、运算和合并与 ie_attn_split 相同）	每步少约 110 次启动

合计：Q4_K_M 写代码（7 个草稿）106.2 → 108.8 个/秒，续写 88.4 → 89.7。

MTP 层保持 Q4_0（2026-10-08）：Q4_K_M 文件自带的 MTP 层（blk.64）是 Q4_0。以前按 Q8_0 加载（Q8_0 有多 token 核函数）， 但 MTP 层每次只算 1 个 token，所以现在保持 Q4_0（MAT_Q4，ie_gemv_q4q8）；只有 ffn_down（K = 17408，超出 Q4 的大小限制）仍转成 Q8_0。 MTP 层 451 → 283 MB。写代码（7 个草稿）108.8 → 112.7 个/秒，聊天（3 个草稿）69.3 → 69.8，续写不变。 MTP 层只决定草稿，不决定输出：三个提示、3 和 7 个草稿，输出仍与不用草稿时逐个 token 相同。 草稿的接受率略有变化（Q4_0 核函数的求和顺序不同，草稿概率有极小差别）。IE_MTP_Q8=1 恢复以前的 Q8_0 加载。

草稿词表和 pmin 的测量（Q4_K_M，7 个草稿，个/秒；中文提示：用三段话介绍上海的历史）：

IE_DRAFT_VOCAB	写代码	续写	聊天	中文
32768	98.2	83.5	63.7	-
65536	113.7	88.9	67.5	-
98304	115.1	91.2	69.5	41.2
131072（默认）	112.8	89.9	68.4	48.9
248320（全部）	105.7	88.3	64.7	47.8

中文输出的 token 有一半 ID ≥ 98304（331 个里 160 个），所以默认值保持 131072。 中文不用草稿 32.9 个/秒，用 MTP 每步只有约 1.8 个 token：pmin 0.3 到 0.75 都在 46.6 到 49.6 之间，默认 0.6 不变。 推测：MTP 层对中文猜得不准，这不是参数能解决的。

采样时的前 K 个在 GPU 上选（2026-10-08）：ie_topk_part（每行 64 个工作组，各选自己那段的前 K 个）和 ie_topk_merge（合并）， 主机只取回 K 对（值，ID），K ≤ 64（--top-k）。候选的顺序固定为“值从大到小，值相同时 ID 小的在前”，主机路径也按这个顺序， 所以 GPU 和主机选出的结果相同（已确认：同一种子，输出逐个 token 相同；MTP 和不用草稿也相同）。 每步的采样时间（Q4_K_M，写代码，7 个草稿，温度 0.7）：top-k 20 + top-p 0.95：1.36 → 0.34 毫秒；top-k 40：1.34 → 0.50 毫秒。 不用草稿时每个 token 0.35 → 0.18 毫秒。只用 --top-p（前 1000 个）时仍在主机上选（每步约 3.4 毫秒）。IE_HOST_TOPK=1：总在主机上选。 预填充时 F32 矩阵仍用每个 wave 一行的 ie_gemv_f32_tw（每行一个工作组在 512 个 token 时更慢：1451 → 1417）。

试过但没有用的：

1. HIP 图（hipGraph）：每步只省 0.45 毫秒（1.5%）。核函数之间约 3.6 微秒的空隙主要在 GPU 端，不在主机端。
2. ie_gemv_q4kq8 的 d、dmin 合成一次读取，一步读 4 个块：都没有变化（大矩阵已到约 600 GB/s，标称值的 94%）。
3. 验证核函数预取下一个块：寄存器 208 → 254，反而变慢。
4. 用 WMMA 做验证的整数点积，再按单 token 核函数的顺序重放浮点累加（32 个虚拟 lane 分给 8 个 wave， 最后按 wave_tree_f 的比特反转二叉树相加）：tools/gemv_eq.c 确认逐位相同，但只在两个形状上快 5% 到 10%， 大矩阵上更慢（每条读取指令访问 16 个缓存行），所以没有采用。
5. Q8_0 验证的同样拆分：基准测试更快，真实运行没有变化，所以没有采用。
6. 验证矩阵向量乘的进一步尝试（都保持逐位相同，tools/gemv_eq.c 确认；8 个 token，Qwen3.8-27B 的 6 个 Q4_K 形状）：
   • LDS 版 WMMA（wl）加深预取到 2 轮：没有变化。去掉全部计算后只剩读取，也只有 400 到 510 GB/s：瓶颈在读取方式本身。
   • ts2 去掉点积和浮点运算：几乎不变；去掉激活值读取：快 5% 到 10%；只算 1 个 token：仍比单 token 核函数慢约 20%。 所以 8 行一个 wave 的结构本身就慢于单 token 核函数（每个 wave 一行、寄存器少、常驻 wave 多）。
   • 权重保持打包、每个 token 再解包，2 到 4 个块同时在途；每个 wave 1、2、4、8 行；token 分给 1、2、4 个 wave； 权重用普通读取代替非时间性读取；d 和 dmin 一次读取：共约 14 种组合，都在 400 到 500 GB/s，没有一种稳定超过 ts2。
   • 注意：同一个核函数在不同次测量之间相差约 ±10%，小于 10% 的差别不可信。 IE_Q6K_AS_Q8=1 时 Q6_K 先还原成浮点数再按 Q8_0 重新量化（近似：Q8_0 每 32 个权重一个缩放，Q6_K 每 16 个一个； Qwen2.5-1.5B 的输出层上，权重的均方根误差约为权重均方根的 0.5%）。 词嵌入如果是 Q8_0，也转成 f32。 验证核心限制矩阵大小：Q4_0 行数 < 2^18，每行块数 < 512（K ≤ 16352）；Q8_0 每行块数 < 256（K ≤ 8160）。 Q4_0 超出限制时程序退出；Q8_0 超出限制时改走 f32 路径。

设计

一步解码 = 一个固定的算子列表。两个后端执行同一个列表，并使用同一组 arena 偏移。 加载时，q、k、v 三个矩阵拼成一个矩阵，gate 和 up 拼成一个矩阵（种类不同时不拼）。 重排布局按行存放，所以拼接只是把行接在后面，索引映射仍然是验证核心的映射。 一个缓冲区可以放几个向量（例如 q、k、v），算子用字节偏移指向其中一个。每层有 9 个算子：

1. RMSNorm，同时写出 Q8 副本。
2. qkv 矩阵向量乘，加上 bias。
3. RoPE + 写 KV 缓存（一个算子）。
4. 注意力，同时写出 Q8 副本（头的大小是 32 的倍数且不超过 256 时）。
5. 输出投影的矩阵向量乘，加上残差。
6. RMSNorm，同时写出 Q8 副本。
7. gate+up 矩阵向量乘。
8. SwiGLU，同时写出 Q8 副本。
9. down 矩阵向量乘，加上残差。

层外还有词嵌入、最后的 RMSNorm、输出层和取最大值。Qwen2.5-0.5B（24 层）每个 token 共 220 个算子。

Q4_0 和 Q8_0 重排：加载时，每个矩阵被重排成 GPU 布局：Q4_0 每块 4 个 u32 字，Q8_0 每块 8 个 u32 字，加一个独立的 f16 缩放数组。 重排只用核心函数（ie_q4_src_*、ie_q4_dst_*、ie_q8_src_*、ie_q8_dst_*、ie_pack4）。

GEMV：网格是 ie_gemv_ngroups(rows) 个工作组（落在 253 到 256 之间时补到 257 个，见下）；wave v 计算行 ie_gemv_row(g, v)； lane l 计算块 ie_lane_blk(l, t)（l, l+32, ...）；块点积是 ie_q4q8_block 或 ie_q8q8_block（各 8 条 dot4 指令）； 每个 lane 一次发出 2 个块的读取，再计算；最后按 ie_tree 的顺序做 wave 归约。 CPU 后端按完全相同的划分和顺序计算，所以两者只差浮点舍入。

已确认的硬件现象：在 R9700 上，矩阵向量乘的工作组数在 253 到 256 之间时，比 252 或 257 个慢 2 到 3 倍 （tools/gemv_bench.c：2048 行 × 48 块，256 个工作组 17.0 微秒，257 个 6.5 微秒）。原因未知。 引擎因此在这个区间补到 257 个工作组，多出的工作组没有行，直接结束。

注意力（ie_attn_split）：工作组 = 查询头 × 位置段（每段 32 个位置）。 每个 wave 对自己的位置做在线 softmax；8 个 wave 在 LDS 中合并；最后完成的工作组（原子计数器）合并各段的结果： 整个工作组并行求各段的最大值和总和，每段的权重只算一次并放在 LDS 中。 （每段 8 或 16 个位置的版本试过：生成 64 个 token 时没有更快，生成 512 个时更慢。）

长上下文的注意力（ie_attn_kv，2026-10-09）：ie_attn_split 里，一个 KV 头被 6 个查询头各读一遍；验证时 T 个 token 再各读一遍。 ie_attn_kv 的工作组 = 一个 KV 头 × 若干段 × 8 个（查询头，token）组合：K 和 V 每次 8 个位置读进 LDS（16 字节读取， 计算这 8 个位置时下一批的读取已经发出），每个 wave 算一个组合。为了和 ie_attn_split 逐位相同，一个 wave 保留 8 路在线 softmax， 第 x 路正好是 ie_attn_split 第 x 个 wave 的位置和运算顺序，最后按原来的 LDS 合并顺序合并，写出同样的分段结果； 各段的最后合并也和原来相同（同样的求和顺序，读取每次发出 8 个）。一个工作组连续做若干段，内存栅栏和原子计数每个工作组只做一次。 位置数大于 768 时使用（更短时 ie_attn_split 更快）。

测量（27B Q4_K_M，R9700）：

注意力（16 层，毫秒/token）	100 个位置	900	7300
ie_attn_split	0.39	1.09	7.43
ie_attn_kv	0.67	1.09	3.44

7300 个位置的上下文	以前	现在
不用草稿	37.7 毫秒/token	34.3 毫秒/token
3 个草稿：验证每步	61.4 毫秒	44.9 毫秒
3 个草稿：速度	48.3 个/秒	65.6 个/秒

短上下文（100 个位置）不变：不用草稿 30.4 毫秒/token，3 个草稿 83.9 个/秒。 检查：两个提示（写代码、7300 个 token）× 四种设置（不用草稿、7 个草稿、7 个草稿加采样、3 个草稿），ie_attn_kv 和 IE_ATTN=split 的输出 token 逐个相同， 也和这次修改之前的输出相同。

预填充的注意力（ie_attn_fa，2026-10-09）：fp16 WMMA 的闪电注意力。工作组 = 一个 KV 头 × 输出的 128 列 × 16 个 token； 6 个 wave 各算一个查询头。每块 32 个位置的 K（全部 256 维）和 V（这 128 列，转置）转成 fp16 读进 LDS， 6 个查询头和 16 个 token 共用；S = Q·Kᵀ 和 O += P·V 用 v_wmma_f32_16x16x16_f16（f32 累加），在线 softmax 用 f32， P 经 LDS 从 D 布局换到 A 布局。输出分成两半（每个 wave 放得下 Q 的 64 个寄存器和半个 O 的 64 个寄存器）， 所以 S 每半算一次。融合的 Q8 拷贝改成单独的 ie_attn_fa_q8（一个头的两半在不同工作组）。 和 ie_attn_pfg 一样不与解码逐位相同（Q、K、V、P 舍入到 fp16）；MTP 与不用草稿时仍逐位相同（两者用同一个预填充）。

预填充（27B Q4_K_M，块 512）	4K	8K	16K	32K	63K
ie_attn_pfg（以前），个/秒	1,151	950	689	452	273
ie_attn_fa，个/秒	1,347	1,293	1,174	980	775

时间拟合 T(n) = a·n + b·n²：b 从 4.7e-5 降到 9.0e-6 毫秒；63K 时注意力占 44%（以前 81%）。 精度：真实文本（README、源代码）的困惑度，8K：4.4846 对 4.4876；32K：2.7656 对 2.7660。 计时实验（16K）：S 只算一半、不读 K/V、不算 P·V，各只省 0 到 0.6 秒，所以剩下的时间主要在每块的同步、softmax 的跨 lane 归约和延迟上。

线程追踪找到的问题（2026-10-10，rocprofv3 --att）：ie_gemm_q4kr 的预取在发出下一块的读取之后立刻把 asum 转成浮点， s_wait_loadcnt 0x0 等完了刚发出的全部读取，预取失效（这一条指令占这个 kernel 15.8% 的周期，等显存共 19.3%）。 改成到写 LDS 时再转换：等显存 19.3% → 5.0%，这个 kernel 的周期少 8.4%，结果逐位相同。 预填充 8K 6.26 → 5.88 秒，32K 29.4 → 27.3 秒（容器里的 ROCm 7.14，交替测三次）。修复后三个 GEMM 的周期分布：

kernel	向量指令	屏障	矩阵指令	等显存	等 LDS
ie_gemm_q4kr	54.6%	22.5%	11.2%	5.0%	2.3%
ie_gemm_q6kr	56.6%	20.8%	8.7%	1.3%	7.5%
ie_gemm_q8r	45.3%	17.7%	16.5%	12.4%（循环之前）	3.7%

ie_attn_fa2（2026-10-09，hd = 256、每个 KV 头最多 6 个查询头时代替 ie_attn_fa；IE_ATTN_FA2=0 关闭）： 工作组 = 一个 KV 头 × 全部 256 列 × 16 个 token，所以每块的 S 只算一次（以前两半各算一次，Q·Kᵀ 做两遍）。 阶段 A：wave 0 到 5 各算一个头的 S 和在线 softmax，把 P 和每行的缩放系数 c 写进 LDS。 阶段 B：6 × 16 = 96 个输出块（头 × 16 列）分给全部 8 个 wave：算 S 的 wave 各 10 块（它们还放着 Q），另外 2 个 wave 各 18 块。 两种 wave 走各自的一份循环代码（屏障相同），这样寄存器分开分配。 LDS 地址在循环里用不透明的 lane 号重新计算，否则编译器把 10 个块的地址一直放在寄存器里，溢出 40 个；现在循环里没有溢出。 每个输出元素的运算和 ie_attn_fa 相同，所以结果逐位相同（困惑度和生成的 token 都相同）。

第二步（同一天）：每块先发出全部 K/V 读取（每个线程 16 个 8 字节读取），再统一写 LDS；K 直接按 fp16 用 8 字节写入， V 按 4 个 key × 4 列在寄存器里转置后用 8 字节写入（以前每个元素一次 2 字节写入）；阶段 B 只在换头时读 P 和缩放系数； 输出块改成 8 个 wave 各 12 块。

计时实验（61K，第一版 ie_attn_fa2，注意力共 24.6 秒）：不读 K/V 省 14.5 秒，不做阶段 A 省 9.9 秒，不做阶段 B 省 14.3 秒。 读取慢的原因是读取循环不展开，每块串行等 10 次显存延迟。没有用的尝试： 把下一块读进寄存器（预取）会让算 S 的 wave 溢出 68 个寄存器；P 放进 K 的 LDS（LDS 从 48 KB 降到 37 KB，每个 WGP 放 3 个工作组）速度不变。

预填充（27B Q4_K_M，fp16 KV，块 512），秒	32K	61K	127K
ie_attn_fa	32.4	77.4	234.2
ie_attn_fa2 第一版	30.2	67.0	191.8
ie_attn_fa2 现在	27.5	59.9	158.6
其中注意力（减去 0.692 毫秒 × n）	9.7 → 4.8	34.9 → 17.4	146.3 → 70.7

结果仍与 ie_attn_fa 逐位相同：32K 困惑度 5.3272；4K 提示后 200 个 token 的输出相同，3 个草稿时也相同。

RMSNorm：工作组数 = 维度 ÷ 256。每个工作组都计算完整的平方和（代码相同，所以值相同）， 然后每个 wave 只归一化并量化一个 32 元素块。

wave 内归约：用 DPP 指令（v_permlanex16 和 row_shl），不用 LDS 的 ds_bpermute。 矩阵向量乘的求和顺序不变（先加 16 之外的 lane，再 8、4、2、1），lane 0 的结果和以前逐位相同。

取最大值：128 个工作组各自找最大值，最后完成的工作组合并。

激活规划：每个缓冲区有大小和生存期。贪心规划器选最低的 256 字节对齐偏移。 然后引擎填好 Plan 结构，运行 ie_plan_check；检查失败就退出。

Qwen3.5 / Qwen3.8（GGUF 架构名 qwen35）：每 4 层中，前 3 层是线性注意力层，第 4 层是全注意力层。 多 token 预测层（nextn_predict_layers）不用于逐个 token 的解码，加载时跳过。

• 线性注意力层（Gated DeltaNet，门控增量规则）：attn_qkv、attn_gate（z）、ssm_beta、ssm_alpha 拼成一个矩阵， 一次矩阵向量乘。然后一个核函数 ie_gdn 做完这一层的其余部分。每个值头一个工作组。步骤：

  1. 4 抽头因果卷积，再 SiLU。卷积状态是 4 个槽的环形缓冲区，第 pos 个输入写到槽 pos % 4。本步只读其他 3 个槽。
  2. q、k 每头做 L2 归一化：x / sqrt(Σx² + eps)。值头 h 用键头 h % 键头数。
  3. beta = sigmoid(b)，衰减 = exp(softplus(a + dt_bias) · A)。
  4. 状态 S（128 × 128，每个值头一个）：S ← 衰减 · S；delta = (v − Sᵀk) · beta；S ← S + k deltaᵀ；o = Sᵀq / sqrt(128)。
  5. 门控 RMSNorm：rmsnorm(o) · 权重 · silu(z)，同时写出 Q8 副本。

  位置 0 时，核函数不读旧状态（当作 0），所以新序列不需要清零状态。

• 全注意力层：q 投影每头输出 [q | gate]。ie_qkn_rope_kv 对 q、k 每头做 RMSNorm，只旋转前 64 维（NEOX 配对）， 写 KV 缓存。注意力输出乘以 sigmoid(gate)（在注意力核函数的最后一步做）。KV 缓存只给全注意力层分配。

• 词嵌入：Q8_0 词嵌入保持 Q8_0（27B 模型 1.3 GB，转成 f32 会是 5 GB），用 ie_embed_q8 取一行。

• 流式加载：GPU 后端在加载时，每个矩阵重排后立刻复制到显存，然后释放主机内存（ie_mat_sink）。 主机内存只需要放下一层。IE_STREAM=0 关闭。

• 验证核心的 Q8 尺寸：每行最多 2047 块，并且行数 × 块数 < 2^26（27B 的 down 投影每行 544 块）。

测试结果

本机（make test，CPU）

• laws_test：所有定律在随机输入上通过。
• kernel_test：Q4_0 和 Q8_0 各 9 种形状。CPU GEMV 与逐元素 double 点积的最大相对误差：Q4_0 约 8e-8，Q8_0 约 5e-8。重排后的权重与 GGUF 一致。
• 两个小模型与 numpy 参考比较（模拟 Q8 激活量化）：相对误差最大约 5e-7。纯浮点模式约 2e-2 到 3e-2（Q8 激活量化本身的误差）。贪心 token 都相同。
• objdump：ie_gemv_q4q8 和 ie_gemv_q8q8 各有 8 条 v_dot4_i32_iu8（前者 A 无符号，后者 A 有符号），没有 scratch 内存。

GPU 主机（Radeon AI PRO R9700，gfx1201，ROCm 7.1 / 7.14 运行时）

提示 20 个 token，正常运行（不计时模式），每个 token 的平均时间：

模型	生成 64 个	生成 512 个	生成 2048 个	每个 token 读取的权重
Qwen2.5-0.5B-Instruct Q4_0（输出层 Q8_0）	1.40 毫秒（714 个/秒）	1.47 毫秒	1.76 毫秒	346 MB
Qwen2.5-1.5B-Instruct Q4_0（输出层 Q6_K → Q8_0）	2.63 毫秒（380 个/秒）	2.71 毫秒	3.14 毫秒	985 MB

1.5B 生成 64 个 token 时，有效带宽 374 GB/s（标称 640 GB/s 的 58%）。 硬件追踪中，大矩阵的核函数接近标称带宽：gate+up（17920 × 1536）585 GB/s，down（1536 × 8960）540 GB/s， 输出层（151936 × 1536，Q8_0）629 GB/s。剩下的时间主要是注意力、几个小矩阵（qkv、输出投影）和核函数之间的间隔。

与 llama.cpp 对比（同一块卡，同样的 GGUF 文件；llama.cpp b11222，Vulkan 后端，RADV 驱动，llama-bench -p 0 -n 64,512 -r 5； 本引擎用 1 个 token 的提示，各跑 5 次）。单位：每秒生成的 token 数。

模型	测试	llama.cpp（Vulkan）	本引擎	倍数
Qwen2.5-0.5B Q4_0	生成 64 个	465 ± 36	716	1.54
Qwen2.5-0.5B Q4_0	生成 512 个	533 ± 28	683	1.28
Qwen2.5-1.5B Q4_0	生成 64 个	280 ± 8	381	1.36
Qwen2.5-1.5B Q4_0	生成 512 个	328 ± 7	371	1.13

注意：llama.cpp 这里用的是 Vulkan 后端，不是 ROCm（HIP）后端，没有测 ROCm 版本。1.5B 的输出层，llama.cpp 直接读 Q6_K， 本引擎读重新量化后的 Q8_0（字节更多，而且是近似）。

Qwen2.5-3B-Instruct Q4_0（输出层 Q6_K → Q8_0，每个 token 读 1891 MB）：生成 64 个 token 226 个/秒（llama.cpp 196 ± 8）， 生成 512 个 219 个/秒（llama.cpp 200.1 ± 0.4）；有效带宽 429 GB/s（标称值的 67%）。

输出质量（困惑度）：wikitext-2 测试集，用 llama.cpp 的分词器分词，前 4 段，每段 512 个 token，只给每段后半的 255 个 token 计分 （和 llama-perplexity -c 512 --chunks 4 相同）。本引擎用 ie-run --tokens-file 段.txt --ppl --ppl-first 256 逐段计算。

模型	llama.cpp（Vulkan）	本引擎（GPU）	差别
Qwen2.5-0.5B Q4_0	15.062	15.125	+0.42%
Qwen2.5-1.5B Q4_0	10.144	10.146	+0.02%
Qwen2.5-3B Q4_0	9.069	9.092	+0.25%

所以 CPU、GPU 和 numpy 参考之间约 2e-2 的 logits 相对差，没有让输出质量变得比 llama.cpp 差。

• 回答 "What is the capital of France? Answer in one sentence." 时，两个模型都输出 "The capital of France is Paris."。
• 两个小模型：GPU 与 CPU 的 logits 相对误差最大 7e-3，中位数约 1e-7 到 3e-7，贪心 token 相同。
• IE_ATTN=check：新旧注意力核函数在 48,456 次调用（位置到 2020）中，输出最大相对差 5.5e-5。
• 已知问题：Qwen2.5-0.5B 上，CPU、GPU 和 numpy 参考两两之间的 logits 相对误差约 2e-2（中位数）。 注意力输出只差 1e-7 到 5e-5，logits 却差 2e-2，说明模型把微小的舍入差放大了。推测原因是 Q8 激活量化遇到离群值。这个问题还没有解决。

Qwen3.5 / 3.8（qwen35 架构）

Qwen3.5-0.8B Q8_0（ggml-org/Qwen3.5-0.8B-GGUF，24 层，词嵌入与输出层共用）：

项目	llama.cpp b11222（Vulkan）	本引擎（GPU）
困惑度（wikitext-2，4 × 512）	15.795	15.826（+0.2%）
生成 64 个 token	233.7 ± 6.6 个/秒	415 个/秒
生成 512 个 token	305.2 ± 13.4 个/秒	415 个/秒

• CPU 后端第 0 段困惑度 9.8839，GPU 9.8883。
• 硬件追踪：ie_gdn 每次 6.9 微秒（先发出状态读取，与卷积重叠；以前 20 微秒）。
• 聊天模板（<|im_start|>、空的 <think> 块）正常；--stop 在结束符处停止。
• 27B 的矩阵形状上（gemv_bench，Q8=1），Q8_0 矩阵向量乘达到 589 到 634 GB/s（标称值的 92% 到 99%）。

Qwen3.8-27B Q8_0（ggml-org/Qwen3.8-27B-GGUF，28.6 GB，64 层：48 个线性注意力层、16 个全注意力层）：

项目	llama.cpp b11222（Vulkan）	本引擎（GPU）
困惑度（wikitext-2，4 × 512）	5.1842	5.1804（−0.07%）
生成 64 个 token	20.03 ± 0.03 个/秒	21.7 个/秒（46.01 毫秒/个）
生成 512 个 token	20.09 ± 0.02 个/秒	21.7 个/秒（46.17 毫秒/个）

• 每个 token 读 27.2 GB 权重，有效带宽 590 GB/s（标称值的 92%）。本引擎各跑 3 次，差别小于 0.1%。
• 每段困惑度（本引擎 / llama.cpp）：4.332 / 4.324，7.752 / 7.765（由累计值推算），4.051，5.294。
• 流式加载：516 个数组（28.57 GB）在加载时复制到显存，加载用 50 秒。显存还剩约 5 GB。
• 两个模型的分词结果相同（wikitext-2 全文的 token 序列逐个相同）。
• 聊天模板下，"Explain in two sentences why the sky is blue." 的回答正确（瑞利散射），--stop 在 <|im_end|> 处停止。

MTP 推测解码（Qwen3.5 / 3.8）

MTP（多 token 预测）：模型自带的一个预测层（GGUF 里的 blk.N.nextn.*）。它先猜后面 D 个 token（"草稿"）， 模型再一次验证这 D 个草稿。用法：--draft D（D = 1 到 3，只用于 GPU）；--mtp 文件：MTP 层在另一个文件里。

一步的做法（src/main.c，spec_generate）：

1. MTP 层补算上一步接受的位置（用模型真实的隐藏状态），再从 (h(P−1), t(P)) 猜出 d1。
2. 用 MTP 层自己的输出 h 和 d1 猜 d2，依此类推到 dD。
3. 模型一次处理 [t(P), d1, …, dD] 这 D + 1 个 token，权重只读一次（ie_gemv_q8q8_t）。
4. 从 d1 开始，草稿等于模型的贪心结果就接受，第一个不等就停止。所以输出和普通贪心解码逐个 token 相同。

为了能退回到任一个被接受的 token：

• 线性注意力状态有 4 个槽。第 t 个 token 之后的状态写到槽 (c + t) % 4。
• 卷积输入环形缓冲区有 8 个槽。
• KV 缓存按位置覆盖，不需要退回。

已确认：多 token 矩阵向量乘的每个 token 与单 token 核函数的算术和顺序相同，结果逐位相同。

27B 的 MTP 层：ggml-org 的 Q8_0 文件没有 MTP 层。unsloth 的 Q8_0 文件有，而且 MTP 张量都在文件末尾。 所以只用 HTTP 分段请求下载了最后 451 MB，再做成一个只含 MTP 层的 GGUF。

结果（Qwen3.8-27B Q8_0，R9700，贪心，每次生成 300 个 token）。单位：每秒 token 数。

提示	草稿数	本引擎	本引擎接受率	llama.cpp b11222（Vulkan，--spec-type draft-mtp）
写代码	0	21.7	-	19.6
写代码	1	38.4	98.7%	33.0
写代码	2	51.3	97.1%	44.2
写代码	3	60.1	93.2%	53.2
续写 wikitext	0	21.7	-	19.4
续写 wikitext	1	36.0	86.9%	30.3
续写 wikitext	2	45.6	81.6%	37.2
续写 wikitext	3	50.6	74.2%	40.0
聊天（天空为什么是蓝的）	3	47.5	67.3%	-

• 所有本引擎的结果都和不用草稿时的输出逐个 token 相同（0.8B 和 27B，草稿数 1 到 3）。

• 更新：最多 7 个草稿，按 MTP 的概率停止。 ie_argmax 同时算出最大值 token 的 softmax 概率（同一遍里求 log-sum-exp）。 草稿的概率低于 --draft-pmin 时，这一步不再继续猜；验证的 token 数随之变化（1 到 8）。线性注意力状态改为 8 份。 27B，300 个 token，贪心（单位：个/秒；所有结果都和不用草稿时逐个 token 相同）：

  提示	3 个草稿	5 个，pmin 0.6	7 个，pmin 0	7 个，pmin 0.6
  写代码	62.2	74.5	77.2	78.3（每步 5.98 个 token）
  续写 wikitext	54.9	59.0	57.5	63.0
  聊天	49.8	49.5	45.9	50.5

  不按概率停止时，草稿越多，接受率越低（聊天 7 个草稿只有 38%），多猜的草稿浪费时间；pmin 0.6 时聊天的接受率是 62%。

• 草稿用 4 位输出矩阵。 每个草稿要读 MTP 层（0.45 GB）和整个输出矩阵（1.35 GB）。加载时把输出矩阵再量化成一份 Q4_0 （0.72 GB，按 ggml 的 Q4_0 规则），只给草稿用；验证仍用原来的 Q8_0，所以输出不变。MTP 补算一批位置时， 也只算最后一个位置的输出（只有它的草稿有用）。IE_DRAFT_HEAD=q8 关闭。27B，7 个草稿，pmin 0.6：

  提示	Q8_0 草稿头	Q4_0 草稿头	接受率（Q8 / Q4）
  写代码	78.5	85.5	84.6% / 84.8%
  续写 wikitext	62.9	66.0	93.9% / 95.0%
  聊天	50.4	53.2	61.6% / 60.8%

• 更新（预填充用 WMMA 之后；验证改用 ie_gemv_q8q8_tr，逐位不变）：写代码 3 个草稿 62.2 个/秒（接受率 92.5%）， 2 个草稿 52.8。续写 wikitext 和聊天的生成文本变了（预填充的数值变了，后面的贪心续写随之不同）， 所以接受率也变了：续写 3 个草稿 46.6 个/秒（61.3%），聊天 49.3 个/秒（65.3%）。仍然和不用草稿时逐个 token 相同。

• 验证步骤和草稿头的读取（2026-10-08）。硬件追踪：7 个草稿时每步约 74 毫秒，其中验证的矩阵向量乘 49.5 毫秒， 草稿（MTP 层 + 草稿头）14.8 毫秒。两处改动：

  1. 验证用的 ie_gemv_q8q8_tr 读权重时用非时间性读取（TH_LOAD_NT），权重不占缓存，8 个 token 的激活值留在缓存里。 基准测试 8 个 token：34816 × 160 从 497 到 558 GB/s。算术不变，结果逐位不变。
  2. 草稿头只保留 token 0 到 131071（BPE 编号小的 token 大致是常见的）和最后的特殊 token：0.72 GB → 0.38 GB。 三个提示的接受率与全部 token 时相同；65536 时聊天的接受率下降（60.8% → 56.4%），32768 更差。 结果（7 个草稿，pmin 0.6，个/秒）：

  提示	以前	现在
  写代码	85.5	92.2
  续写 wikitext（prose 提示）	54.2	55.9
  聊天	52.3	55.3

  三个提示、3 和 7 个草稿，输出都与不用草稿时逐个 token 相同。现在各部分都接近带宽：单 token 矩阵向量乘约 630 GB/s， 验证约 580 GB/s。试过但没有用的：只给 8 个 token 的验证专门写一个核函数（累加器更少、每步两个块）：基准测试快 3%，实际不快。

• llama.cpp 的数字来自 llama-server 的 predicted_per_second（同一请求跑两次，取第二次）。 它用的是把 ggml-org 权重和 unsloth 的 MTP 层合并成的一个文件（qwen35.block_count = 65，nextn_predict_layers = 1）。 llama.cpp 的聊天回答在 64 个 token 处结束，所以聊天一行没有对比。

• 推测：草稿数为 3 时每步约 63 毫秒（不用草稿时每个 token 46 毫秒）。多出的时间主要是 3 次 MTP 预测 （每次读 MTP 层 0.45 GB 和输出层 1.35 GB）。

从上下文起草（prompt lookup，默认开启，IE_LOOKUP=0 关闭，2026-10-18）

思路来自 HyperQwen（syv-ai/HyperQwen）。改写、引用、修改代码时，输出大段复述输入。 这时可以不靠 MTP 层，直接从已有的文本里找后续 token 当草稿（src/gen.c 的 lookup_draft）。

1. 在已有 token（提示加已生成的）里找最长的匹配：某个更早的位置，它前面最多 16 个 token 与当前末尾相同。
2. 这个位置后面的 token 就是上下文草稿，每步最多 15 个。
3. 只在 MTP 同意时使用。原因：16 个 token 的验证约 78 毫秒，8 个 token 的验证约 48 毫秒。猜错的草稿很贵。
   • 匹配 ≥ 16 个 token：MTP 链的前 2 个草稿与上下文草稿相同时，用上下文草稿，MTP 链的其余步骤不再运行。
   • 其他情况：MTP 链照常运行；它的全部草稿（至少 2 个）都与上下文草稿相同时，把上下文草稿后面的 token 接上。
4. 验证步骤不变，所以输出不变，只是草稿不同。

为了一次验证最多 16 个 token：线性注意力状态从 8 份改为 16 份（显存多 1.2 GB）， ie_gdn1 也用于 9 到 16 个 token 的验证（以前 8 个以上用 ie_gdn，大核函数里没有它）。

结果（Q4_K_M，7 个 MTP 草稿，大核函数开，400 个 token，个/秒）：

提示	只用 MTP	加上下文起草	每步 token 数
写代码	103.4	102.6	5.87 → 5.87
续写（prose）	92.3	104.4	4.43 → 5.78
中文	47.7	47.7	1.79 → 1.79
聊天	71.6	71.5	3.15 → 3.15
改写英文代码（重命名函数、加注释，1069 个 token 的提示）	108.3	138.8	7.00 → 10.50
改写中文说明（替换一个词，979 个 token 的提示）	105.5	170.0	6.88 → 14.25

• 六个提示的输出都与只用 MTP 时逐个 token 相同；改写中文和写代码两个也与不用草稿时逐个相同。

• 采样（温度 0.7，top-k 20）也能用：改写英文代码 118.6 个/秒，每步采样 0.4 毫秒。

• 试过但没有用的：

  • 只要 MTP 第一个草稿相同就用上下文草稿：写代码慢 4% 到 10%。
  • MTP 不同意时也接上上下文草稿：写代码慢 3.7%。
  • 不让 MTP 链提前结束（IE_LOOKUP_LONG=99）：改写中文 170 → 155。

• ie_attn_kv 加进大核函数（2026-10-18）：位置数大于 768 时，注意力用 ie_attn_kv。以前它不在大核函数里， 所以长提示的解码和验证都改回逐个启动。现在：

  • 大核函数的 LDS 只有 16 KB：K 和 V 在 LDS 里保持 fp16（8 KB，和 at_o 用同一块 LDS），读出时转成 f32（转换精确，结果逐位相同）； 合并的权重写在 part 里 m_t 的位置（和 attn_split_body 一样）。
  • 它在大核函数里是函数调用（KERNEL_CALL，不内联）。内联时大核函数多溢出寄存器，短提示的解码慢 1.2%（29.09 → 29.43 毫秒/token）； 不内联时 29.13，和以前相同。

  测试	以前	现在
  改写中文，MTP + 上下文起草	170.0	178.0 个/秒
  改写英文代码，MTP + 上下文起草	138.8	145.7 个/秒
  改写中文，只用 MTP	105.5	113.4 个/秒
  979 个 token 的提示后普通解码	32.14	31.25 毫秒/token
  写代码，MTP + 上下文起草（短提示）	102.6	102.4 个/秒

  输出都与以前逐个 token 相同；大核函数开和关时，长提示的解码输出也逐个相同。

DFlash 2 起草器（--dflash，2026-10-18）

DFlash 2（z-lab / Inco AI，Apache-2.0）是一个 5 层的小模型（Q4_K_M 文件 1.1 GB，incoai/Qwen3.8-27B-DFlash2-GGUF）。 它一次前向就给出一整块草稿，而 MTP 层每个草稿要单独运行一次。

1. 主模型每次运行时，把第 5、19、33、47、61 层的输出（共 5 × 5120 个数）复制到 df_feat（OP_COPY）。
2. 上下文部分（第 2 段）：fc（25600 → 5120）和 RMSNorm 之后，算出每个起草器层的 K、V，存进起草器自己的 KV 缓存（f32）。
3. 块部分（第 3 段）：块 = [t(P), 掩码 × 15]，位置 P 到 P+15。每层：RMSNorm → 两抽头动态卷积（ie_df_conv）→ 注意力 （ie_df_qkv 做每头 RMSNorm、RoPE 和存 K/V；ie_df_attn 不分因果，滑动窗口 2048）→ 卷积加回残差；FFN 同样前后各一次卷积。 最后用主模型的输出矩阵算出每个位置的 logits，在 GPU 上取前 16 个候选。
4. 选择器（主机）：从 t(P) 出发，每个位置的分数 = logit + (前一个 token 的码本行 · 隐藏投影 · 候选的码本行)，逐位置取最大，得到一条路径。
5. 草稿在概率（前 16 个候选上的 softmax）低于 0.3 时截断。验证步骤不变，所以输出与不用草稿时逐个 token 相同。
6. 自动选择：每一步都更新两个起草器的上下文（MTP 层只跑 KV 部分），按衰减后的“token 数 / 时间”比较两者的速度， 另一个快 10% 以上时才切换，每 16 步试一次另一个。文本里有长匹配（≥ 16 个 token）时用 MTP 加上下文起草。

结果（Q4_K_M，400 个 token，个/秒）：

提示	MTP + 上下文起草（以前的默认）	只用 DFlash（15 个草稿）	DFlash 自动选择
写代码	102.6	130.5	128.3
续写（prose）	104.4	100.1	106.5
中文	47.7	37.4	47.1
聊天	71.5	80.0	77.8
改写英文代码	145.7	-	145.3
改写中文说明	178.0	-	172.0
7949 个 token 的长提示	50.0	44.5	47.7
写代码，采样（温度 0.7，top-k 20）	99.3	-	136.8

• 所有提示的输出都与只用 MTP 时逐个 token 相同。
• 中文：两个起草器每步接受的 token 都只有约 0.7 个；MTP 不确定时少起草，验证更便宜，所以自动选择几乎都用 MTP。
• 一次验证 16 个 token 约 73 毫秒，8 个约 48 毫秒，所以块长 15 只在草稿可信时才划算（概率截断）。
• 选择器有用：15 个草稿时每步 9.97 个 token，不用选择器 8.67 个。
• 起草器每步约 8.9 毫秒；长上下文（窗口填满）时约 11.2 毫秒，主要是 ie_df_attn（每个查询头各读一遍窗口内的 K/V）。
• ie-serve --dflash（2026-10-19）：同一个写代码请求（400 个 token，温度 0，--ctx 65536），只用 MTP 3.98 到 4.06 秒， 加 DFlash 3.26 到 3.37 秒（快约 20%），回答文字相同。起草器的 KV 缓存改为环形（窗口 + 32 个位置，85 MB），与上下文长度无关。
• 显存：上下文起草让线性注意力状态从 8 份变成 16 份（多 1.2 GB），--ctx 131072 因此放不下了。现在分配前先查空闲显存， 放不下 16 份时用 8 份（一次验证最多 8 个 token，日志里有一行说明）。加 DFlash 时，32 GB 里上下文最多 65536。
• 还没做：起草器注意力按 KV 头共享读取；MTP 层在 DFlash 步骤里也要更新（每步约 1.5 毫秒）。

提示批量处理（预填充）

GPU 上，提示按每块最多 256 个 token 处理（--chunk，1 到 512）。一块只读一次权重。按块的大小分两条路径：

块中 token 数	矩阵乘	注意力	线性注意力	和单 token 结果的关系
2 到 4（推测解码的验证）	ie_gemv_q8q8_tr（每 wave 4 行）	每个 token 启动一次	ie_gdn1（每个头一个工作组，归一化在核函数内）	逐位相同
5 到 16	ie_gemv_q8q8_tr	每个 token 启动一次	ie_gdn（每个头 4 个工作组）+ ie_gnorm	矩阵乘逐位相同；线性注意力求和顺序不同
17 以上	ie_gemm_q8r（WMMA，权重直接进寄存器；每行块数不是 4 的倍数时用 ie_gemm_q8）	ie_attn_pfg（一次启动）	分块并行：ie_gdn_prep + ie_gdn_wy + ie_gdn_seq + ie_gnorm	接近（矩阵乘相对差约 1e-7，线性注意力约 2e-6）

另外：RMSNorm、量化、SwiGLU、q/k 归一化 + RoPE、门控归一化都是一次启动，网格的第二维是 token。

WMMA 矩阵乘（ie_gemm_q8）：

• 指令：v_wmma_i32_16x16x16_iu8（int8 × int8 → int32）。lane 布局由 tools/wmma_test.c 在 GPU 上确认。
• 工作组：128 行 × 64 个 token；K 每次 4 个块（128 个权重）。
• 权重块、激活块和它们的缩放先放进 LDS（每次 16 字节的合并读取）；读取下一段 K 的同时计算当前段（寄存器预取）。
• 每个 wave 算 32 行 × 32 个 token（2 × 2 个 WMMA 块）。每个 32 权重的块：两次 WMMA 得到精确的 int32 块点积，再乘 (权重缩放 × 激活缩放) 累加到 float。
• tools/gemm_bench.c 单独计时。

新结构 ie_gemm_q8r（默认）：

• 工作组 256 行 × 64 个 token。每个 wave 只管自己的 32 行（2 个行块）和全部 64 个 token（4 个 token 块）。 别的 wave 不用它的权重，所以权重不经过 LDS：每个 lane 每个块读一次 16 字节，直接进寄存器；下一段 K 的权重在当前段计算时预取。
• 只有激活值（所有 wave 都要用）和缩放经过 LDS，双缓冲，每段 K 一次同步。
• 块内 K 的顺序：lane 的半区 h 在第一条 WMMA 用字节 16h..16h+7，第二条用 16h+8..16h+15。A 和 B 用同一顺序， 每个字节只用一次，所以块点积还是同一个精确的 int32。已确认：困惑度与 ie_gemm_q8 完全相同（0.8B 9.8314，27B 4.3227）。
• 编译器会把 4 个块的指令交错，寄存器不够而溢出到显存。循环里的一个分支（k >= nk 时退出）把块分开，避免溢出。

tools/gemm_bench.c（512 个 token，T MAC/s）：

形状（行 × 每行块数）	ie_gemm_q8	ie_gemm_q8r
34816 × 160	49.0	63.5
5120 × 544	49.9	63.9
14336 × 160	49.1	62.4

已确认（同一核函数的调试版本，34816 × 160）：去掉缩放运算后 101 T MAC/s；去掉 WMMA、只留缩放运算时 1259 微秒（完整版 1437 微秒）。 推测：剩下的限制是缩放运算。每个块每个输出元素要 3 条向量指令（减常数、乘两个缩放、累加）， 这些指令和 WMMA 在同一个 SIMD 上，不能同时执行。

已确认：WMMA 矩阵乘和矩阵向量乘的结果相对差约 1e-7（IE_GEMM_CHECK）。困惑度：0.8B 第 0 段 9.84（逐位路径 9.89）， 27B 第 0 段 4.3172（逐位路径 4.3322；llama.cpp 4.3243）。差别来自 Q8 激活值舍入翻转后被模型放大，不是矩阵乘的误差。

Qwen3.8-27B Q8_0，512 个 token 的提示	token/秒
一次一个 token（最早）	21.6
每块 16 个，每 wave 4 行	172
每块 16 个，小核函数按 token 合并	217
每块 256 个，WMMA 第 1 版（每 wave 一个 16 × 16 块，直接从全局内存读）	276
WMMA 第 2 版（LDS 暂存，每 wave 2 × 2 块）	725
加上一次启动的预填充注意力	866
加上缩放的向量化	938
加上寄存器预取下一段 K	1120
加上线性注意力每个头分给 4 个工作组	1157
加上按 KV 头分组的预填充注意力、按 token 的 RMSNorm、只给最后一个 token 算输出层	1351
加上分块并行的线性注意力（2026-10-08 测：每块 256 个 1390，每块 512 个 1480；不用它时 1320 / 1393）	1390
加上新结构的矩阵乘 ie_gemm_q8r（每块 256 个 1598，每块 512 个 1758）	1598
llama.cpp b11222（Vulkan），pp512	1283

试过但没有用的：fp16 WMMA（权重和激活在 LDS 中转成 fp16，IE_GEMM=h）：更慢（每块 209 对 188 毫秒）； K 每次 8 个块：更慢。

27B 预填充的时间（512 个 token 一块，硬件追踪，2026-10-08）：矩阵乘 270 毫秒（79%，约 48 T MAC/s）， 线性注意力 28 毫秒（以前逐 token 的 ie_gdn 是 46.5 毫秒），注意力 15 毫秒，SwiGLU 11 毫秒。

分块并行的线性注意力（提示块，17 个 token 以上）。把块分成 32 个 token 一段。设一段开始时的状态是 S0， γ_t 是这一段第 1 到 t 个 token 的衰减之积。门控增量规则可以写成：

• Δ = U − W S0，其中 U = T diag(β) V，W = T diag(βγ) K，T = (I + L)⁻¹，L[t][s] = β_t (γ_t/γ_s) k_t·k_s（s < t）。U 和 W 与 S0 无关。
• 输出 O = diag(γ) Q S0 + M Δ，其中 M[t][s] = (γ_t/γ_s) q_t·k_s（s ≤ t）。
• 新状态 S = γ_n S0 + Kᵀ diag(γ_n/γ) Δ。

已确认：tools/gdn_chunk.py 在 numpy 中把这个等式与逐 token 递推比较（误差约 1e-16）。三个核函数：

1. ie_gdn_prep：每个 token 一个工作组。卷积 + SiLU、q/k 的 L2 归一化、β 和 log 衰减，写进临时缓冲区。
2. ie_gdn_wy：每个（值头，段）一个工作组，全部并行。算 L、M，再用前代法算 U 和 W（每个线程一列）。
3. ie_gdn_seq：每个值头 4 个工作组（每个管 32 列），按顺序处理各段。wave w 管 4 列；lane l 管 4 个键维度。 X S0（X 是 32 行 × 128）先在每个 lane 内对它的 4 个键维度求部分和，再用 DPP 做 lane 间的“归约-分散”， 最后每个 lane 得到一行的结果。wave 之间不用交换数据。下一段的输入在当前段计算时预取到寄存器。

已确认：IE_GDN_CHECK=1 下，0.8B 和 27B 的各层、块大小 17 到 512，输出和状态与 ie_gdn 的最大相对差 ≤ 4e-6。 0.8B 第 0 段困惑度 9.83 到 9.87（随块大小变，不用分块算法时 9.875）；27B 4.3227。 MTP 推测解码与普通解码的输出仍然相同（27B 三个提示，3 和 7 个草稿；0.8B 两个提示，1、3、7 个草稿）。 推测：普通解码的输出和以前不同的位置（例如 27B 对话提示第 7 个生成的 token）是两个候选 logit 只差约 0.03 的地方。

512 个 token 的提示，token/秒	逐 token ie_gdn	分块并行
0.8B，每块 256	15734	21452
0.8B，每块 512	18433	27117
27B，每块 256	1320	1390
27B，每块 512	1393	1480

试过的版本（27B，48 层合计）：ie_gdn_seq 每个头一个工作组、状态 64 个寄存器：56 毫秒；每头 4 个工作组：43 毫秒； 加寄存器预取：29 毫秒；改成 lane 管键维度 + DPP 归约：14 毫秒。ie_gdn_wy 的 LDS 行宽从 128 改成 132（列读取无冲突）：42 → 8.5 毫秒。 推测：矩阵乘离硬件峰值还远。每个块每 wave 要做 96 条普通向量指令（整数转浮点、乘缩放、累加），而 WMMA 只有 8 条。

采样

--temp、--top-k、--top-p、--seed。K ≤ 64 时前 K 个在 GPU 上选，主机读回 K 对（值，ID）；否则主机读一行 logits（1 MB）再选。抽样在主机上做。 和 MTP 一起用时是"采样后比对"：模型在每个位置从自己的分布采样一个 token，草稿等于这个采样结果才接受。 草稿是 MTP 的最大值（固定的），所以每个输出的 token 都是模型分布的一个样本，输出分布和不用 MTP 时相同。

已确认：每输出一个 token 只用一个随机数，顺序和不用 MTP 时相同，而且 logits 逐位相同。所以同一个种子下， 用 MTP 和不用 MTP 的输出完全相同（0.8B：300 个种子，每个比较 3 个 token；27B：种子 42，300 个 token）。

Qwen3.8-27B，写代码，300 个 token，--temp 0.7 --top-k 20 --top-p 0.95	token/秒
不用 MTP	21.7
3 个草稿（接受率 92.1%）	50.7

贪心时 3 个草稿是 60.1 个/秒。采样多出约 3 毫秒/token：每步从 GPU 读 4 行 logits（4 MB）， 再在主机上从 248320 个 logits 里取最大的 K 个。推测：把取前 K 个放到 GPU 上可以省掉大部分。

27B 解码的硬件追踪

rocprofv3，每个 token 47.1 毫秒（含追踪开销；不追踪时 46.0 毫秒）：

部分	毫秒/token	说明
矩阵向量乘	43.7	平均约 627 GB/s（标称值的 98%）；5120 × 6144 只有 568 GB/s
RMSNorm（129 次）	0.71	每次 5.5 微秒
线性注意力（48 次）	0.47	每次 9.8 微秒
其他核函数	0.30
核函数之间的间隔	1.9	532 次启动

只读权重的下限是 27.2 GB ÷ 640 GB/s = 42.5 毫秒。剩下的约 3.5 毫秒（8%）分散在间隔、RMSNorm 和较小的矩阵上， 没有一个大的单项。

大核函数（megakernel，默认开启，IE_MK=0 关闭，2026-10-15）

做法：解码一步原来有 677 次核函数启动，验证一步也有约 690 次。主机把它们记录成命令表（每条 256 字节：核函数编号、 工作组数、参数），异步复制到显存，然后用一次协作启动运行 ie_mk。128 个常驻工作组（每个计算单元 2 个）按顺序执行 各条命令：工作组 p 做第 p、p + 128、… 个工作组的工作，命令之间做全局同步。互不依赖的矩阵向量乘（q/k/v/z/beta-alpha， gate/up；多 token 时按 token 的跨步区间判断）在同一段里执行，中间不同步。每个核函数的代码不变，所以输出逐位相同 （解码比较了全部 logits；MTP 比较了 4 个提示的全部 token）。

两个版本：ie_mk.hsaco（解码）和 ie_mk_mtp.hsaco（验证）。验证版里多 token 的矩阵向量乘用每个 wave 4 行的 ie_gemv_q4kq8_tr / ie_gemv_q6kq8_tr（与 8 行的 ts2 / tr8 逐位相同）：8 行的版本单独就要约 192 个寄存器， 放进大核函数会大量溢出，反而更慢。草稿步骤仍是逐个启动（IE_MK_MTP=1 也让草稿走大核函数：更慢）。

测量（Q4_K_M，贪心，300 个 token；输出与逐个启动相同）：

项目	逐个启动	大核函数
解码	30.49 毫秒/token	28.93 毫秒/token（快 5.1%）
MTP 写代码，7 个草稿	94.2 个/秒	102.7 个/秒（快 9.0%）
MTP 续写，7 个草稿	82.5	87.5（快 6.1%）
MTP 中文，7 个草稿	46.2	48.3（快 4.5%）
MTP 聊天，3 个草稿	68.0	70.2（快 3.2%）
Coder390 Q6_K 解码	37.80 毫秒/token	36.15 毫秒/token
Coder390 Q6_K 写代码，7 个草稿	92.2	97.5

其他测量：

项目	结果
工作组之间单向交接	0.2 微秒
全局同步，8 到 128 个工作组	0.3 微秒
等待同步时预取下一个矩阵（IE_MK_PF）	在测量误差内

要点：

• 常驻工作组数由 LDS 和寄存器决定。注意力合并的权重原来放在 16 KB 的 LDS 里，现在写回 part（全局内存）； 大核函数的 LDS 从 30 KB 降到 14 KB，每个计算单元能放 2 个工作组。硬件按计算单元放置工作组：占用率查询按 WGP 计算，256 个寄存器时说能放 96 个，实际不能，全局同步超时。
• 全局同步必须先让每个 wave 做释放（release），再在工作组内会合，最后由线程 0 计数。只有线程 0 做释放时， 其他 wave 的写入可能还没完成：验证时输出不同。
• 每个等待都有上限（MK_SPIN 次轮询）。超时后所有工作组退出，主机报错，不会让 GPU 卡住。
• 预填充（超过 16 个 token 的提示块）仍是逐个启动。
• 加载时用一次空的协作启动检查工作组数；不能常驻时减半，到 32 还不行就关闭。一次运行里遇到大核函数没有的核函数时， 这一类运行改回逐个启动（屏幕上有一行说明）。等待同步时预取（IE_MK_PF）默认关闭：s_prefetch_data 会读到地址之后， 小矩阵时越过了分配的末尾（Qwen3.5-0.8B 非法访问）。

小模型（解码，128 个 token，输出与逐个启动相同）：

模型	逐个启动	大核函数
Qwen2.5-0.5B Q4_0	1.485 毫秒/token	1.236（快 20%）
Qwen2.5-1.5B Q4_0	2.612	2.382（快 10%）
Qwen2.5-3B Q4_0	4.357	4.017（快 8%）
Qwen3.5-0.8B Q8_0	2.176	1.958（快 11%）
Qwen3.8-27B Q4_K_M	30.45	29.07（快 5%）

核函数内的事件时间线（prof/，2026-10-15）

prof/sm_profiler.h 是 leepoly/sm-profiler（CUDA）的 AMD 版，用 C 写， 接口和缓冲区布局与原版相同：核函数里每个（工作组，wave）记录区间事件和瞬时事件，不用原子操作；主机读回后写成 Chrome 追踪 JSON。每个 WGP 是一个进程，每个（工作组，wave）是一个线程。

CUDA 原版	AMD 版
%globaltimer（纳秒）	s_sendmsg_rtn MSG_RTN_GET_REALTIME（100 MHz 恒定时钟，各计算单元一致；gfx9/10 用 s_memrealtime），存成纳秒，精度 10 纳秒
%smid	s_getreg HW_REG_HW_ID1：SE、SA、WGP 编号合成一个 WGP 编号；另存 SIMD 编号
%laneid、warp	mbcnt；wave 大小 32 或 64
CUDA 运行时	dlopen libamdhip64.so，不需要 ROCm 头文件
设备端 C++	C（clang amdgcn，本仓库的核函数）；定义 SMP_HIP 后也可在 HIP C++ 里用
B/E 两条记录	一条 X 记录（开始和时长）；另有 smp_print_stats 按事件统计

新增 *_at 接口：常驻核函数可以传入逻辑工作组编号。已确认：R9700 上 128 个工作组分布在 32 个不同的 WGP 编号上 （这块卡有 32 个 WGP）。

示例：make build/smp_demo build/smp_demo.hsaco，然后 build/smp_demo build/smp_demo.hsaco 写出 demo_trace.json。

大核函数：IE_MK_TRACE=文件 记录第 IE_MK_TRACE_AT（默认 10）次大核函数启动里每个工作组的时间线（每段工作和 每次全局同步等待），打印统计并写出文件。27B 解码一步约 60 万个事件。记录时这一步约慢 6%。

被验证（Lean 证明，对所有输入成立，make proofs 约 1 分钟）：core/ie_core.h 中的函数满足 laws/ie_laws.cpp 的 35 条定律（make proofs 检查 37 个定理）： GEMV 网格和 lane 划分的覆盖与唯一性；Q4_0 和 Q8_0 重排地址在数组内且是双射；半字节和字节解包符合 GGUF； Q4×Q8 和 Q8×Q8 的块点积公式等于精确的整数点积； Q4_K（MAT_Q4K）：重排地址和 GGUF 源地址在数组内、字地址是双射、半字节重排符合 GGUF（q4k_nibble）、 子块点积 ie_q4k_dot 等于 Σ q·a（q4k_dot；单 token、多 token 核函数和 CPU 都调用它）、预解包的半字节点积 ie_q4k_dot_u 等于同一个和（q4k_dot_u）； Q6_K（MAT_Q6K）：地址在数组内（q6k_addr）、高位打包和解包还原 GGUF 的 6 位值（q6k_unpack）、 无符号和有符号子块点积等于 Σ (q − 32)·a（q6k_dot、q6k_dots）、预解包的点积相同（q6k_dots_s）；跨 lane 求和与 wave 树归约等于普通求和（模 2^32）； 被 ie_plan_check 接受的规划：缓冲区在 arena 内、256 对齐、同时存活的缓冲区不重叠。

只被测试：GGUF 解析、模型加载、贪心规划器（它的输出由验证过的检查函数检查）、 两个后端的完整前向计算、所有浮点部分（rmsnorm、rope、softmax、swiglu、合并进核函数的 bias、残差和量化副本）、 注意力和取最大值核函数里的跨工作组同步（原子计数器）。

没有测试：长于 4020 个位置的上下文、Qwen2.5-0.5B 和 1.5B 以外的真实模型、头的大小超过 256 的模型（会走旧的注意力核函数）。

信任边界

C 代码调用的是同一个验证头文件 core/ie_core.h；cpp-lean 把同一个文件翻译成 Lean。 下面这些是被信任的，没有被证明：

1. C 语义与 Lean 模型之间的联系：EDG 前端 + cpfe-lean 翻译是正确的。
2. 编译器：gcc / clang（主机与 amdgcn）正确地编译了这些函数。
3. v_dot4_i32_iu8 指令的行为等于模型 ie_dot4_us（AMD RDNA4 ISA 文档）。
4. 调用方给出的前提：例如 ie_sizes_ok 成立、缓冲区真的有核心假设的长度。引擎在加载时检查大小，但缓冲区分配本身没有被证明。
5. 浮点计算不在验证范围内：只有整数部分（索引、点积、归约顺序、内存规划）被证明。 GPU 上的浮点 wave 归约使用同样的顺序，但浮点加法不满足结合律，所以“与 CPU 只差舍入”是测试结论，不是证明结论。