[Bug] TileOP-API TileDType 硬编码 4KB 导致 tile buffer 溢出

[lvhao7896]

[Bug] TileOP-API TileDType 硬编码 4KB 导致 tile buffer 溢出

摘要

Linx-TileOP-API 仓库 commit f9a30a69 将 Tile::TileDType 从按维度派生的 DType tile_size(Rows*Cols) 改为硬编码的 int32_t __attribute__((ext_vector_type(1024)))（恒定 4096 字节 = 4KB）。后端 calculateVCallSizeMask 据此生成 B.IOT/B.IOS 指令的 TSize 字段，导致所有实际数据 > 4KB 的 tile 在 gfrun 运行时因 tile buffer 容量不足而断言失败。

复现环境

组件	版本
gfrun	bb91c18e（SuperScalarModel，feat/pto-v058-adaptation 分支）
工具链	linx-toolchain-build-latest，clang 15.0.4，llvm 86959776b
工具链源码分支	temp/shared-tload-integration-20260811
Linx-TileOP-API	commit a96d4e9（a96d4e9 分支，2026-08-17）
运行方式	/Users/blacktraker/Programming/gitproj/DV4/SuperScalarModel/bin/gfrun -t 1 -f <elf> -s softcore.multiThreadNum=4

问题定位

错误代码

文件: linx-toolchain-build-latest/src/Linx-TileOP-API/include/common/pto_tile.hpp
行号: 703

#ifdef __linx
  using TileDType = int32_t __attribute__((ext_vector_type(1024)));   // ← 罪魁祸首
#else
  using TileDType = DType[Rows * Cols];
#endif

TileDType 是 tile 的 LLVM 向量承载类型，后端从它的字节大小推导 TSize。此处写死 1024 个 int32 元素 = 4096 字节 = 4KB，完全无视模板参数 Rows_、Cols_、DType。

引入此 bug 的 commit

