[TROWSUM] physical Columns=1 reduction vector is rejected by Tile alignment check

[wangyuascend-spec]

问题背景

硬件约束已经确认发生变化：TROWSUM 等 row-reduction 接口的输出 descriptor 必须设置：

physical Columns = 1
ValidColumns     = 1
physical Rows    = tile capacity / (Columns * element bytes)

同一 reduction vector 经过 TMULS/TADDS/TRECIP 等操作并作为 TROWEXPANDMUL 的 broadcast source 时，也需要保持 physical/valid Columns=1。

当前 TileOP API 无法表达这个合法 descriptor。对 float row-reduction vector 使用 Cols=1 会被普通 RowMajor 每行 32-byte 对齐的静态断言拒绝。

版本

• Linx-TileOP-API: a795b973020de012fa2a8e2c29b26adfc83d5d29
• SuperNPUBench: ops-20260823 / 244d088fdc56dc0b41cc77b60e03dc974b67ea93
• 用例：one-level normalization/rms_norm/rms_norm_static
• 配置：DType=__half G_A=16 G_R=512 TILE_A=2 TILE_R=512 PE_NUM=1

最小复现

创建 /tmp/repro_rowsum_col1.cpp：

#include <common/pto_tileop.hpp>

using namespace pto;

// TROWSUM result: two valid rows, one physical/valid column.
using row_reduce_t =
    Tile<Location::Vec, float, 2, 1, BLayout::RowMajor, 2, 1>;

void instantiate_row_reduce_tile() {
    row_reduce_t value;
    (void)value;
}

使用包含当前 TileOP API 的 linx toolchain 编译：

export COMPILER_DIR=/path/to/linx_blockisa_llvm_musl/bin
export TILEOP=/path/to/Linx-TileOP-API

"$COMPILER_DIR/clang++" -c -mlxbc -fenable-matrix -std=c++20 \
  -D__linx -DENABLE_TENSOR_INSTR \
  -I"$TILEOP/include" \
  /tmp/repro_rowsum_col1.cpp -o /tmp/repro_rowsum_col1.o

实际结果

编译失败，首个错误为：

include/common/pto_tile.hpp:721:3: error: static assertion failed due to requirement
'BLayout::RowMajor ... Cols * type_traits<float>::bits % (32 * 8) == 0 ...'

in instantiation of template class
'pto::Tile<pto::Location::Vec, float, 2, 1,
           pto::BLayout::RowMajor, 2, 1, ...>'

Cols=1、dtype=float 时一行只有 4 bytes，因此当前检查：

Cols * type_traits<DType>::bits % (32 * 8) == 0

必然失败。随后出现的 GetValidCol()、GetValidRow()、data() 不存在以及 TADDS substitution failure 都是 Tile 类实例化失败产生的连锁错误，不是独立根因。

SuperNPUBench 端到端复现

将 RMSNorm 的 reduction vector 改为：

using tile_v =
    Tile<Location::Vec, float, tA, 1, BLayout::RowMajor, tA, 1>;

然后从 SuperNPUBench/benchmark/one-level-arch/test/kernel/normalization/rms_norm 执行：

make TESTCASE=rms_norm_static DType=__half \
  G_A=16 G_R=512 TILE_A=2 TILE_R=512 PE_NUM=1 \
  COMPILER_DIR=/path/to/linx_blockisa_llvm_musl/bin \
  res_check=on all

会在实例化：

Tile<Location::Vec, float, 2, 1, RowMajor, 2, 1>

时得到同一个 pto_tile.hpp:721 静态断言。

预期行为

TileOP API 应提供符合新 row-reduction contract 的表示和 lowering：

• descriptor 发布 physical Columns=1、ValidColumns=1；
• physical Rows 根据 tile capacity/TSize 和 element size 推导，而不是要求一行本身达到 32-byte；
• TROWSUM -> TMULS/TADDS/TRECIP -> TROWEXPANDMUL 全链路保持同一个合法 reduction-vector shape；
• 普通 matrix tile 的 RowMajor 对齐检查不应误拒绝 row-reduction vector；如果需要新的 layout、carrier type 或显式 API，请在接口中提供并说明用法。

建议验收测试

1. 增加 float/fp16、ValidRows=1/2/4、physical/valid Columns=1 的 Tile 构造编译测试。
2. 增加 TROWSUM -> TMULS/TADDS/TRECIP -> TROWEXPANDMUL 编译测试，并检查生成 descriptor 的 physical Columns 确实为 1。
3. 使用 SuperNPUBench static RMSNorm TILE_A=2 做端到端编译回归。

[wangyuascend-spec]

尝试修复后的结论：底层 Tile 已可表达，carrier 需要保留 TSize 容量

不能直接把旧 carrier：

Tile<Vec, float, tA, 128, RowMajor, tA, 1>

机械改成 Tile<Vec, float, tA, 1, RowMajor, tA, 1>。后者逻辑存储只有 tA * 4B（tA=2 时 8B），既不满足当前对齐规则，也无法编码原来的 tile capacity/TSize。

保留原容量并把 carrier 转成 one-column ColMajor 后可以编译：

using row_reduce_t =
    Tile<Location::Vec, float,
         tA * 128, 1,
         BLayout::ColMajor,
         tA, 1>;

例如 tA=2：

旧 physical shape = 2 x 128 = 256 floats = 1KB
新 physical shape = 256 x 1 = 256 floats = 1KB
new valid shape   = 2 x 1

验证结果：

• dynamic RMSNorm（tA=1）：编译通过，gfrun golden 精度通过，max_abs=0.001953125。
• static RMSNorm（tA=2）：编译通过。
• static 非 RES_CHECK ELF：gfsim 通过。
• static RMSNorm gfrun：仍 R2=0 后精度失败，与 SuperScalarModel #341 的 descriptor publication/preflight 未更新一致。

因此当前首个静态断言不是简单删除即可修复。TileOP 侧剩余建议是提供明确的 row-reduction carrier alias/helper 和测试，避免调用方把 physical Columns=1 误写成丢失 TSize 容量的 skinny RowMajor tile；同时修正/覆盖 TRowExpandMul.cpp 中类似 row x 1 RowMajor 的用法。

[VV0003]

验证结论：TileOP 断言问题已解决

已按当前 active PTO-ISA descriptor 规则复核并验证该问题：

• ISA 对 row-reduction descriptor 的要求是：physical Cols=1、ValidCols=1，physical Rows 根据 TSize / (Cols * element_bytes) 推导；并没有要求所有 descriptor 都满足普通 RowMajor 的“每行 32B”规则。
• TileOP 当前 linx 分支已通过 commit 2e4d695 移除过严的普通 RowMajor 32B 静态断言，因此可以表达 Issue 中的 Cols=1 row-reduction descriptor。
• 使用 Issue 中的最小测例和真实 Linx clang 编译后，已成功生成 Linx 目标 ELF relocatable object：/tmp/repro_rowsum_col1.o。
• TileOP make check 通过：32 个 Python contract tests、C++ syntax-only tests、shell syntax checks 和 git diff --check 全部通过。

之前真实 clang 仍报 pto_tile.hpp 静态断言的原因，是工具链 resource include 中残留旧版 TileOP headers；同步当前 linx 分支的 include/ 后，Issue 测例可以正常编译。

因此，Issue 中由 TileOP alignment assertion 导致的编译失败已解决，当前实现与 active PTO-ISA 的 shape/capacity contract 一致。

本次只验证了 Issue 的最小编译测例并生成 .o；完整 SuperNPUBench kernel 的 link/gfrun 运行属于后续端到端验证，不作为本 issue 的直接阻塞项。