Repository navigation
Dev/f16 fix - #19
Merged
Merged
Dev/f16 fix#19
Conversation
按 docs/todo/24-f16-training-defects.md 的 F1/F2/F3 落地。 - F3（根因订正 + 修复）：GPU `stable=f16` 的「loss 冻结 / inf」**不是**回读链 （cast/异步槽位）问题，而是 CE 的 (1,1) loss 归约按 `stable=f16` 落盘溢出 ——该标量是跨 batch×seq 的 log_prob 求和（≈ -log(vocab)·N），大 token 数下 超 f16 上限 65504；14 号台账的 `4.0137 ≈ 65504/num_valid` 即同一溢出在部分 硬件上饱和为 f16 上限的表现（故「与权重/数据无关、跨进程逐位相同」）。 修复 = `compute_loss.hpp` 稀疏/稠密两条链的 (1,1) 归约输出恒 F32（内部本就 f32 累加，调用方随后本就 cast 到 F32 回读）。GPU 复测原 inf 配置 9.0056→8.6922； stable=f16 连跑 120 步无 NaN。确认无退化（cast 仅 0.4%、无 trace_miss）。 - F1：危险精度组合启动期告警（仅警告，不拒绝）。新增 `include/neuralnet.cpp/cli/cli_precision_warning.hpp`，在 text_train/mnist_train 解析后、进训练循环前打印一次到 stderr；触发 = optimizer=f16（含四字段全 f16）； 退出码与训练行为不变。门禁 `src/cli_precision_warning_test.cpp`（组合矩阵单测 + 子进程跑真实 CLI 捕获 stderr、断言告警出现且退出码不变）。 - F2：覆盖缺口。`f16_precision_test` 的 profile 组合矩阵补 `all=f16` 行、提为默认 ctest（CPU + GPU 两分支都跑），新增 `test_ce_loss_stable_f16_no_overflow` 回归。 - 附带发现：`NN_PROFILE` 显示每步约 20% 在 `copy_from`（与精度无关，f32/f16 同现） → 立项 `18-roadmap` **P2-9**（step 内固定 I/O 开销归因与消除）。 - 文档订正：14（状态改「已修」+根因订正）、24（结论表/§2/§3.1/§4.4/§5）、 18（P1-3 完成、P2-9 立项）、05-mixed-precision、precision.hpp 注、history.md。 验证：CPU build-ab ctest 41/41；GPU build/gpu ctest 43/43（含全部 *_gpu 变体与 ce_fusion_test）；bench/gui_cli_audit.py PASS。 注：本提交一并带入先前工作区中与本主题相关但未提交的改动 （CONTRIBUTING.md / docs/README.md / docs/todo/README.md / docs/todo/25-release-requirements.md）。
- 新增 docs/todo/26-f16-speedup-plan.md：F4 方案报告。口径按用户要求改为
「任意设备都要提速」，并**允许原生 f16**（不再排除）：
· 存储路径 S0–S4（带宽优先、算术留 f32、不依赖 shaderFloat16）：
S0 去 step 内固定开销（先决）→ S1 主导张量 f16 + 融合 CE 消费者 →
S2 f16x2 打包读写 → S3 f16 GEMM（storage-only，f32 累加）→ S4 CPU F16C。
· S5 原生 f16 算术（`#a` / shaderFloat16）：用足 + 验收（A/B 口径
`NN_VULKAN_NO_16BIT_ALU=1/0`）。
实测归因（850M）：步内 88% 是与精度无关的固定开销（flush_batch/begin_batch/
copy_from；根因 = start_frame 复用帧槽阻塞等 fence，PIPELINE_FRAMES=6 < 每步
~7 次 flush）；`--f16` 下 LM head 强制 stable ⇒ logits 仍 f32（最大激活没省字节）；
f16 路径另多出 import ×1.86 / cast 585 次。热路径命中 storage-only `#x` 变体，
证明底座不依赖 f16 ALU。
- 18-roadmap：新增 P2-9（step 内固定 I/O 开销归因与消除）、P2-10（原生 f16 算术
用足）；P2 计数 8→10。
- 24：F4 行与 §4.4 更新（口径 + 指向 26 号；原生 f16 = S5）。
- todo/README：索引加入 26 号。
根因订正：f16 慢的最大单项不是"固定开销掩盖"也不是"logits 仍 f32"，而是
运行期精度分派本身——`#x` 热循环里每次读写都带
`((prec >> i & 1u) ? float(bNh[…]) : bN[…])` 三元，编译器无法化简。
落地（只需 16 位存储，**不需要** shaderFloat16 → 任意设备成立）：
· key#s = "全输入 f16 + 输出 f16"，无条件发射（兜底）
· key#xxxx = 扫描期真实观察到的签名：dry-run 与模型 pass 各跑 f32/f16 两遍，
scan_register（dsl::compute*）+ CPU 引擎 eval_expr* 入口（fold / 直接求值）
按真实张量精度 + 真实 P 记入 ExprRegistry::variants
运行期查找序 = #a（原生 f16）→ #xxxx → #s → #x（任意签名兜底）→ 边界 cast。
特化集合只是优化清单：未枚举到的签名照旧命中 #x，语义/闭合世界口径不变。
实测（850M，无 shaderFloat16；d64/h4/L4/ff256/seq256/b16；15 步 ×3 轮）：
· --f16 step 墙钟 2260 → 2055 ms（f32 2037 → 由 +10.6% 变为 +0.9%）
· 设备 kernel 合计 2163 → 1962 ms；峰值显存 818 → 606 MiB（−26%）
· 单 kernel：CE 融合注意力 146.1 → 120.0 ms（f32 117.2）；
LM head 154.4 → 137.4（f32 137.1）；wgrad 143.9 → 138.8（f32 140.6）
40HX（b64）仍 +30%，残差全部来自一个结构：`#f`/`#ks`/`#kr` 按设计只发 V0
全 f32 → 网格饥饿的 block wgrad 在 f16 下掉出 K 切分（25→79 ms/块 ×4）；
配对旁证 NN_MM_NO_KSPLIT=1 → 细 tile 84 ms/块（细 tile 不够，只有 K 切分能救）
→ 立项 P2-13；优化器批量派发 #b 同理 → P2-12。
文档：26 号新增 S1s 并订正 §2.1「固定开销掩盖」的旧判断；24 号新增 §4.5 四段式
实测（含跨设备矩阵）；18 号 P2-11 完成 + P2-12/P2-13/P2-10 口径；AGENTS 补变体
选路约定并回正两处 stale 的 f16 冻结条目；history 新增一条（说明本条目是对
2026-10-01「签名降级为运行期参数」的**部分回退**，以及为何不冲突）。
回归：ctest CPU 41/41 + GPU 43/43（含 f16_precision_test / ce_fusion_test /
expr_fuse_test）；融合结构 126 条 / 精度变体 289 → 377 条。
缺口：`#f`（细 tile）与 `#ks`/`#kr`（K 切分）只对 V0 全 f32 发射，于是网格饥饿形状
（大 batch 的 wgrad）在 f16 下没有 K 切分可用 → 掉回 64-tile。40HX 实测该结构
f32 `#ks` 25 ms/块、f16 79 ms/块（4 块 +368 ms = 该机 f16 的全部残差）；
配对旁证 NN_MM_NO_KSPLIT=1 → 细 tile 84 ms/块 ⇒ 只有 K 切分能救。
改法（sig 穿进三条生成路径，每个静态签名变体各发一份，键如 key#10003#ks）：
· 部分和缓冲恒 f32（后端 use_ks 按 f32 分配）→ #ks 输出声明钉成 float、写原始
部分和；注入 sig 只决定 A/B 读取类型；#kr 跨段求和后按 sig 落盘（含 RNE）
· 后端「网格变体无需 16 位存储」的豁免收窄到纯 V0 键（key#10003#ks 里仍有
float16_t 缓冲，误豁免会在无 16 位存储的设备上建出非法 pipeline）
· #kr 补发 nn_f16_rne 声明（glsl_f16_store 依赖）
· 运行期 #x 哨兵不叠加网格变体（它多一个 prec PC 槽、绑定翻倍）
实测：
· 40HX（b64）该 wgrad f16 16.0 ms/块 < f32 25.1 ms/块
· 40HX 端到端 --f16 1013 vs f32 1027 ms（−1.4%）
· 850M（b16）--f16 2005 vs f32 2038 ms（−1.6%）
· loss 对拍（40HX）f32 8.6757/8.2646/7.8743 vs f16 8.6676/8.2569/7.8643（<0.15%）
· ctest GPU 43/43；细 tile 14→46、K 切分 14→46（fused shader 617→713）
文档：18 号 P2-13 完成；26 号 S3 补「网格粒度部分已完成」；24 号 §4.5 更新跨设备
矩阵（两台机器 f16 均转为更快）；AGENTS 记录两条易错点；history 新增一条。
缺口：批量派发变体 #b 只对 V0 全 f32 发射 → --f16 下优化器退化为 4 个逐 kernel
融合（980 次调用/个）vs f32 的 1 个批量 kernel。
改法：
· #b 对每个 **extras 为空**的静态签名变体各发一份（键形如 key#0002#b）；
extras 的元素类型不在 ExprPrecSig 编码里 → 带 extras 的 f16 结构仍逐 kernel 降级
· 生成器：批量变体的 buffer reference 元素类型按 sig 选（f16 槽走 float16_t[]）
· 运行期：组键必须带精度身份（否则同组混签名会把 f16 缓冲按 f32 读）；
实例缓冲按 t.precision() 取槽位（gpu_tensor() 默认实参是 F32，f16 会命中
“无 P 精度存储”断言）；批量变体查找与 find_prec_variant_ 同序（精确签名 → #s 兜底）
· 后端「网格/批量变体无需 16 位存储」的豁免同样收窄到纯 V0 键
实测（850M，f16，15 步）：优化器 4 个结构里 2 个（fd3bd5767952c708#0002#b /
a3baa338349afd75#0002#b）从 980 次/个降到 42 次/个；NN_IRC_NO_BATCH=1/0 配对
A/B（3 轮）wall 2019 → 2008 ms（−0.5%，设备侧持平，噪声内）。
残差：另 2 个结构是降级的复合结构（f16 参数 + f32 动量 → 混精度成员），其逐成员
派发绕过批量队列 → 需改 IR-C 降级路径（≈1 ms/step ≈0.5%，另案记录）。
ctest GPU 43/43；批量变体 71 → 183 条。
此前 RAPT/CNN 在 --f16 下仍全程 f32：RAPT 的 27 处 dsl::compute* 都不传精度
实参（落默认 F32），两个残差分支硬编码 Precision::F32；Conv2D 的权重/梯度/
偏置硬编码 F32、融合算子也不传 P。
改法：
· RAPT：27 处 dsl::compute* 补 p_.compute；残差分支改 p_.compute（与
GPTBlock/TransformerEncoderLayer 同口径）；token_emb_/grad_token_emb_ 改
p_.param；rms_norm_forward_/backward_ 的输出张量改 p_.compute
—— 这一条是**必须的**：它此前默认 F32，而内部逐头 op 在 f16 下产 f16，
insert_rows 按 dst 精度选槽位 ⇒ dst/src 错配直接断言
· Conv2D：w_/b_/grad_w_/grad_b_ 改 p_.param；4 处 dsl::compute* 补 p_.compute
· CnnConfig 增 precision 字段；build_cnn_model(_from_spec) 在 add 之前
注入 profile（新增可选形参；默认 profile 逐位不变）
门禁：cnn_smoke_test / rapt_smoke_test 增 --f16；ctest 新增 4 条
cnn_test_f16(_gpu) / rapt_test_f16(_gpu)（断言 loss 有限且下降、梯度无 NaN）。
实测（850M）：CNN f32 2.57538→0.23991 vs f16 2.57543→0.239591；
RAPT f32 3.90125… vs f16 3.90134…（同一模型、同一种子，差异在 f16 舍入内）。
ctest 47/47 全绿。
原判断：每步约 20% 墙钟在 copy_from、三者合计 ≈88% host 墙钟 ⇒ 可消除的
固定开销，且会把计算侧收益吞掉（26 号曾因此把 S0 定为"先决项"）。
实测（850M，d64/h4/L4/ff256/seq256/b16，15 步）：
① staging ring 64MB×4 → 8MB×16（消除 StagingRing::acquire 的 fence 阻塞）：
copy_from 392 → 3 ms，但 flush_batch 955 → 1213 ms —— 两者之和恒 ≈1347 ms，
墙钟 3 轮 2200/2059/2053 vs 原 2044（噪声内）→ 改动已回退
② --flush-interval 1 → 4（flush 105 → 20 次）：墙钟 2044 → 2096 ms（+2.5%，更慢）
③ 设备忙率 95.9% ⇒ host 几乎全程在等 GPU，等待点可互相搬家
结论：无可回收收益，保留默认配置；真天花板 = 墙钟 − 设备 kernel ≈ 86 ms/15 步 ≈ 4.2%
（f32/f16 同幅）。同时作废"S0 是先决、阻塞其余所有项"的判断——S1s 在 S0 未做时
已实测把 --f16 从 +10.6% 拉到 +0.9%。
文档：18 号 P2-9 / 26 号 S0 / 24 号 §4.4 附带发现 全部改为"已评估-不做"并附证据；
history 新增一条（含"否决方案"类型的 A/B 数据）。无代码改动（staging 改动已回退）。
做法：LM head 计算精度从"强制 = stable（--f16 下即 F32）"改为跟随 compute， 即 --f16 下 logits 落 f16（全模型最大张量减半）。历史顾虑（边界 cast 给 loss 链 每个算子各物化一份 f32 副本 → OOM）已由 in-kernel 变体消除，原理上成立。 实测（850M，f16，15 步）却是净亏： · LM head 前向 138.4 → 130.5 ms（输出写量减半，−7.9） · CE 反向 77.4（V0 全 f32）→ 116.4 ms（掉回 #x 运行期分派，+38.9） · 净 +31 ms（端到端 2018 → 2042） 根因：CE 反向链的签名从未被扫描观察到 —— scan_exprs 只调 ce.forward_sparse， 从不调 ce.backward() ⇒ 该结构没有编译期特化变体 ⇒ 带 prec 三元的 #x。 数值：loss 轨迹偏离 ~3%（f32 8.6764/8.2419/7.8675 vs f16 8.7111/8.3892/8.1161） —— f16 logits 使 CE 梯度（softmax−onehot）也按 f16 落盘，收敛变慢。 结论：当前形态下不采用（实现已回退）；若重开，前置条件是让扫描覆盖 CE 反向 （跑一次真实 forward→loss→backward→optimizer）并重新评估质量影响。 文档：26 号 S1 改为"已评估-不做"并附完整证据 + 进度总览；24 号 §4.4 同步。
detail::float_to_half_bits / half_bits_to_float 增 `#if defined(__F16C__)` 快路径：
· 运行时走硬件单指令（_cvtss_sh / _cvtsh_ss）
—— imm8 = _MM_FROUND_TO_NEAREST_INT：恒 RNE、与 MXCSR 舍入模式无关（确定性）
· 常量求值（std::is_constant_evaluated()）与 NN_ENABLE_NATIVE=OFF 的可移植
构建仍走软件位操作 → constexpr 语义与旧行为不变
· <immintrin.h> 放在文件 include 块（放 namespace 内会把 __m128i 等重声明进 nn::）
新增门禁（P2-8 的验收项）：precision_test 增「f16 转换快路径一致性」段——
用 23 条 f32→f16（±0、次正规、2^-25 平局、2^-14 进位、65504/65519/65520/2^16
上溢边界、±inf、NaN 负载）+ 15 条 f16→f32 位型，把**编译期软件路径**与
**运行期硬件路径**逐位对比（constexpr 表项 = 软件，循环内调用 = 硬件）。
实测（CPU，GPT d64/L2/seq64/b4，10 步）：f16 3.6 s → 1.9 s（f32 1.9 s）。
两态验证：NN_ENABLE_NATIVE=ON（F16C）与 OFF（build/portable，纯软件）均通过；
ctest 47/47。文档：18 号 P2-8 / 26 号 S4 标完成并附数据。
根因（bisect 定位）：--precision-optimizer f16 下第 1 步更新后即 NaN（CPU/GPU 同）。 --f16 健康、+stable=f16 健康、+optimizer=f16 复现。机制 = Adam 二阶矩 v ≈ (1-β2)·g² ~ 1e-10 在 f16 下**下溢为 0**，而一阶矩 m ~ 1e-8 仍在次正规区 ⇒ m/(√v+eps) 退化 ⇒ 更新暴涨 ⇒ 参数越界 ⇒ NaN。 修法：Adam::optimizer_state_scale_() —— 状态存储精度 < 4B 时 S = 4096（2 的幂， 与 g 相乘在 f16 下无损），把 (1-β1)·S / (1-β2)·S² / eps·S 作为 K1/K2/K3 的 rparam。 Adam 的更新尺度不变 ⇒ 更新量数学不变（f16 舍入内），但 m/v 回到正常区。 比训练循环里的 loss scaling 更内聚：不改 CLI、不改梯度/裁剪语义、无"是否反缩放" 分支；rparam 恒 f32 传递 ⇒ 原方案担心的"eps flush-to-0"不存在（eps 的相对比例 由 eps·S 保持）。S 只由精度决定 ⇒ 确定性；checkpoint 存的就是缩放后状态。 门禁与实测： · f16_precision_test 新增「四字段全 f16 可训练」（tiny GPT 20 步：loss 有限且下降） · profile 组合矩阵 all=f16（原为 NaN 形态）现 ok=1：3.4778→3.1780（f32 3.4772→3.1796） · 端到端（850M）CPU 100 步 7.78→5.96、GPU 100 步 7.73→6.16、GPU 500 步 6.05→4.05 残留边界：|g| ≳ 0.0625 时 v 溢出 f16 ⇒ 配 --max-norm；F1 启动期告警保持有效。 文档：24 号 F5 标完成 + 新增 §4.6；26 号进度总览更新。
S5（原生 f16 资格扩张）——原生 f16 **允许且已在用**（用户口径）：#a 在查找序里
优先级最高，40HX 实测 9 条结构命中。NN_VULKAN_NO_16BIT_ALU=1/0 配对 A/B
（40HX，b64，f16，15 步 ×3 轮）：ALU 开 1005 ms vs 关 1021 ms ⇒ 现有覆盖贡献
≈1.5%。不再扩到 fold/matmul/reduce 的理由：
· 这三类的耗时主体是**归约累加循环**，累加按数值要求必须保持 f32
（§7.2 / 业界 fp16 GEMM 同口径）⇒ 扩资格只覆盖其中**逐元素**片段（≈1%）
· 其余排除项（rparams / 微常量 / 行列批索引）是原则性的：放宽会把整链隐式提升
回 f32（收益归零）或让微值 flush / 索引失精（静默错值）
· 故"最重的几条链拿不到原生 f16 ⇒ 收益大"不成立
S2（f16x2 打包读写）——前提是"f16 热变体受**事务数**支配"。实测不支持：
若受事务数支配，存储减半的 f16 变体就该显著快于 f32；而实测同结构只换变体时
f16 与 f32 已在 ±2.5% 内（注意力 +2.4%、LM head +0.2%、wgrad −1.3%、
6996c1ccdb7656fc +0.6%；部分结构 f16 反而快 11%）⇒ 字节/事务不是限制项
（注意力核 shared/barrier 主导、matmul 核 f32 累加主导）。打包读预期收益 <3%
且只作用于地址连续快路径（本仓热核多为跨步/广播访问），代价是重写三条生成路径
⇒ 不做（若将来出现真正带宽受限的形状，按四段式重新立项）。
文档：18 号 P2-10 / 26 号 S2 与 S5 改为"已评估-不做/不再扩张"并附数据与理由；
26 号进度总览同步。代码零改动。
落地：Layer::init NVI 在 init_impl 之后统一执行 check_param_precision_() ——
层声明的 param 槽是存储精度时，逐个校验 parameters() 的张量实际存储精度与之一致
（precision_bytes == 0 的 T1_58 等非存储精度声明跳过；compute/stable 只影响中间量、
不要求同精度；optimizer 槽作用于不在 parameters() 里的状态张量）。复合层的
parameters() 递归下探子层 → 每个子层各自校验一次即可。
**它当场抓到 3 处真实静默降级**（数值都对、只是 --f16 的省字节落空）：
· Conv2D 的 w_/b_/grad_w_/grad_b_ 硬编码 F32（P2-3 已修）
· RAPT 的 token_emb_/grad_token_emb_ 硬编码 F32（P2-3 已修）
· **ReLULinearAttention 从未把 profile 转发给它自持的 4 个 Linear**
→ RAPT 注意力投影权重在 --f16 下一直是 f32（本轮补 set_precision_profile；
RAPTBlock 的注释把这条记为"既有状态、P2-3 未落地"，实际漏了）
扫描期直接报错："参数存储精度与 param 槽不一致: 层 Layer 的 (16x16) 实际 f32，
但 param=f16" → 修完 fused_registry 生成恢复正常。
验证：ctest 47/47（含 cnn/rapt 的 f16 门禁；RAPT f16 轨迹 3.90206 起步、
f32 3.90125，差 <0.1%）；三种精度配置（--f16 / 四字段全 f16 / --t1_58）实跑无
误报。文档：18 号 P1-6 标完成并记录三处命中；25 号第 1 项订正"全 f16 已可训练"。
新增的 rapt_test_f16_gpu 用例在纯 CPU 构建（build-ab）里返回 1 而不是 77， 暴露了 rapt_test 一直缺少 cnn_test 那样的 "--gpu 且设备不可用 → 77 = ctest SKIP" 预检（此前 rapt_test 没有注册任何 --gpu 用例，所以从未暴露）。补齐后： · 纯 CPU 构建：build-ab ctest 45/45（4 条 f16/gpu 相关用例记 Skipped） · GPU 构建：rapt 相关 4/4 通过 同时把最终的 40HX 对照数（3 轮：f32 1046 ms vs f16 1013 ms，−3.2%）记入 24 号 §4.5 的跨设备矩阵。
F4（28 号）：跨设备矩阵最终值 · 850M（无 shaderFloat16）--f16 2005 vs f32 2038 ms（−1.6%），峰值显存 606 vs 818 MiB（−26%） · 40HX（原生 f16 ALU） --f16 1013 vs f32 1046 ms（−3.2%，3 轮） · CPU：f16 不再慢于 f32（3.6→1.9 s，P2-8） · 全 f16（四字段）已可训练（F5） S3：原"仍未做"的小 N GEMV 实为**早已覆盖**——matmul_gemv_f16 是 op-level f16 原语 （destroy.inc Phase C2e 建 pipeline），而方案里写的 matmul_q_gemv 是三值打包 GEMM 的 专属原语（T1_58，非 f16 稠密路径）。K 重的形状由 #f/#ks/#kr 在静态签名变体上各发一份覆盖。
缺口（P2-12 残差）：--f16 下 Adam 的每参数复合 kernel {K3,K4} 运行期落降级
→ 逐成员派发（70 参数 = 140 dispatch/步），既不融合也不批量。
根因（两处，均为构建期 scan 与运行期不同源）：
① scan 的优化器 dry-run 未把当前 pass 的 precision profile 透传给
create_optimizer（文本路径 text_train 是透传的）→ f16 pass 里优化器仍按
默认 f32 profile 建状态/定 upd_prec → 融合出的复合结构与运行期不同 →
运行期命中不到 → 降级为逐成员派发。
② 即便结构被登记，复合 kernel（≥2 成员）的**运行期精度签名**在 scan 侧没有
登记入口（单节点结构有 scan_register / eval_expr* 入口；融合结构只在 GPU
end_expr 执行时才算出签名）→ 只有 V0// 载体，拿不到 key#xxxx
静态特化，也拿不到 key#xxxx#b 批量变体。
改法：
· scan 的优化器 dry-run 透传 scan_prof（f32 pass 不变；f16 pass 与运行期同源）；
· scan 的 CPU end_expr 对复合 kernel 补记 (结构, 输入精度位图 + tail 输出
精度)——与运行期 expr_prec_sig_of(d.ins, out.precision()) 同源。
实测（850M，f16，GPT d64/h4/L4/ff256/seq256/b16）：
· [ir-c][degrade] 70 → 0；优化器每步派发 142 → 3（K1/K2 批量 + 复合 #10004 批量）
· 观察签名 126 → 127、精度变体 388、批量变体 193
· f16 loss 轨迹 9.009→8.594 与 f32 9.006→8.591 同步收敛（健康、无 NaN）
· AGENTS §2 编译期特化条：批量变体段标 P2-12 已完成，补两条不变量 （scan 优化器 dry-run 必须透传 profile；复合 kernel 的运行期签名必须在 scan 的 CPU end_expr 补记）与边界（只覆盖 scan 真跑过的 profile）。 · 18-roadmap P2-12 行 → ✅ 已完成（含 40HX 配对实测与 e2e 无收益的如实记录）。 · 24 号台账 §4.5 次要项：残差根因/修法/实测/边界。 · 26 号方案：S1s 残留项与进度总览同步。 · history：两则（构建期工具组 = dry-run 未透传 profile；CPU 引擎组 = 复合 kernel 运行期签名无登记入口），目录计数 22→23 / 12→13。
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.
No description provided.