问题摘要
当前 Linx-TileOP-API linx 分支生成三操作数 TEPL 汇编：
BSTART.TEPL mode, function, dtype
当前 LinxV5 LLVM 汇编器接受的 canonical 形式为：
BSTART.TEPL selector, dtype
因此使用最新 TileOP 与最新 LLVM 编译 SuperNPUBench RMSNorm / GroupNormGrad 时，在汇编阶段出现 Match Instruction Error，无法生成 ELF。
环境
Linx-TileOP-API: 943311c52f166296e5d25017651dc97e8c41811e (linx)
llvm-project: 8a4bb5bda33acaf6152deabc63922b3d1249cdc9
SuperNPUBench: d387ec09917aafdedcb83fb179847f72b17b5ef8 (PR #84)
TileOP make check 通过，LLVM inline-asm dtype modifier 回归也通过；但两者组合编译真实 normalization kernel 失败。
复现
在 SuperNPUBench/benchmark/one-level-arch 下：
COMPILER_DIR=/path/to/linx_blockisa_llvm_musl/bin
make -C test/kernel/normalization/group_norm_grad \
TESTCASE=group_norm_grad COMPILER_DIR="$COMPILER_DIR" \
DType=__half N_BATCH=2 C_CH=16 G_GRP=4 HxW_SZ=16 TILE_HW=8 \
res_check=on diss
典型错误：
template_asm.hpp:7487:5: error: Match Instruction Error!
BSTART.TEPL 1, 27, %D1
instantiated into assembly here:
BSTART.TEPL 1, 27, FP32
同一 kernel 还会触发：
BSTART.TEPL 0, 0, FP32
BSTART.TEPL 0, 1, FP32
BSTART.TEPL 2, 5, FP32
RMSNorm 会触发同类错误，例如：
BSTART.TEPL 0, 27, FP16
BSTART.TEPL 2, 0, FP32
BSTART.TEPL 2, 7, FP32
根因
include/jcore/template_asm.hpp 仍普遍输出分离的 Mode/Function 字段，而当前 LLVM MC/CodeGen 测试使用单一 TEPL selector。数值 selector 的编码为：
selector = mode * 32 + function
示例：
0,27 -> 27
1,27 -> 59
2,0 -> 64
2,5 -> 69
2,7 -> 71
将安装目录中本次算子涉及的形式临时改为合并 selector 后，RMSNorm 以及 GroupNormGrad 的动态/静态、HxW>1/HxW=1 四个精度 ELF 均可成功编译。GroupNormGrad 四个 gfrun 均正常结束、R2=0，dx/dgamma/dbeta 在 atol=0.02、rtol=0.02、mse_tol=0.001 下全部通过。
这说明失败发生在 TileOP 生成汇编与 LLVM parser 的接口处，不是算子数值实现或模型执行问题。
建议修复
- 在 TileOP 模板的生成源中统一将 Mode/Function 合成为 canonical selector，不只修本次出现的几个实例。
- 重新生成 include/jcore/template_asm.hpp，避免手工维护生成产物。
- 增加真实 Linx clang syntax/codegen 回归，覆盖至少 Mode 0、1、2，并包含 RMSNorm/GroupNormGrad 用到的 selector。
- 检查所有 BSTART.TEPL emitter，避免仍残留三操作数形式。
- 与 LLVM 当前命名 selector和数值 selector 的 round-trip 测试对齐。
验收标准
相关：Linx-TileOP-API#29、Linx-TileOP-API#32、PTO-ISA/SuperNPUBench#84。
问题摘要
当前 Linx-TileOP-API linx 分支生成三操作数 TEPL 汇编：
当前 LinxV5 LLVM 汇编器接受的 canonical 形式为：
因此使用最新 TileOP 与最新 LLVM 编译 SuperNPUBench RMSNorm / GroupNormGrad 时，在汇编阶段出现 Match Instruction Error，无法生成 ELF。
环境
TileOP make check 通过，LLVM inline-asm dtype modifier 回归也通过；但两者组合编译真实 normalization kernel 失败。
复现
在 SuperNPUBench/benchmark/one-level-arch 下：
典型错误：
同一 kernel 还会触发：
RMSNorm 会触发同类错误，例如：
根因
include/jcore/template_asm.hpp 仍普遍输出分离的 Mode/Function 字段，而当前 LLVM MC/CodeGen 测试使用单一 TEPL selector。数值 selector 的编码为：
示例：
将安装目录中本次算子涉及的形式临时改为合并 selector 后，RMSNorm 以及 GroupNormGrad 的动态/静态、HxW>1/HxW=1 四个精度 ELF 均可成功编译。GroupNormGrad 四个 gfrun 均正常结束、R2=0，dx/dgamma/dbeta 在 atol=0.02、rtol=0.02、mse_tol=0.001 下全部通过。
这说明失败发生在 TileOP 生成汇编与 LLVM parser 的接口处，不是算子数值实现或模型执行问题。
建议修复
验收标准
相关：Linx-TileOP-API#29、Linx-TileOP-API#32、PTO-ISA/SuperNPUBench#84。