[TROWSUM] Columns=1 full RMSNorm path fails on dtype modifier and sub-128B TSize

[wangyuascend-spec]

Summary

After applying the updated hardware contract from #29 (the output of TROWSUM uses physical Columns = 1 and ValidCols = 1), the full SuperNPUBench RMSNorm kernel still cannot compile.

The minimal type-instantiation test in #29 (comment) does not instantiate tile operations, so it does not cover the failures below.

Versions

• Linx-TileOP-API: dacedc2
• llvm-project / installed clang: 76044f43610027d016191453b9a60bbdc329e3e9
• clang version: 15.0.4 (linx64v5-musl-local 76044f...)
• SuperNPUBench: tag 20260823 plus the RMSNorm change below

Required RMSNorm change

The row-reduction output is declared with one physical column:

Tile<Vec, float, tA, 1, RowMajor, tA, 1> row_sum;

For the reproducer parameters below, this becomes Tile<float, 2, 1, RowMajor, 2, 1>.

Reproduction

From:

SuperNPUBench/benchmark/one-level-arch/test/kernel/normalization/rms_norm

run:

make TESTCASE=rms_norm_static \
  DType=__half G_A=16 G_R=512 TILE_A=2 TILE_R=512 PE_NUM=1 \
  COMPILER_DIR=/home/wangyu/Code/PtoNpu/linx-toolchain-build/output/linx_blockisa_llvm_musl/bin \
  res_check=on all

Actual result

Compilation fails in TileOP inline assembly. Two concrete problems are visible:

1. %D dtype modifier is not printed

include/jcore/template_asm.hpp emits, for example:

BSTART.TLOAD %D[SrcType]

The installed compiler reports an invalid operand, and the rendered instruction has no dtype:

template_asm.hpp:1781: invalid operand for instruction
BSTART.TLOAD

LinxV5AsmPrinter::PrintAsmOperand at LLVM commit 76044f4 handles %S and %Z, but does not handle %D.

2. A narrow Rows x 1 tile has no valid physical TSize

Tile<float, 2, 1, ...> currently has StorageBytes = 8. However, the active B.IOT tile size range starts at 128B. TilesizeCode is therefore invalid and the emitted assembly contains an empty size:

B.IOT ... ->u<>

This also contradicts the stated hardware contract that physical rows are inferred from TSize / (physical Columns * element bytes): with physical Columns fixed to 1, TileOP still needs a representable physical TSize (presumably at least 128B), while retaining logical/valid rows as 2.

Expected result

The complete TLOAD -> TROWSUM -> subsequent tile ops/TSTORE RMSNorm path should compile when the reduction output has physical Columns = 1 and ValidCols = 1.

In particular:

1. TileOP and the supported LLVM revision should agree on how dtype operands such as %D are rendered.
2. Non-cube tiles smaller than 128B should receive a legal physical carrier/TSize without changing their logical shape and valid dimensions.

Why this is different from #29

#29's original alignment assertion was removed, but the comment's verification only compiled a minimal tile declaration to an object file. It did not instantiate operations that require dtype rendering or B.IOT TSize encoding. The full normalization kernel exercises both and remains blocked.

[VV0003]

Fixed in:

• TileOP: fae5b1ab4e8cfcdfcb2e891a439e64a3230ab3d0 on linx
• LLVM: 118d038ab435c992275decc4114e83ef9e91561e on dev-llvm15_56

The issue had two independent causes:

1. The LinxV5 LLVM backend did not implement the %D inline-assembly operand modifier. TileOP expressions such as BSTART.TLSU TLOAD, %D[SrcType] therefore could not emit the canonical data-type spelling. LLVM now maps the target DataType encoding to the canonical assembly name, including FP32, BF16, integer types, and DTYPE_NONE.

2. An ordinary Tile whose logical payload was smaller than 128 B could produce an invalid TilesizeCode and empty ->u<> syntax. PTO-ISA defines SizeCode=1 as the minimum Local destination capacity of 128 B. TileOP now rounds the physical carrier capacity up to the next legal SizeCode while preserving the logical/valid shape. For example, Tile<float, 2, 1> remains logically 2 x 1, but uses a 128 B physical carrier with SizeCode=1. This matches the TROWSUM contract: the destination is source.ValidRow x 1, with capacity-derived physical rows.

Regression coverage includes:

• LLVM %D CodeGen checks for FP32 and BF16;
• a TileOP static contract for the narrow 2 x 1 Tile;
• a narrow TLOAD -> TROWSUM -> TSTORE RMSNorm path.

Validation passed:

• LLVM llc/clang build and inline-asm-dtype-modifier.ll;
• TileOP make check (32 contract tests plus C++ syntax checks);
• PTO 0.58.3 MC/disassembly/negative contract verification.

The two fixes are consistent with the active PTO-ISA operand and SizeCode contracts, so the compiler/API failures reported here are resolved.