commit f9a30a69f17e6a0ec204074d2ce525d77f39684c
Author: Kevin Zhou <zhoubot@gmail.com>
Co-authored-by: RuoyuZhou <ruoyu.zhou@hisilicon.com>
Date:   Tue Aug 11 18:57:39 2026 +0000

    tileop: align target headers with active contracts (#3)

diff（pto_tile.hpp）：

 #ifdef __linx
-  using TileDType = DType tile_size(Rows *Cols / (sizeof(DType) * 8 / type_traits<DType>::bits));
+  using TileDType = int32_t __attribute__((ext_vector_type(1024)));
 #else
   using TileDType = DType[Rows * Cols];
 #endif

同一个 commit 还新增了正确的 LogicalTileBytes/TilesizeCode 常量（从 Rows*Cols*bits 计算），但这些常量从未被用来构造 TileDType——正确计算成了死代码。

机制分析

tile_size 是编译器内置宏

-mlxbc 自动引入编译器内置头文件 linx_blkc.h，其中定义：

// lib/clang/15.0.4/include/linx_blkc.h:6
#define tile_size(n) __attribute__((ext_vector_type(n)))

因此 baseline 代码 DType tile_size(Rows*Cols) 展开为 DType __attribute__((ext_vector_type(Rows*Cols)))，向量长度随 tile 维度变化。

完整责任链

pto_tile.hpp:703  TileDType = <1024 x i32> (固定 4KB, f9a30a69 后)
       │
       │  Tile::data() 返回 TileDType&，作为内联汇编输出操作数 "=Tr"(dst.data())
       ▼
template_asm.hpp  每条 PTO 指令（TLOAD/TMATMUL/TEXPANDS…）
       │  asm volatile("...B.IOT ... ->%[Dst]<%Z[TileSize]>"
       │              : [Dst]"=Tr"(dst.data())           // 类型 = <1024 x i32>
       │              : [TileSize]"i"(tile_type_traits<TileDType>::TilesizeCode)  // sizeof=4096 → 6
       ▼
LinxV5ISelLowering.cpp:1100  calculateVCallSizeMask(EVT Type)
       │  SizeBytes = Type.getFixedSizeInBits()/8 = 32768/8 = 4096
       │  return Log2(4096) - 6 = 12 - 6 = 6
       ▼
B.IOT 指令 Inst{11-9} = TSize = 6  →  反汇编打印 <4KB>
       ▼
gfrun: encodedOutputBytes = 128 << (6-1) = 4096  →  8KB 数据塞不进 4KB → 断言失败

baseline（改前）的展开效果

tile_size 宏展开后，公式 Rows*Cols / (sizeof(DType)*8/type_traits<DType>::bits) 对标准类型恒等于 Rows*Cols：

tile	DType	Rows×Cols	baseline 展开	字节大小	TSize	结果
fa shared (128×16)	float	2048	<2048 x float>	8192 = 8KB	7	✅
reduction (32×64)	int32	2048	<2048 x i32>	8192 = 8KB	7	✅
reduction (32×64)	__half	2048	<2048 x half>	4096 = 4KB	6	✅
—（f9a30a69 后，全部写死）	int32	1024	<1024 x i32>	4096 = 4KB	6	❌（>4KB 的 tile 全挂）

影响范围

失败的算子（实际 tile > 4KB）

算子	tile 维度	dtype	实际字节	容器	断言	状态
fa (multi_thread)	128×16	FP32	8KB	4KB	cooperative TMATMUL, SoftCore.cpp:575	❌ FAIL
reducemax_col	32×64	int32	8KB	4KB	TLOAD, AccumulateBlockInfo.cpp:74	❌ FAIL
reducemax_row	32×64	int32	8KB	4KB	TLOAD	❌ FAIL
reducesum_col	32×64	int32	8KB	4KB	TLOAD	❌ FAIL
reducesum_row	32×64	float	8KB	4KB	TLOAD	❌ FAIL

通过的算子（实际 tile ≤ 4KB，碰巧相等）

算子	tile 维度	dtype	实际字节	容器	状态
matmul_shared (multi_thread)	32×32	FP32	4KB	4KB	✅ PASS (R2=0)
reducemax_col	32×64	__half	4KB	4KB	✅ PASS (R2=0)
matmul (single)	16×16	FP32	1KB	4KB	✅ PASS (max_abs=8.94e-8)

断言详情（fa multi_thread 为例）

• 触发块: B17，BPC 0x1144e，BSTART.CUBE TMATMUL FP32
• 最后一条指令: 0x1146e B.IOT mask=1111, last, ->m<4KB>（T3 执行完后触发）
• B.DIM: lb0=16(m), lb1=16(n), lb2=128(k)
• B16 TLOAD（0x11438）: 向 S1<4KB> 装填 128×16 的 FP32 tile，需 8KB，但 tile buffer 仅 4KB
• 失败条件:

  // SoftCore.cpp:575 — cooperative TMATMUL assertion
  rightShared.data.size() >= ((k - 1) * rightCol + n) * elementBytes
  // 4096 >= ((127×16)+16)×4 = 8192  →  FALSE

  另：rightInfo->validRow == k（128）也可能失败，因为 4KB tile 只能装 64 行（4096÷(16×4)=64）。

时间线

日期	事件	TileOP-API 版本	TileDType	fa 结果
2026-06-22	baseline 91e3d7a 创建	DType tile_size(Rows*Cols)	维度派生 ✅	✅
2026-08-04	旧工具链构建，安装 baseline headers	正确	维度派生 ✅	✅ fa 能过
2026-08-11	f9a30a69 Kevin Zhou 改成 ext_vector_type(1024)	写死 4KB	恒定 4KB ❌	❌
2026-08-14	最新工具链构建（含 bug）	写死 4KB	恒定 4KB ❌	❌
2026-08-17	旧工具链 TileOP-API 也被重建到 a96d4e9	写死 4KB	恒定 4KB ❌	❌ 两个工具链都过不了

fa 之前能跑过，是因为 Aug 4 构建的旧工具链安装的是 baseline 版本（91e3d7a）的 TileOP-API，TileDType 大小正确派生自 tile 维度。Aug 11 的 f9a30a69 引入 bug 后，Aug 14 构建的新工具链含 bug；Aug 17 旧工具链 headers 被重建到最新版，现在两个工具链都受影响。

根因总结

f9a30a69 的提交信息是 "align target headers with active contracts"。改动原因可能是将 TileOP-API 头文件与编译器实际支持的能力对齐——但 tile_size 宏在 linx_blkc.h 中一直存在且工作正常，该替换是不必要的且破坏性的。

讽刺的是，同一个 commit 还新增了正确的 LogicalTileBytes 和 TilesizeCode 常量（从 Rows*Cols*bits 计算正确的 size class），但这些常量从未被用于构造 TileDType——正确计算结果成了死代码。

修复建议

方案 A：回退 TileDType（最小改动，推荐）

将 pto_tile.hpp:703 回退到 baseline 形式：

#ifdef __linx
  using TileDType = DType tile_size(Rows *Cols / (sizeof(DType) * 8 / type_traits<DType>::bits));
#else
  using TileDType = DType[Rows * Cols];
#endif

tile_size 宏已在编译器内置头文件 linx_blkc.h 中定义（#define tile_size(n) __attribute__((ext_vector_type(n)))），展开后 TileDType 为 DType __attribute__((ext_vector_type(Rows*Cols)))，字节大小 = Rows*Cols*sizeof(DType)，后端 calculateVCallSizeMask 会正确计算 TSize。

验证：static_assert 确认展开后各类型大小正确：

• float tile_size(2048) → <2048 x float> → 8192B = 8KB → TSize=7 ✅
• int32_t tile_size(2048) → <2048 x i32> → 8192B = 8KB → TSize=7 ✅
• _Float16 tile_size(2048) → <2048 x half> → 4096B = 4KB → TSize=6 ✅

方案 B：后端增加维度校验（防御性）

在 LinxV5ISelLowering.cpp 的 calculateVCallSizeMask 调用点（lowerTLoad/lowerTemplateBLK* 等），从 B.DIM 操作数（lb0/lb1/lb2）× 元素类型重新校验 TSize，而非仅依赖 Op.getValueType()。

注意事项

• 后端要求 tile 字节数必须是 2 的幂且在 128B~8KB 范围内（calculateVCallSizeMask 会 report_fatal_error 否则）。方案 A 回退后需确保所有 tile 维度满足此约束（当前 fa 128×16=8KB、reduction 32×64=8KB 均满足）。
• pto_tile.hpp 的 #else 分支（非 __linx 平台）使用 DType[Rows*Cols]（host 数组形式），不受影响。

涉及文件清单

文件	说明
Linx-TileOP-API/include/common/pto_tile.hpp:703	Bug 所在：TileDType 硬编码
Linx-TileOP-API/include/jcore/template_asm.hpp	PTO 指令内联汇编，使用 TileDType 作为输出操作数和 TileSize 立即数
Linx-TileOP-API/include/jcore/type.hpp:84-121	tile_type_traits 从 sizeof(TileDType) 计算 TilesizeCode
llvm-project/llvm/lib/Target/LinxV5/LinxV5ISelLowering.cpp:1100	calculateVCallSizeMask：从 EVT 字节大小算 TSize
llvm-project/llvm/lib/Target/LinxV5/LinxV5InstrInfo.td:2765-2979	B.IOT/B.IOS 指令编码，Inst{11-9} = TSize
linx_blkc.h:6	#define tile_size(n) __attribute__((ext_vector_type(n)))（编译器内置宏）
SuperScalarModel/emulator/SoftCore.cpp:575	cooperative TMATMUL 断言（fa 失败点）
SuperScalarModel/isa/AccumulateBlockInfo.cpp:74	TLOAD 断言（reduction 失败点）

[VV0003]

已修复并推送到 linx 分支：

• commit: 8b2ee78 ([tileop-api] Restore logical Tile register size)
• 恢复 f9a30a69 之前按 Rows / Cols / DType 派生 TileDType 的实现，不再固定为 <1024 x i32> / 4KB；
• 同步更新回归检查，验证 FP32 2x1024 和 S32 32x64 的 carrier TSize 均为 8KB，FP16 2x1024 为 4KB；
• 更新 ReinterpretedTileView 注释，明确 reinterpret 保留 source carrier，而不是假定 carrier 固定为 4KB。

验证结果：

1. Linx 编译器 -mlxbc 可正确处理 float tile_size(2048) / int tile_size(2048)，sizeof 均为 8192 字节；
2. 使用最新仓库头文件编译 TCvt.cpp 成功；
3. llvm-objdump 中实际生成了 3 处：

B.IOT mask=1111, last, ->t<8KB>

说明后端已从恢复后的 TileDType 正确得到 TSize=7，不再把 8KB tile 编码为 4KB。

仓库完整 test_v058_engine_contract.py 仍有 9 个既有失败项（canonical asm/旧接口清理等），与本次 TileDType 回退无关；本次新增的 carrier 定向测试通过。