PLAT=linx standalone build fails with in-repo toolchain: triple mismatch + missing tile-asm/SIMT flags

[Auyuir]

Summary

make ... PLAT=linx does not produce an ELF with the in-repo toolchain (compiler/toolchain/2026-06-22/linx_blockisa_llvm_musl). The triple hardcoded in test/common/Makefile.common (linx64-linx-none-elf) is rejected by the bundled clang, and after working around that, both the tileop_api tests and the accelerator/fusion cases (the ones consumed by gfsim via simall.py) fail to compile due to missing tile-asm / SIMT flags.

Commit c9cb855 ("Route SuperNPUBench tileop smoke toward current Linx clang") already notes this is in-progress ("Not-tested: SuperNPUBench TAdd ELF production; current first failure is benchmark-owned source/toolchain surface mismatch"), so I'm filing this to capture the concrete reproduction and ask what the intended standalone invocation is.

Environment

• Repo: PTO-ISA/SuperNPUBench @ a54fc74 (2026-06-23)
• Toolchain: compiler/toolchain/2026-06-22/linx_blockisa_llvm_musl (LFS-pulled, extracted)
  • clang --version → clang version 15.0.4 (linx64v5-musl-local ...), default triple linx64v5-unknown-linux-musl
  • Registered targets: linx64v5, linx64v5be only
  • Bundled musl sysroot at <tc>/sysroot
• Host: x86_64 Linux

Reproduction

1. Triple mismatch (Makefile vs in-repo toolchain)

cd test/tileop_api
make TESTCASE=TAdd PLAT=linx COMPILER_DIR=<tc>/bin

error: unknown target triple 'linx64-linx-none-elf', please use -triple or -arch

Makefile.common sets CC_O = -c -target linx64-linx-none-elf -fenable-matrix -O2 and CC_LINK ?= -target linx64-linx-none-elf -nostdlib, but the bundled clang only knows the linx64v5 arch (no linx64).

A triple sweep confirms linx64v5-* is required:

linx64-linx-none-elf             FAIL  (unknown target triple)
linx64v5-linx-none-elf           OK    (ELF relocatable, but no sysroot -> <cstddef> not found)
linx64v5-unknown-linux-musl      OK    (ELF relocatable, sysroot found, <cstddef> OK)

i.e. *-none-elf cannot find the bundled musl sysroot; only linx64v5-unknown-linux-musl resolves standard headers.

2. tileop_api tests fail at the shared exit inline asm

With the triple overridden to linx64v5-unknown-linux-musl:

make TESTCASE=TAdd PLAT=linx COMPILER_DIR=<tc>/bin \
  CC_O="-c -target linx64v5-unknown-linux-musl -fenable-matrix -O2" \
  CC_LINK="-target linx64v5-unknown-linux-musl -nostdlib"

test/tileop_api/src/TAdd.cpp:17:10: error: inline-asm should start from BSTART or Tile Call.
        "addi t#1, 1365, ->t\n"
...
        "c.swi t#1, [u#1, 0]\n"
        "BSTOP\n"
6 errors generated.

