[LinxV5] rms_norm_binary：dyn 插入 TMOV(INVALID) 致 gfsim Deadlock；static 把 Local tile cache spill 成 TSTORE/TLOAD S64

[wangyuascend-spec]

代码分支 / 版本

本机 checkout（不是远程 default），用于复现：

仓	分支	HEAD (short hash + subject) / 版本说明
SuperNPUBench	drop-rms-norm-binary-workspace / 复现包 issue/rms-norm-binary-dyn-tmov-gfsim	kernel：465fd5d Drop unused GM workspace from rms_norm_binary. 包：d77b8c8 Add rms_norm_binary dyn TMOV/gfsim deadlock repro pack.
SuperScalarModel	main	3e1ebf2f Merge pull request llvm#286 from LinxISA/chore/untrack-local-exec-logs
Linx-TileOP-API	本机 src linx	c02dae6 [TileOP] Enforce Group TMATMUL shape contracts。toolchain 默认 pin 分支 linx。本 repro 编译用 overlay 6369707（clang++ -resource-dir …/clang_res_6369707）。tip c02dae6 与当前 clang 组合会在 TQUANT %Z 上编不过，故 kernel 用 overlay。
llvm-project	verify/dev-llvm15_56（跟踪 origin/dev-llvm15_56）	b945a5d01323 [handoff] Record Group TMATMUL Issue 18 acceptance。clang++：15.0.4 (linx64v5-musl-local b945a5d01323de2f38989288d0a4c423abff3ad1) Target linx64v5-unknown-linux-musl
linx-toolchain-build	main	e6a31ef macOS: install GNU tar for the package step
pto-spec	main	0b8ce516 arch: add extension first-use profile hook (llvm#101)

• gfrun mtime：2026-08-20 14:15:51（SuperScalarModel/bin/gfrun）
• gfsim mtime：2026-08-20 19:27:02（SuperScalarModel/bin/gfsim）
• TileOP include 来源：linx-toolchain-build/src/Linx-TileOP-API 的 linx 分支；实际 -resource-dir 指向 commit 6369707 的头文件树。
• 复现包（含 ELF / diss / 源码快照）：https://github.com/wangyuascend-spec/SuperNPUBench/tree/issue/rms-norm-binary-dyn-tmov-gfsim/Daily/2026-08-21/rms_norm_binary_dyn_tmov_gfsim
• 预编译 ELF：https://github.com/wangyuascend-spec/SuperNPUBench/raw/issue/rms-norm-binary-dyn-tmov-gfsim/Daily/2026-08-21/rms_norm_binary_dyn_tmov_gfsim/rms_norm_binary_dyn.elf
• kernel 源（同一份）：https://github.com/wangyuascend-spec/SuperNPUBench/blob/drop-rms-norm-binary-workspace/benchmark/one-level-arch/kernels/normalization/rms_norm_binary/rms_norm_binary_pto.hpp
• 对照 PR：Drop unused GM workspace from rms_norm_binary PTO-ISA/SuperNPUBench#72

相关但不相同：#60（编译期 TMOV ICE）、#61（整数块 ctzll BFI 物化错误，gfrun 精度）。本 issue 是 功能 gfrun 已 PASS 之后，timing gfsim 因插入的 TMOV 缺元数据 Deadlock。

问题描述与分析

症状

rms_norm_binary dynamic 入口（fp16 [1,8192]，tile_r=1024，pow_r=4096，PE1）：

仿真	结果
gfrun -f <res_check ELF>（无 -t 1）	PASS R2=0，precision max_abs=0
gfsim -f <无 res_check ELF> --pto-v02 true -s core.singleTierMode=true	FAIL Deadlock cycles=4147，abort signal 6

同算法的 static 模板（rms_norm_binary_static，同 shape）gfrun / gfsim 都 PASS（gfsim Total Cycles=3959）。

源码不写 TMOV

Local cache 拷贝明确用 TMULS，注释写了 device 无 TCOPY、TMOV 运行时下标会断言：

namespace rms_bin {
constexpr int kWsCols = 32;
constexpr int kMaxLevels = 8;

inline int64_t GetCacheId(int64_t idx) {
    return static_cast<int64_t>(
        __builtin_ctzll(static_cast<unsigned long long>(idx + 1)));
}

template <typename TileVec>
inline void update_cache(TileVec (&upd)[kMaxLevels], TileVec &cur, int64_t &r) {
    const int cid = static_cast<int>(GetCacheId(r));
    for (int j = 0; j < cid; ++j) {
        TADD(cur, cur, upd[j]);
    }
    TMULS(upd[cid], cur, 1.0f);
    ++r;
}
}

dyn 入口里：

using tile_v = Tile<Location::Vec, float, tA, rms_bin::kWsCols,
                    BLayout::RowMajor, 1, 1>;
tile_v updateTile[rms_bin::kMaxLevels];
// Pass1 每块：TLOAD → TCVT → TMUL(x,x) → TROWSUM(cur) → update_cache(updateTile, cur, r)

完整文件见复现包 rms_norm_binary_pto.hpp / 测试入口 rms_norm_binary.cpp。

编译器实际发出的 TMOV

llvm-objdump -dl（与 gfsim 用的无 res_check ELF 一致，.text 起址 0x112e4）。GetCacheId 的 SWAR 刚算完 srli 56，立刻插入 TMOV：

   1162c:  srli    t#1, 56, ->s0
   11630:  c.setc.eq s0, zero
   11632:  BSTART.TLSU  TMOV,
   11636:  B.IOT        t#1, mask=1111, last, ->t<128B>
   1163a:  C.BSTART.STD

没有 B.DATR（dtype），没有 B.DIM。 同文件里合法的 TADD/TMULS 是：

   BSTART.TEPL  TADD, FP32
   B.DIM s2, 0 ->lb0
   B.DIM s2, 0 ->lb1
   C.B.DIMI 32 ->lb2
   B.IOT t#1, u#1, mask=1111, last ->t<128B>

这份 ELF 里这样的 TMOV 至少 18 处（tmov_sites.diss）。

gfsim 解码 / Deadlock

[C:4146] Deadlock detected at thread 0, cycles=4147
[C:4146] B16 STID0 BPC 0x11632 [TMA FALL] rdy:0  TMOV, lb0:1, lb1:1, lb2:1 INVALID  TLSU allocated
[C:4146] B21 STID0 BPC 0x1166c [TMA FALL] rdy:0  TMOV, lb0:1, lb1:1, lb2:1 INVALID  TLSU allocated
...
ERROR: Deadlock execution at:
  B9 STID0 BPC 0x11524 TCVT FP16
  Cycle:4147  Wait cycles:3999
  0 instructions correctly executed so far
ASSERTION FAILED: Deadlock detected: simulation stalled, no new blocks committed for too long

0x11632 正好是上面那条 BSTART.TLSU TMOV。DIM 被当成 1,1,1，dtype INVALID。功能模型 gfrun 不走这条 timing 检查，所以 gfrun 仍 PASS。

期望

• 不要为 TMULS/TADD 的 tile rename 插入缺元数据的 TMOV
• 若必须插 TMOV：带上与源 tile 一致的 dtype（FP32）和 DIM（本 kernel Valid=1,1，Cols=32 → TSize 合法），gfsim 能像 static 一样跑完
• dyn 的 gfsim 应到达 benchmark end，而不是 cycle 4147 Deadlock

实际

clang 在 tile_v updateTile[8] 跨 R-split / update_cache 循环保活时插入 无 DATR、无 DIM 的 BSTART.TLSU TMOV；gfsim 报 TMOV INVALID 后 stall。

复现方式

不要 gfrun -t 1 / gfsim -t 1。PE1，不要 --conf fourpe。gfsim 用 无 res_check ELF（带 host I/O 的 res_check ELF 会让 gfsim SIGSEGV，与本 bug 无关）。

0. 最快：直接跑预编译 ELF

# clone 复现分支
git clone git@github.com:wangyuascend-spec/SuperNPUBench.git
cd SuperNPUBench
git checkout issue/rms-norm-binary-dyn-tmov-gfsim

GFSIM=/home/wangyu/Code/SuperScalar/SuperScalarModel/bin/gfsim   # 改成本机 gfsim
ELF=Daily/2026-08-21/rms_norm_binary_dyn_tmov_gfsim/rms_norm_binary_dyn.elf
"$GFSIM" -f "$(pwd)/$ELF" --pto-v02 true -s core.singleTierMode=true
# 期望：Deadlock cycles=4147，BPC 0x11632 TMOV INVALID

1. 从源码重编（与本机相同 overlay）

export COMPILER_DIR=/home/wangyu/Code/SuperScalar/linx-toolchain-build/output/linx_blockisa_llvm_musl/bin
export CFLAGS="-resource-dir /home/wangyu/Code/SuperScalar/SuperNPUBench/Daily/2026-08-20/rms_norm_verify_main/clang_res_6369707"
# overlay 头文件树也在本机 Daily；若无，用 toolchain 自带 TileOP（需能编过 TQUANT）

cd SuperNPUBench/benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary
make TESTCASE=rms_norm_binary COMPILER_DIR="$COMPILER_DIR" DType=__half
# ELF:
# .../output/kernel/normalization/rms_norm_binary/elf/kernel_normalization_rms_norm_binary_rms_norm_binary_DType__half.elf

$COMPILER_DIR/llvm-objdump -dl <elf> | rg TMOV
# 应看到：BSTART.TLSU TMOV + 单独一条 B.IOT，没有 DATR/DIM

Daily/2026-08-21/rms_norm_binary_dyn_tmov_gfsim/compile.sh 会编 ELF 并拷到该目录。

2. 对照：gfrun 功能（应 PASS）

需要 res_check=on ELF + golden（gen_rms_norm_binary_data.py）。

make TESTCASE=rms_norm_binary COMPILER_DIR="$COMPILER_DIR" DType=__half res_check=on
python3 src/gen_rms_norm_binary_data.py \
  -o ../../../compare/kernel_normalization_rms_norm_binary_rms_norm_binary_DType__half
# 无 -t 1
/home/wangyu/Code/SuperScalar/SuperScalarModel/bin/gfrun -f <res_check.elf>
# 期望：Success to Reach the End of Benchmark! R2 = 0

3. 对照：static 同 shape（gfsim 应 PASS）

make TESTCASE=rms_norm_binary_static COMPILER_DIR="$COMPILER_DIR" DType=__half \
  G_A=1 G_R=8192 TILE_A=1 TILE_R=1024 POW_R=4096
"$GFSIM" -f <static.elf> --pto-v02 true -s core.singleTierMode=true
# 本机：PASS，Total Cycles=3959

测试 main 只有调用：

int64_t tiling_info[5] = {1, 8192, 1, 1024, 4096};
rms_norm_binary<dtype>(input, tiling_info, output, 1e-6f);

[wangyuascend-spec]

补充：static 把 Local updateTile[] spill 成 TLOAD/TSTORE S64

同一份 update_cache（源码只有 TADD / TMULS，没有对 cache 做 TLOAD/TSTORE）。static 模板 gfrun / gfsim 都能跑完（本机 gfsim Total Cycles=3959），但 clang 没有把 tile_v updateTile[8] 留在 tile 寄存器里，而是 spill 到栈。

源码（static 与 dyn 相同）

using tile_v = Tile<Location::Vec, float, tA, rms_bin::kWsCols,
                    BLayout::RowMajor, 1, 1>;  // Cols=32, Valid=1,1, FP32
tile_v updateTile[rms_bin::kMaxLevels];        // Local，期望一直在 tile reg
// ...
TROWSUM(cur, sq);
rms_bin::update_cache(updateTile, cur, r);
// update_cache: TADD(cur, cur, upd[j]); TMULS(upd[cid], cur, 1.0f);

main：

rms_norm_binary<dtype, 1, 8192, 1, 1024, 4096>(input, output, 1e-6f);

实际发出的 spill（无 res_check ELF）

kernel 入口 subi sp, 1536。cid 来自 GetCacheId 的 SWAR 后：

1. 读档（应对 TADD(cur, cur, upd[j])） — 从栈 TLOAD, S64，步长 160：

114bc:  BSTART.TLSU  TMOV,                 # 同样缺 DATR/DIM
114c0:  B.IOT t#1 ->t<128B>
114ca:  BSTART.TLSU  TLOAD, S64
114ce:  B.DATR NORM.normal, Zero
114d2:  B.IOT ->u<128B>
114d6:  B.IOR [a6],[]
114da:  C.B.DIMI 16 ->lb0
114dc:  C.B.DIMI  1 ->lb1
114de:  C.B.DIMI 16 ->lb2
114e4:  BSTART.TEPL  TADD, FP32
...
114f6:  BSTART.TLSU  TMOV,                 # 又一条无 DATR/DIM
114fe:  C.BSTART.STD COND, .LBB4_13
11500:  addi a6, 160 ->a6                  # spill slot stride

2. 写档（应对 TMULS(upd[cid], cur, 1)） — TMULS 之后 TSTORE, S64 到 sp+224+cid*160：

11518:  addi zero, 160 ->t
1151c:  mul  a5, t     ->t                 # cid * 160
11520:  addi sp, 224   ->t
11524:  add  t#1, t#2  ->t
11526:  addi t#1, 32   ->a5
11530:  BSTART.TEPL  TMULS, FP32
        B.DIM 1,1 ; C.B.DIMI 32 ; B.IOT ->u<128B> ; B.IOR [a6]
11546:  BSTART.TLSU  TSTORE, S64
1154a:  B.DATR NORM.normal, Null
1154e:  B.IOT u#1
11552:  B.IOR [a5],[]
11556:  C.B.DIMI 16 ->lb0
11558:  C.B.DIMI  1 ->lb1
1155a:  C.B.DIMI 16 ->lb2

要点：

• 真正的输入/输出是 TLOAD, FP16 / TSTORE, FP16（GM）。上面的 S64 对的是 栈上的 cache tile，源码没有这条路径。
• spill 把 FP32 tile_v 当成 S64 存（DIM 16×1×16 = 128B，和 1×32 float 物理大小一致，类型不对）。
• static 里同样有无 DATR/DIM 的 TMOV（和 dyn 同一类），只是 gfsim 这次没 Deadlock。

期望

• tile_v updateTile[8] 应分配在 tile 寄存器，不要用 TLOAD/TSTORE S64 做 spill
• 若必须 spill：应保持 FP32 tile 语义，而不是 S64；也不该靠缺元数据的 TMOV 做 rename

复现

# 预编译 ELF
git clone git@github.com:wangyuascend-spec/SuperNPUBench.git
cd SuperNPUBench && git checkout issue/rms-norm-binary-dyn-tmov-gfsim
llvm-objdump -dl Daily/2026-08-21/rms_norm_binary_dyn_tmov_gfsim/rms_norm_binary_static.elf | rg 'TLOAD, S64|TSTORE, S64|TMOV'

# 从源码重编
export COMPILER_DIR=.../linx_blockisa_llvm_musl/bin
export CFLAGS="-resource-dir .../clang_res_6369707"
cd benchmark/one-level-arch/test/kernel/normalization/rms_norm_binary
make TESTCASE=rms_norm_binary_static COMPILER_DIR="$COMPILER_DIR" DType=__half \
  G_A=1 G_R=8192 TILE_A=1 TILE_R=1024 POW_R=4096

摘录 / ELF：

• https://github.com/wangyuascend-spec/SuperNPUBench/blob/issue/rms-norm-binary-dyn-tmov-gfsim/Daily/2026-08-21/rms_norm_binary_dyn_tmov_gfsim/static_tile_spill.diss
• https://github.com/wangyuascend-spec/SuperNPUBench/raw/issue/rms-norm-binary-dyn-tmov-gfsim/Daily/2026-08-21/rms_norm_binary_dyn_tmov_gfsim/rms_norm_binary_static.elf

[wangyuascend-spec]

对照更新（plain rms_norm，不是 binary）：

同一批 README pin（llvm b945a5d / TileOP c02dae6 / SoftCore 78adfe32）上，rms_norm dyn/static PE1 的 Newton 循环也有同样的 DATR-less TMOV（f8211181 + 仅 B.IOT，无 B.DATR/B.DIM，PC 0x11586），但 gfsim PE1 PASS。把同一编码编进 *_PE4.elf 后：

• gfrun -s softcore.multiThreadNum=4 PASS
• gfsim --conf fourpe Deadlock，ROB 上 TMOV … INVALID TLSU allocated rdy:0

所以 #64 的「缺 DATR 的 TMOV」在 1PE 上并不总死锁；rms_norm PE4 这条更像 fourpe TLSU 没走 gfrun ExecuteTMA 那套缺省推断。已另开 SoftCore 票：https://github.com/LinxISA/SuperScalarModel/issues/327

本票仍建议编译器给 Local TMOV 补上 B.DATR/B.DIM（binary dyn PE1 仍会因同类编码死锁），但不要把 PE4-only 失败只归因到 clang。

日志/ELF：/home/wangyu/Code/SuperScalar/SuperNPUBench/Daily/2026-08-21/rms_norm_readme0820/