[PTO v0.58][TLOAD/TSTORE] B.IOR 行跨度元素与字节单位不一致(基于tag:20260823)

[luguo23187]

验证环境

tag：20260823

LLVM/Clang:       611105f2be11fab9a8ef20bd02b740f2c5d786b3
Linx-TileOP-API:  a795b973020de012fa2a8e2c29b26adfc83d5d29
SuperScalarModel: a5dca25a5a6802d047573ae71cf39e2615d5356b
Target:           LinxV5 / PTO v0.58
Data type:        FP16

问题描述

PTO v0.58 的 B.IOR 为 TLOAD/TSTORE 提供 GM 基地址和行跨度。新版
SuperScalarModel 从提交 2d467114 开始按照字节解释该跨度：

// SuperScalarModel/emulator/engine/TMAEngine.cpp
stride = block->hasBIOR ? tCopy.stride
                        : tCopy.totalCol * eleSize;
srcAddrRow = srcAddrStart + i * stride;

但 Linx-TileOP-API/include/jcore/template_asm.hpp 的通用 Local TLOAD
仍直接传入 GetStride(3)：

template <is_tile_data_v tile_shape, is_global_data_v gm_shape>
void TLOAD(tile_shape &dst, gm_shape &src) {
  // ...
  asm volatile(
    // ...
    "B.IOR [%[s0],%[GmStride]], []\n"
    // ...
    : [s0]"r"(src.data()),
      // ...
      [GmStride]"r"(src.GetStride(3))
    : "memory");
}

GetStride(3) 是按元素计数的行跨度。对于 FP16 [Rows, Cols] 矩阵，
连续行的正确字节跨度应为：

Cols * sizeof(__half) = Cols * 2 bytes

当前模板却把 Cols 直接放入 B.IOR。model 按字节执行后，下一行从
Cols 字节处开始，而不是从 Cols * 2 字节处开始。

同文件中的通用 TSTORE 也使用相同写法，因此输出为多行矩阵时存在
同类风险。

已有正确实现可作为参照

include/jcore/TLoadBackend.hpp 的 NORM 路径已经进行字节换算：

(src.GetStride(3) *
 type_traits<typename gm_shape::DType>::bits + 7) / 8

include/jcore/TStoreBackend.hpp 也采用相同模式：

(dst.GetStride(3) *
 type_traits<typename gm_shape::DType>::bits + 7) / 8

这表明问题不是 model 对字段单位理解错误，而是不同 TileOP API 路径没有
统一遵守同一个字节 stride 契约。

最小复现

使用一个 FP16 矩阵乘即可稳定观察输入行跨度错误：

#include <common/pto_tileop.hpp>

using namespace pto;

constexpr int M = 32;
constexpr int N = 32;
constexpr int K = 64;

using GmA = global_tensor<__half, RowMajor<M, K>>;
using GmB = global_tensor<__half, RowMajor<K, N>>;
using GmC = global_tensor<float, RowMajor<M, N>>;

using TileA = TileLeft<__half, M, K>;
using TileB = TileRight<__half, K, N>;
using TileC = Tile<Location::Vec, float, M, N, BLayout::RowMajor>;

TileA a;
TileB b;
TileC c;

TLOAD(a, gm_a);
TLOAD(b, gm_b);
TMATMUL(c, a, b);
TSTORE(gm_c, c);

建议令 A 为前 32 行单位矩阵，并令：

B[k,n] = n * 10 + k

数学上 C[0,n] 应等于 B 的第 0 行：

0, 10, 20, 30, 40, 50, ...

问题 ELF 中，FP16 右矩阵的 B.IOR 为 32；model 将其解释为 32 字节，
而正确行跨度应为 32 * 2 = 64 字节。详细 trace 会显示相邻逻辑行从
半行位置开始读取。

同样，左矩阵 [32,64] 的正确行跨度应为 128 字节，但问题 ELF 只编码
64。该现象与 QSMLA 新版环境下第一轮 MM1/第二轮 P@V 同时产生大范围精度
偏差相符。

QSMLA 中的复现现象

QSMLA 基线规格：

B=1, S1=64, S2=128, N1=1, N2=1, D=512
Tm=32, Tk=32, Td=64
dtype=FP16

旧 ELF 在新版 SuperScalarModel 中可以正常结束，R2=0，但精度比较为：

passed   = 897/32768 (2.737427%)
failed   = 31871
max_abs  = 0.954180509
mean_abs = 0.195974510
nan_npu  = 0

该结果不是 NaN、softmax 状态合并或 QSMLA BSND 地址公式导致的。最早共同
失败点是 Local CUBE 前的矩阵输入读取；最小矩阵乘用例进一步把异常缩小到
TLOAD 的 GM 行跨度。

推荐修正

通用 TLOAD 应将元素 stride 转换为字节：

-      [GmStride]"r"(src.GetStride(3))
+      [GmStride]"r"(
+          (src.GetStride(3) *
+           type_traits<typename gm_shape::DType>::bits + 7) / 8)

通用 TSTORE 应采用对应修正：

-      [GmStride]"r"(dst.GetStride(3))
+      [GmStride]"r"(
+          (dst.GetStride(3) *
+           type_traits<typename gm_shape::DType>::bits + 7) / 8)

如果同一文件还有列主序或其他使用 GetStride(4) 的 GM 路径，也应按相同
规则审计，但 Issue 的最小修复和回归范围可以先限定为当前可复现的通用
Local NORM TLOAD/TSTORE。

不建议在 model 中恢复以下旧行为：

srcAddrRow = srcAddrStart + i * stride * eleSize;

这样会使已经正确传入字节 stride 的 TLoadBackend/TStoreBackend 路径再次
出错，并造成同一条 B.IOR 指令的单位依赖调用来源。

建议回归测试

至少增加以下覆盖：

用例	目的
FP16 NORM TLOAD, 32×64	验证 2 字节元素的 GM 行跨度换算
FP32 NORM TLOAD, 16×32	验证 4 字节元素，不把问题限定为 FP16
FP16 NORM TSTORE, 32×32	验证输出 GM 行跨度
UINT8 NORM TLOAD/TSTORE	验证 1 字节类型行为不回退
带 padding 的 GM tensor	验证 stride 大于 validCol 时仍按字节寻址
QSMLA FP16 小规格	验证两次 CUBE 路径和最终 golden 精度

测试不能只覆盖 UINT8。UINT8 的元素 stride 与字节 stride 数值相同，无法
暴露该问题，这也是问题在 HiF8/UINT8 类用例中可能被隐藏、而 FP16 适配中
稳定出现的原因。

Issue 核心诉求

1. 请确认 PTO v0.58 TLOAD/TSTORE 的 B.IOR GM row stride 单位为字节。
2. 请统一 template_asm.hpp 与 TLoadBackend.hpp/TStoreBackend.hpp 的单位。
3. 请修正通用 Local TLOAD/TSTORE，在写入 B.IOR 前把元素 stride 换算成字节。
4. 请审计同文件中所有将 GetStride(3/4) 直接传给 GM 搬运类 B.IOR 的接口。
5. 请增加 FP16、FP32 和带 padding stride 的编译加执行回归；不能只使用 UINT8。
6. 请在验证说明中强调必须重新安装 TileOP API 并重新生成 ELF。