Repository navigation
fix(ffi): validate model geometry at the C ABI load boundary - #6
Merged
Merged
Conversation
C ABI 路径 tinyllm_load 此前完全不校验模型几何:validateModelConfig 只在 InferenceEngine::Load 被调用(全仓仅此一处)。 而 attention kernel 里 kv_head(q_head) = q_head / (num_heads / num_kv_heads) 是整数除法。num_heads 不被 num_kv_heads 整除时 kv_head 会超出 [0, num_kv_heads),最后一个 token 的 K/V 读越过缓冲末尾。已用最小复现 + Compute Sanitizer 证实(Hq=14/Hkv=3:Invalid __global__ read,13 errors, cudaDeviceSynchronize 返回 unknown error —— CUDA 上下文被毒化,不只是当次 请求失败)。 GGUFParser::extractModelConfig 对这类元数据是"补默认值"而非报错(并且 head_dim = hidden_dim / num_heads 也是不校验整除性的截断除法),所以必须在 进入 kernel 前显式拒绝。修复在 extractModelConfig 之后、loadGGUF 之前调用 Validator::validateModelConfig:一处覆盖整条 C ABI 路径(attention prefill/ decode、RoPE、FFI 缓冲尺寸),对合法模型零行为变化。 不改 kernel 层:kernels/ 没有 Result 通道,在那里"校验"只能静默 return, 比现状更糟。也不只修 forwardPaged:那样 prefill 分支与策略 2 仍然暴露。 新增 tests/test_validator.cpp(纯 host,无 GPU 也运行): - validateModelConfig 9 项单元测试(Validator 此前零覆盖),含 1..32 全部 非整除 head 组合的穷举,以及 14/3 这个越界回归锚点; - 2 项 C ABI 边界回归,用字节级构造的 GGUF 断言 tinyllm_load 拒绝 14/3 且 不误拒 14/2。 验证:全量 198 passed / 11 skipped(11 项均需真实 GGUF 模型);新增 11 项全过; clang-format 18.1.8 全仓 0 violation。变异检验:移除 tinyllm_load 中的校验 调用后边界测试失败(错误信息变成 "load weights: ..."),确认它不是装饰性测试。 Co-authored-by: CommandCodeBot <noreply@commandcode.ai>
holtwood
added a commit
that referenced
this pull request
Sep 14, 2026
…ackage 按 owner 委托的决议,关闭 §4.2 与 §12 的全部待决项: - Q1 保持 flat 参数签名;补上真正的理由——逐元素相等门禁已覆盖"传参顺序" 这一失败模式,原先"POD 字段顺序成为跨 TU 契约"的反方论证站不住(C++17 下 逐字段命名赋值本就与顺序无关)。 - Q2 抽取共享 tile loop;PR-1 增加硬性要求:必须附 attention_decode 的前后 kernel_bench 对比证明无性能回归,出现可测回归则退回复制实现。 - Q3 冻结「非法块 id = 零行且参与 softmax」为稳定契约;可观测计数留作 follow-up。 - Q4 保持要求 direct 与 legacy 逐元素严格相等。 - Q5 修正原设计错误:原文称"连续版静默取整,属既有缺口"并提议在 forwardPaged 校验,实际 Validator::validateModelConfig 早已实现该校验,只是不在 C ABI 路径上; 正确修复位置是 tinyllm_load,已拆为独立 PR #6。新增 §7.1 记录证据链。 - Q6 保持"kernel 级收益不得外推为 TTFT/TPOT"的表述边界。 - Q7 新增:显式记录 direct 路径令 k_scratch/v_scratch 冗余(本任务保留分配, 避免改 ffi.cpp;显存收益记为 follow-up)。 - Q8 删除不可达的"几何不支持则回落"分支——触发需 head_dim > 7800,既不可测试 也不构成真实 fallback;fallback 只由 TLLM_PAGED_ATTENTION 开关表达。 §12 增加「决议」「独立性声明」「本轮修正的设计缺陷」三节,如实记录本包由作者 编写也由作者汇总决议,不满足独立 reviewer 要求,并指出 PR-1 是唯一有真实爆炸 半径的决定,合并前应经第二方复核。 验证:仅文档改动;docs 本地 npm run build exit 0(vitepress 1.6.4)。 Co-authored-by: CommandCodeBot <noreply@commandcode.ai>
…lidation # Conflicts: # CHANGELOG.md
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
问题
Validator::validateModelConfig实现了正确的模型几何校验(num_heads % num_kv_heads、hidden_dim % num_heads、head_dim偶数),注释还写明"静默截断会导致 kv_head 映射错位"。但它全仓只在
inference_engine.cpp:65被调用一次,C ABI 路径tinyllm_load(
ffi.cpp,分页/策略 1 与 paged-serving 走的那条)完全不调用 —— grep 计数为 0。而 attention kernel 用整数除法做 GQA 映射:
num_heads不被num_kv_heads整除时kv_head越界,最后一个 token 的 K/V 读越过缓冲末尾。证据(不是推理)
最小复现
Hq=14, Hkv=3, D=64, visible=8,把 K 放在显存池尾部以避免相邻分配掩盖越界:kv_head = 13/4 = 3,合法范围[0,3),越界读 64 个 half = 128 字节,正好落在分配边界后1 字节。
unknown error说明 CUDA 上下文被毒化 —— 不只是当次请求失败。补充:
GGUFParser::extractModelConfig对这类元数据是"补默认值"而非报错,且head_dim = hidden_dim / num_heads同样是不校验整除性的截断除法。修复
在
extractModelConfig()之后、loadGGUF()之前调用Validator::validateModelConfig:不是只堵 attention;
kernels/没有Result通道,在那里"校验"只能静默 return,比现状更糟;
forwardPaged:那样 prefill 分支和策略 2 仍然暴露。测试
新增
tests/test_validator.cpp(纯 host,无 GPU 的 CI 机器同样运行):validateModelConfig9 项单元测试(Validator此前零覆盖),含 1..32 全部非整除 head 组合的穷举,以及
14/3这个越界回归锚点;tinyllm_load拒绝14/3几何,且不误拒合法的
14/2。验证
变异检验:移除
tinyllm_load中的校验调用后,ModelConfigBoundaryTest.TinyLlmLoadRejectsNonDivisibleHeadCounts失败(错误信息变成load weights: ...),确认该测试是真门禁而非装饰。范围
不是 attention 实现变更。
paged-serving。extractModelConfig的"补默认值"行为本身(改动面更大、可能影响宽松解析的既有行为),仅在其下游加校验;如需收紧可另开任务。
权限说明
本 PR 只提交,不合并(
merge_pr: false)。