tileop-api: SuperNPUBench 所需的 4 处增强（THISTOGRAM 操作数修正 / TCmp 类型 / TCast 崩溃 / 新增 TXOR·TSLL·TSRL）

[liujiang833]

tileop-api 增强需求：SuperNPUBench 算子所需的 4 处改动

-mlxbc 会注入 -I <ResourceDir>/include/tileop-api（驱动代码见
clang/lib/Driver/ToolChains/LinxV5Linux.cpp:77 与 LinxV5Toolchain.cpp:118），
这套头文件由组件仓库 LinxISA/Linx-TileOP-API（分支 linx，当前
make install CLANG_PREFIX=... 安装进 toolchain）提供。

为让 SuperNPUBench 的 control/hashtable_lookup_simd 与 sort/topk 端到端跑通，
我们此前在 SuperNPUBench 本地 include/ 里临时打了补丁并用 symlink 覆盖
toolchain 的 tileop-api。现已确认 tileop-api 的归属在编译器侧，故撤掉
SuperNPUBench 的本地副本与 symlink hack，把以下增强上游到编译器自带的 tileop-api。

当前 baseline（Linx-TileOP-API commit 91e3d7a / "Create fa-56 TileOP import
baseline"，2026-06-22）仍存在下列 1~3 的缺陷，且缺少第 4 项的算子。

---

1. jcore/template_asm.hpp：THISTOGRAM 内联汇编操作数编号错位（off-by-one）

THISTOGRAM 模板的操作数编号比实际少 1，导致汇编到错误寄存器。

-    "B.DIM zero, %c4, ->LB2\n"
-    "B.IOT [%5, %6], last, ->%0<%c7>\n"
+    "B.DIM zero, %c5, ->LB2\n"
+    "B.IOT [%6, %7], last, ->%0<%c8>\n"

SuperNPUBench 的 sort/topk 直接调用 tileop-api 的标准 THISTOGRAM wrapper
（不再自带修正副本），因此 topk 阻塞在此项修复上。

---

2. jcore/TCmp.hpp：类型约束过严，拒绝合法的整型/浮点类型

现仅允许 int64/int32/float/__half，u16/u32 等会触发 static_assert。应放宽为
所有整型与浮点：

