Skip to content

feat(normalization): add tree-reduction static RMS Norm - #144

Draft
wangyuascend-spec wants to merge 3 commits into
PTO-ISA:mainfrom
wangyuascend-spec:feat/rms-norm-simt-static-m-r-tree
Draft

wangyuascend-spec wants to merge 3 commits into
PTO-ISA:mainfrom
wangyuascend-spec:feat/rms-norm-simt-static-m-r-tree

Conversation

@wangyuascend-spec

Copy link
Copy Markdown
Collaborator

Summary

  • add rms_norm_simt_static_m_R_tree for the fixed [512, 8192] 4PE RMS Norm case
  • split R as [16, 32, 16] and process paired outer-R blocks before reduction
  • assemble eight [32, 1] pair-reduction results into one [32, 8] Tile
  • reduce the assembled Tile with TROWSUM followed by TCOLSUM
  • apply gamma using the RMS affine path already present on main
  • add the dedicated test driver and compile.all registration

Base

Validation

  • git diff --check: PASS
  • Test shape: DType=__half, G_A=512, G_R=8192, PE_NUM=4
  • LLVM: 73cbdf34fbd0cc58b263b3da7942ec34ad29323a
  • Linx-TileOP-API: aa3b538614ab60f3a8fd072e9abd708c40c3dcc2

Compilation currently reaches TileOP TROWSUM/TROWSUM_ASS lowering but is blocked by LinxISA/llvm-project#94:

B.IOT tile_t1, mask=1111, last
error: Match Instruction Error!

The earlier TROWSUM physical-shape assertion is resolved by TileOP PR #136. No ELF or gfrun result is available for this new testcase until llvm-project#94 is fixed.

Dependency

@VV0003

VV0003 commented Sep 16, 2026

Copy link
Copy Markdown
Collaborator

The two blockers for this PR are now resolved on the default branches:

  1. LLVM Add four-PE multi_thread operators, migrate matmul to CubeTile, refresh eval/ binaries #94B.IOT tile_t1 Match Instruction Error: fixed by LinxISA/llvm-project#95, merged into dev-llvm15_56 as ae8cd75a5259. The MC parser now accepts the absolute physical Tile register spellings (tile_t1/tile_u1/tile_m1/tile_n1 …) that the register allocator emits for long-lived inline-asm Tile operands. The attached rms-tree.s assembles successfully with this build.

  2. TileOP — reduction _ASS missing session metadata: TROWSUM_ASS/TROWMAX_ASS now emit the destination-only B.ASSEMBLE (INIT/LAST/RegSrc/Offset/ParentSize) after the destination binder. Fixed in Linx-TileOP-API kernels: add matmul_quantize (matmul -> group-32 quantize -> data+scale) #142, already merged into linx (414546f).

Rebuild against current dev-llvm15_56 + linx and this PR should compile through. Note when rebasing the reproduction: the three pair blocks in square_pair_reduce pass pair_index as the second argument of range::assemble_last/middle, which is the range base GPR parameter, not the column offset — with the metadata now actually encoded in B.ASSEMBLE, offsets likely need to be spelled explicitly (e.g. range::assemble_middle<1, N>) so each [32,1] result lands in its own column of the [32,8] parent. Worth double-checking during runtime validation.

XrXie added a commit that referenced this pull request Sep 16, 2026
Replace the three mxquant variants (mxquant [128,64], mxquant_1024_assembly
and mxquant_1024_streaming [128,1024]) with a single [512,256] operator that
matches the dynamic_mx_quant solution's tail OCP shape: four PEs split M
(128 rows each), each 32-row block is quantized one 32-wide MX group at a
time (8 groups/block, 4 blocks/PE), and the eight per-group E8M0 columns are
packed four-at-a-time into two U32 words per row (two full-CELL 128 B scale
stores per row block instead of eight 1/4-full ones).  The byte order stays
plain row-major [512,8] (scale(row, group) = scales[row*8 + group]) and the
test/reference contract is unchanged; the tests are reworked to the same
single variant.

Why max(x,-x) instead of abs: on the current gfrun a unary TEPL op writes its
CUBE_M32 destination densely instead of through the CELL indexer, which
permutes the [32,32] tile (LinxISA/SuperScalarModel issue #678).  TSUB/TMAX
are CUBE-cell aware and max(x,-x) is bit-exact for every input, so the
scale/quant result is identical to an abs-based one.  The removed variants
used TABS and were only correct on other model builds.

Performance: this kernel is TLSU bound.  Each 32-wide group is TLOADed/TSTOREd
with its own strided view into the [32,256] row block, i.e. 64 B payload
bursts at a 512 B stride (32 B at 256 B on the store side), both well under
the 256 B cacheline, so per-access efficiency is burst/256; the non-contiguous
stores additionally pin the SCB.  The fix is one contiguous [32,256] TLOAD
(512 B rows) plus TPARTVIEW subviews for the per-group reductions and
TASSEMBLY/B.ASSEMBLE on the store side, which is blocked today: TPARTVIEW's
SubTileView is not yet a general tile operand (LinxISA/Linx-TileOP-API issue
#144) and gfsim does not model B.ASSEMBLE.  The disabled #if 0 variant in
mxquant.hpp records the blocked contiguous form.

Bandwidth model (from kernels/multi_thread/tlsu_bandwidth_usage.md):
  * one TLSU op costs ~ fixed overhead + bytes/256; load ~15 cyc, store ~31 cyc;
  * only burst >= CACHELINE(256 B), or stride == burst, fills a cacheline; a
    smaller burst costs burst/256 efficiency and still pays a tag/beat per
    cacheline;
  * measured on the same [512,256] MX quant: the contiguous [1,16384] reshape
    is ~1367 cyc total, while the per-group [512,256] strided form is ~12676
    cyc (TLSU Active Wall ~99%, STQ peak 15, SCB 30-46 of 64, store-side
    back-pressure into the frontend);
  * store is dearer than load (31 vs 15) and is pinned by the SCB
    global-visibility window, so cutting store count / making stores
    contiguous has the largest payoff.

Verification (res_check=on; gfrun -s softcore.multiThreadNum=4 -f <elf>):
  payload 131072/131072, scale 4096/4096 bytes, byte-exact;
  input_readback.bin == input.bin; both goldens non-zero;
  perturbing one golden byte is detected at the expected offset;
  the default kernel-only build compiles.
@wangyuascend-spec
wangyuascend-spec force-pushed the feat/rms-norm-simt-static-m-r-tree branch from 4240b5e to 6618d6a Compare September 16, 2026 11:29
@wangyuascend-spec
wangyuascend-spec force-pushed the feat/rms-norm-simt-static-m-r-tree branch from 6618d6a to d32897c Compare September 16, 2026 11:31
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants