[linx-toolchain][ops-20260904] clang-15 对内联 bf16 CUBE matmul codegen SIGABRT

[ziyang-cheng]

组件版本清单

组件	仓库	分支 / tag	commit
SuperNPUBench(看护 demo)	ziyang-cheng/SuperNPUBench	tileop-guard-batch1(PR #96)	7163e60
SuperScalarModel	LinxISA/SuperScalarModel	codex/consolidate-post-main-fixes-20260903	49547742
Linx-TileOP-API	LinxISA/Linx-TileOP-API	linx	0566283
llvm-project	LinxISA/llvm-project	dev-llvm15_56	67d3ac9
SuperNPUBench(release 参考)	PTO-ISA/SuperNPUBench	ops-20260904	a0ddcc3

看护 demo 见 SuperNPUBench PR #96：PTO-ISA/SuperNPUBench#96
工具链指纹：clang++ md5 e427d1429c0e、gfrun md5 0c433cd11c00。

通用复现步骤

git fetch origin tileop-guard-batch1 && git checkout 7163e60
source microbenchmark/tileop-guard/env.sh   # COMPILER_DIR / GFRUN
bash microbenchmark/tileop-guard/run_guard.sh <domain> <case>

以下每条均已用「唯 pto-spec 合规」判据自证 demo 写法无误，缺口在模型/工具链侧。

---

linx-2 · clang-15 对内联 bf16 CUBE matmul codegen SIGABRT

涉及接口：TMATMUL + fixp::bf16()（F322BF16，bf16 输出）。

问题：把 CUBE bf16 matmul 直接写在 main() 内联时，clang-15 在中端 "Function Pass Manager" 阶段
abort()（exit 134），非诊断错误：

#8 ... abort ./stdlib/abort.c:81:7
clang-15: error: clang frontend command failed with exit code 134

复现：bash run_guard.sh cube tmatmul_bf16（内联版）。

自证 demo 合规 + 是编译器 bug：

• demo 符合 spec：fixp::bf16()=F322BF16=码 16，在 B.FPATR 合法 PreQuantMode（16..20）内；matrix-postprocess.asl
  把 BF16 作合法目标转换。用法与 release microbenchmark/fixp BF16 模式一致（run_single<__half,__bf16>）。
• 编译器对任何输入都不应 SIGABRT——把同一段逻辑包进 noinline 函数即正常编译（release fixp 正是 noinline 结构），
  证明代码语义合法，是 clang 在 __bf16 CubeAccumulatorM32 内联实例化下的 codegen 健壮性 bug。
• 不用 noinline 绕过（会失去看护意义）；demo 保持内联以暴露该 bug。

建议：修 LinxV5 后端/clang 对内联 bf16 CUBE 累加器的 codegen（不应 abort）。

[VV0003]

复现与定性（2026-09-09，LLVM dev-llvm15_56@3682dca968a5 + TileOP linx）

用本 issue 引用的 SuperNPUBench PR#96 原始测例（tmatmul_bf16.cpp @ 940fb45）+ 配套 guard_common.hpp/benchmark.h 完整复现，并做了三层最小化。结论：TMATMUL + fixp::bf16() 的 CUBE lowering 本身没有 bug，崩溃全部来自测例中的标量 BF16/FP16 初始化与读回转换，与 TileOP #50 的早期分析一致，但本轮把根因定位到了后端的明确缺口：

崩溃 A（原始测例形态）

hd[i] = (__bf16)0.0f;   // 或 ha[i] = (__half)(0.01f * i)

blkc 头把 fp32→bf16 实现为内联汇编 v.cvt.fp322bf16（返回 i16）。janus（非 SIMT）后端没有注册 i16 寄存器类（只有 i64/f64/f32/f16），类型合法化需要 promote 这个 CopyFromReg(i16, glue, divergent) 节点，而 PromoteIntegerResult 不支持 CopyFromReg → UNREACHABLE（即本 issue 的 exit 134 abort）。

崩溃 B（干净版读回形态）

return (int)hd[argc & 31];   // bf16 → s32 标量转换

同样的窄类型缺口，v.cvt.bf162s32 输入 i16 经 getCopyToParts 断言 "Copying to an illegal type!"。

纯 CUBE 路径验证通过

把初始化换成 __builtin_memset、去掉 __bf16 标量读回后，同一 TMATMUL + fixp::bf16() 编译完全正常，产物：

BSTART.CUBE  TMATMUL, FP16
B.FPATR      16, 0, 0, ...    # F322BF16 post-quant

修复评估（重要）

尝试了两条修复路线后评估如下：

1. 给 janus 注册 i16/i8/i32 寄存器类（upstream 目标的常规做法）：实测不可行——janus 的 CodeGen 围绕"i64 是唯一合法整数类型"设计，i16 一旦合法，store(s16)/setcc(i32)/any_extend 等整族节点失去 promote 路径且无选择 pattern，连锁缺口逐层暴露（CC_LinxV5 断言 → setcc → store s16…）。等于要实现完整的窄标量类型后端，是大工作包。
2. 在 LegalizeIntegerTypes 加 CopyFromReg 的 promote：受限于 glued CopyFromReg 的 vreg 类型不可单侧加宽，上游 LLVM（截至 main）同样没有此支持——所有上游目标都注册全部窄整型从而避开此路径。

建议

• 短期：SuperNPUBench guard 测例改用 guard_read_bin/memset 初始化并避免 __bf16 标量算术（PR#96 侧调整），即可全部编译通过——CUBE 数值路径不受影响；
• 长期：将"janus 窄标量类型支持（i8/i16/i32 寄存器类 + 选择 pattern + ABI）"列为独立工作包。在那之前，任何含 fp32→bf16/half 或 bf16→int 标量转换的 -mlxbc TU 都可能触发同类崩溃，这与 TMATMUL 无关。

本票建议改题为"janus 窄标量 inline-asm 转换类型合法化崩溃"或转移焦点至标量板块；bf16 CUBE matmul 部分可标记为已验证无缺陷。

[VV0003]

定性更正（补充上一条评论，请以本条为准）

上一条把根因归为"janus 缺窄标量类型后端支持（大工作包）"，经架构侧确认后该定性作废。正确分析如下：

标量转换与 tile 转换是两条不同路径，本 issue 只涉及前者

• tile ↔ tile dtype 转换：走 TileOP TCVT wrapper（BSTART.TEPL TCVT, <dtype> bundle），本 issue 的 TMATMUL + fixp::bf16() CUBE 路径（BSTART.CUBE TMATMUL, FP16 + B.FPATR 16）已验证完全正常；
• 标量值 C++ cast（测例第 8-10 行 ha[i] = (__half)(0.01f * i) 等对普通数组的赋值）：走编译器 builtin 头 linx_blkc.h 的 DEFINE_CONSTRUCT/DEFINE_OPERATOR 宏。

真正的根因：linx_blkc.h 的 cast 宏用错了指令域

宏展开为内联汇编 v.cvt.*（SIMT/向量指令）：

asm volatile("v.cvt.fp322bf16 %1.fs, -> %0.h" : "=vr"(short) : "vr"(float));

按当前架构约束，janus（block 编程）模式下所有点开头的 v./l. 向量指令都不应被使用——v.cvt 的 instruction pattern 只存在于 SIMT 寄存器类（LinxV5InstrInfoSIMT.td），janus 下出现即非法。此前观察到的两个崩溃（CopyFromReg(i16) promote UNREACHABLE / getCopyToParts illegal type）都只是这个非法写法的症状："vr" 约束 + short 存储让非法指令先死在类型合法化，而不是编出非法 v.cvt。

修复方向（修正后）

1. linx_blkc.h cast 宏改为标量实现：janus 下生成标量 fcvt.*（FSU，支持 FP64/FP32/FP16/E4M3 互转与定点族）或纯位运算（BF16=FP32 高 16 位截断，两条整数指令即可表达）；v.cvt.* 形式在当前架构下无合法消费者，应移除或仅保留给未来 simt 模式；
2. janus 后端无需窄标量类型支持工作包（上一条的该建议作废）；
3. 可选的后端加固：janus 下对 v.* 内联汇编给明确编译 diagnostic，替代当前的 UNREACHABLE 崩溃，让同类误用 fail-closed 且可诊断。

对 PR#96 测例的即时建议（不变）

初始化改用 memset/guard_read_bin、避免 __bf16/__half 标量 C++ cast，即可全部编译通过；CUBE 数值路径无缺陷。linx_blkc.h 源头仓修复落地后，原写法也应恢复可用。

[VV0003]

修复已合入 dev-llvm15_56（commit 10abd6eea670）

按前两条评论的定性，修复落在 clang/lib/Headers/linx_blkc.h（v.cvt 写法的源头）：

• fp8_e4m3 及 fp64/fp32/fp16 互转：改用标量 FSU 转换指令 fcvt.<src>2<dst> / fcvtn.<src>2<iN> / scvtf/ucvtf，内联汇编全走 64 位 "r" 约束 carrier——asm 操作数从头到尾保持合法标量类型；
• bf16：无标量 FCVT 编码，改为软件位运算（fp32→bf16 = 高 16 位截断 + RNE 舍入；反向 = 低 16 位左移），并补齐 explicit 整型 conversion operator 消除歧义；
• tf32/hf32：与 FP32 共享存储，cast 纯透传（零指令）；
• fp8_e5m2/e8m0/e6m2/hif8/fp6/fp4 及跨格式 cast：无标量编码且当前无标量使用面，cast 宏显式禁用并注明使用 tile 级转换；
• v.cvt.* 全部移除，头文件顶部注明架构规则（block 模式下点开头向量指令不可用，不得重新引入）。

验证

• SuperNPUBench PR#96 tmatmul_bf16.cpp 原测例（本 issue 主场景）+ 两个缩减变体（bf16 数组初始化循环 / bf16 标量读回）全部编译通过；
• 生成汇编 v.cvt 零出现：bf16 初始化循环降为 srli/slli 位运算，浮点转换走 fcvt，BSTART.CUBE TMATMUL, FP16 + B.FPATR 16 的 CUBE lowering 完全不变；
• LinxV5 CodeGen/MC 回归与基线一致；TileOP fixture 的 4 个失败经原版头对照确认全部 pre-existing（retired TQUANT、旧 MATMUL 拼写、SIMT constructs、TileRegion getReg 断言）。

PR#96 测例可恢复原写法（(__bf16)0.0f 初始化），无需再改 guard_read_bin。