The same linx_supernpu_exit block (BSTART.STD / lui / addi t#1 / c.swi / BSTOP) is shared by ~10 tileop_api cases listed in test/tileop_api/compile.all, so the whole tileop_api suite is blocked at this point. BSTART.STD and the lui ..., ->t lines parse, but addi t#1 / c.swi / BSTOP are rejected.

The old flags -mlxbc and -mllvm -enable-all-vector-as-tilereg=true (per the c9cb855 diff) were removed because the new clang rejects them; it looks like the flag that enabled tile-register (t#/u#) asm notation has not been re-wired under a new name.

3. accelerator/fusion (gfsim-consumed) cases fail on SIMT launch syntax

cd test/accelerator/fusion
make TESTCASE=fa1 PLAT=linx COMPILER_DIR=<tc>/bin \
  CC_O="-c -target linx64v5-unknown-linux-musl -fenable-matrix -O2" \
  CC_LINK="-target linx64v5-unknown-linux-musl -nostdlib"

include/accelerator_fusion.h:189: error: expected expression
  flashsoftmax_src_exp_dn<...><<<tileSrc::ValidRow, tileSrc::ValidCol, 1>>>(...)
                                          ^
include/accelerator_fusion.h:210: error: unknown type name '__vbuf__'
  __vbuf__ typename tileOut::DType *rescale_out_ptr = blkv_get_tile_ptr(rescale_out);

The CUDA-style <<<grid,block>>> launch syntax and the __vbuf__ qualifier are not recognized — they appear to need a SIMT/vector-core compilation flag that is not present in Makefile.common nor in test/accelerator/fusion/program.list.

4. program.list CFLAGS= is not wired into the build

test/accelerator/fusion/program.list entries pass CFLAGS='-DENABLE_TENSOR_INSTR ...', but Makefile.common only consumes CC_OPTS:

CC_O_ALL = $(CC_O) $(CC_VER) $(CC_OPTS)

CC_OPTS is never assigned and CFLAGS is never read, so when program.list is run via plain make the -D defines do not take effect. (simall.py reads COMPILER_DIR from the env but otherwise just shells out to the make ... lines verbatim, so it inherits the same gap.)

Ask

What is the intended standalone invocation (triple + flags) to produce Linx ELFs from this repo with the bundled toolchain? Specifically:

1. Should Makefile.common use linx64v5-unknown-linux-musl (or is linx64-linx-none-elf expected to resolve on a different/newer toolchain build)?
2. Which flag replaces -mllvm -enable-all-vector-as-tilereg=true to enable the t#/u# tile-register asm used by linx_supernpu_exit?
3. Which flag enables the <<<>>> launch syntax and __vbuf__ qualifiers for the accelerator/fusion cases?
4. Should program.list use CC_OPTS= (or should Makefile.common append CFLAGS)?

If these are expected to be injected by an external runner (referenced in c9cb855 as tools/bringup/run_ai_workload_flow.py, not in this repo), it would help to document the required flags / driver in the README so the repo can be built standalone.

Happy to test any suggested invocation against TAdd and fa1 and report back.

[Auyuir]

更新（2026-06-26）：origin/main 已被 force-push 回退到 aff9508 lineage（旧 a54fc74 lineage 保留在 origin/backup/main-latest），并继续前进到 a8dbf4b。本 issue 最初针对的 a54fc74 状态，现状如下：

• 第 1 点（linx64-linx-none-elf triple 被工具链拒绝）：已过时。main 回到了 -mlxbc -fenable-matrix -mllvm -enable-all-vector-as-tilereg=true（默认 triple，无 -target），工具链接受该组 flag。
• 第 2 点（tileop_api 退出 asm should start from BSTART 报错）：已解决（-enable-all-vector-as-tilereg 重新启用后 tile 寄存器 asm 正常）。另外 tileop_api 单算子测试已在重组中删除，算子搬到 test/kernel/{matmul,fa,...}。
• 第 3 点（fusion <<<>>> / __vbuf__ 不识别）：已解决（-mlxbc 支持）。
• 第 4 点（program.list 的 CFLAGS 未接入 Makefile.common 的 CC_OPTS）：仍成立，main 上 Makefile.common 仍只读 CC_OPTS、不读 CFLAGS。

当前 main（a8dbf4b）构建状态：用 LinxISA/linx-toolchain-build（dev-llvm15_56）本地编出的工具链，compile_all.sh 能编出 46 个 LinxV5 ELF（覆盖 broadcast / concat / fa / gather / matmul / reduction / transpose）。仍失败的：control、sort 因缺 .data 文件（incbin）失败；部分 matmul / element_wise 配置有源码问题（kernels/matmul/matmul_mx.hpp 中 gen_ND2ZZ_offset_Impl 未声明，以及 tileop-api/jcore/TCopyIn.hpp 的 no viable overloaded '='）。

另外：用这些 ELF 喂 gfsim（LinxCoreModel）会在首个块段错误。根因已定位并提到 LinxCoreModel：https://github.com/LinxISA/LinxCoreModel/issues/18 —— 与本仓无关，是 gfsim 参考仿真器对压缩 C.BSTART.STD 入口的建块顺序 bug（GetSrcData 在 AccumulateBlockInfo 建块前解引用空 currentBlock）。