-  if constexpr (std::is_same<typename tile_shape_in::DType, int64_t>::value ||
-                std::is_same<typename tile_shape_in::DType, int32_t>::value ||
-                std::is_same<typename tile_shape_in::DType, float>::value ||
-                std::is_same<typename tile_shape_in::DType, __half>::value) {
+  if constexpr (std::is_integral<typename tile_shape_in::DType>::value ||
+                std::is_floating_point<typename tile_shape_in::DType>::value) {

（对应的 static_assert 分支同样放宽。）

---

3. jcore/TCast.hpp：SIMT <<<>>> 路径会让编译器在 ClockHands pass 崩溃

现有实现走 TCast_*_Imp<...><<<...>>>(...) 的 SIMT 路径，编译期崩溃。建议改为
tileblock-asm 的 TCVT（BSTART.TEPL TCVT + B.DATR SrcType/DstType）。ISA
v0.56 的 DataType 枚举（UINT16=26、UINT32=25、INT16=18…）支持 u16<->u32 等整型转换：

+#include "template_asm.hpp"
 ...
-  size_t row = src.GetValidRow();
-  size_t col = src.GetValidCol();
-  static constexpr size_t row_lines =
-      tile_shape_in::Rows / (LaneNum / tile_shape_in::InnerCols);
-  if constexpr (is_Nz_layout<tile_shape_in>::value && is_Nz_layout<tile_shape_out>::value) {
-    TCast_NzLayout_Imp<tile_shape_out, tile_shape_in><<<LaneNum, row_lines, 1>>>(dst.data(), src.data());
-  } else if constexpr (tile_shape_in::isRowMajor && tile_shape_out::isRowMajor) {
-    TCast_RowMajor_Imp<tile_shape_out, tile_shape_in><<<col, row, 1>>>(dst.data(), src.data());
-  } else if constexpr (!tile_shape_in::isRowMajor && !tile_shape_out::isRowMajor) {
-    TCast_ColMajor_Imp<tile_shape_out, tile_shape_in><<<row, col, 1>>>(dst.data(), src.data());
-  } else {
-    static_assert(tile_shape_in::isRowMajor && tile_shape_out::isRowMajor, "Storage layout type not supported");
-  }
+  static_assert(tile_shape_in::isRowMajor == tile_shape_out::isRowMajor &&
+                is_Nz_layout<tile_shape_in>::value == is_Nz_layout<tile_shape_out>::value,
+                "TCAST requires identical layout on src and dst (pure type cast, no layout convert)");
+  TCVT_T(dst, src);

---

4. 新增 TXOR / TSLL / TSRL 三个按位/移位 tile op（control 算子需要）

ISA 有 v.xor / v.sll / v.srl，但 tileop-api 没有对应 wrapper。
control/hashtable_lookup_simd（MurmurHash3 纯 tile-op 版）需要它们。
新增内容（仿照已有 TAND/TOR）：

• jcore/TXor.hpp / jcore/TSll.hpp / jcore/TSrl.hpp（含 RowMajor/ColMajor 实现）
• cpu_sim/TXor.hpp / cpu_sim/TSll.hpp / cpu_sim/TSrl.hpp（host 参考实现）
• common/tileop_api.hpp：加 TXOR/TSLL/TSRL wrapper
• common/tileop_api_impl.hpp：include 上述实现

TSRL 走 make_unsigned_t 保证逻辑右移。完整实现见 SuperNPUBench 历史 commit
804e112（include/jcore/{TXor,TSll,TSrl}.hpp、include/cpu_sim/{...}）。

---

影响 / 现状

• SuperNPUBench PR control: pure-tile SIMD hashtable lookup (remove SIMT; add TXOR/TSLL/TSRL) PTO-ISA/SuperNPUBench#12（control）依赖第 4 项（TXOR/TSLL/TSRL），
  上游补齐前 blocked。
• SuperNPUBench PR sort/topk: THISTOGRAM radix-select top-K (gfrun+gfsim PASS) PTO-ISA/SuperNPUBench#13（topk）调用标准 THISTOGRAM wrapper，
  依赖第 1 项，上游修好前 blocked。
• 相关模型侧改动见 LinxISA/LinxCoreModel#20（THISTOGRAM 建模 + ADDTPC + srcRType）。

[liujiang833]

Follow-up: item 4 的 TSLL/TSRL 落地为 tile–tile SIMT 形式，与 ISA 的统一标量移位（v.sll/v.slli）不符，control 仍无法编译。已拆出 #28 跟踪「标量/立即数 block-ISA TSLL/TSRL」。THISTOGRAM（item 1）已在重建 toolchain 上验证通过：gfrun exit 0 / R2=0，gfsim exit 0。

[liujiang833]

补充 item 4 的验收口径（原 #28 并入此处）：TSLL/TSRL 需要的是 统一标量/立即数移位 的 block-ISA 形式 TSLL/TSRL(dst, src, unsigned shamt) → v.slli SrcL, shamt / v.sll SrcL, SrcR（SrcR 标量寄存器，掩码 mod 64/32，见 isa/v0.56/semantics_conventions.json shifts.reg_shifts）。当前 75eaf77 落地的是 tile–tile 逐元素 SIMT (__vec__ src0[i]<<src1[i], <<<>>>) —— ISA 无逐元素 tile 移位指令，且不应走 SIMT 路径。control/hashtable_lookup_simd 用的全是标量量（>>16/>>13/运行期 r），故仍无法编译。TXOR(tile–tile) 不变。

[liujiang833]

澄清（更正前两条评论，请以本条为准）

忽略之前关于「SIMT <<<>>> 路径」的说法 —— 那不是问题（TXOR 等 tile-op wrapper 同样用 __vec__/<<<>>> 内部实现，能正常编译成 tile 指令）。

唯一的问题是 TSLL/TSRL 的移位量操作数形状：

• 需要（tile-op，移位量为标量/立即数）：

  TSLL(Tile &dst, Tile &src, unsigned shamt);   // dst = src << shamt （整 tile 统一移位）
  TSRL(Tile &dst, Tile &src, unsigned shamt);

  对应 ISA v.slli SrcL, shamt（立即数）/ v.sll SrcL, SrcR（标量寄存器，运行期量）。
  shamt 为常量时走 SLLI，运行期变量时走 SLL。

• 当前 75eaf77 落地的是「第二个操作数是 tile」的逐元素形式
  TSLL(dst, src0, src1) → src0[i] << src1[i]。这不符合需求，也没有对应的逐元素
  tile 移位 ISA 指令。

control/hashtable_lookup_simd 的调用全是统一移位：TSRL(tmp, h, 16u)、
TSRL(.., h, 13u)、TSLL(tA, x, r)（r 运行期标量）。TXOR(tile,tile,tile) 不变。