[one-level][normalization] Use tile subviews for RMSNorm binary cache

[wangyuascend-spec]

Change

• Replace the GM level cache selected by GetCacheId with an eight-slot FP32 RowMajor parent tile.
• Read cache slots through TPARTVIEW, merge lower slots in tile, and rebuild/update the selected slot through TileArray plus TASSEMBLY.
• Remove the GM workspace allocation and kernel argument from the rms_norm_binary test.
• Keep the cache exact FP32; no BF16 conversion is introduced.

Baseline

• SuperNPUBench tag: ops-20260908
• Base commit: a3fa598
• gfrun model commit: 07e9c6611270827303e1435e1199bfc410fbec53

Validation

Configuration: DType=__half, G_A=16, G_R=16384, PE_NUM=4, res_check=on.

• Compile and link: PASS.
• gfrun: BLOCKED before output/golden precision comparison.
• Error: gfrun: illegal instruction at 0x0: illegal TMULS operand or descriptor contract

The current ops-20260908 model rejects the RowMajor tile-subview operand/descriptor path used by this implementation, so this PR does not claim a passing accuracy result yet.

[chatgpt-codex-connector]

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: c714c17fcd

ℹ️ About Codex in GitHub

Codex has been enabled to automatically review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

When you sign up for Codex through ChatGPT, Codex can also answer questions or update the PR, like "@codex address that feedback".

[chatgpt-codex-connector]

Avoid unsupported vector tile assembly

In the default G_A=16, G_R=16384 path, execution reaches this assembly before producing any output, but both the fragments and parent are Location::Vec/RowMajor; the current ASL/model supports local B.ASSEMBLE only for Matrix/CUBE layouts, as documented in rowsum_subview_gfrun_issue.md:130-140. Consequently compilation succeeds but gfrun rejects the generated descriptor path with an illegal instruction, matching the failed validation recorded for this commit. Keep the cache in GM or lower it through an ISA-supported representation until Vector/NORM assembly is supported.

Useful? React with 👍 / 👎.

[wangyuascend-spec]

Updated in commit 22a7f6f:

• Cache fragment now uses the native RowSum physical/valid shape: 32x1 FP32 (128 B).
• The parent cache tile is 32x8 FP32 (1 KiB), containing eight native slots.
• Removed cache-only TCOLEXPAND, TROWEXPAND, TROWMAX, and TCOLMAX operations.
• GetCacheId selects TPARTVIEW slots directly; TileArray plus TASSEMBLY rebuilds the parent because borrowed SubTileView is source-only.

Validation with DType=__half, G_A=16, G_R=16384, PE_NUM=4:

• compile/link: PASS
• gfrun: still blocked before accuracy comparison with illegal TMULS operand or descriptor contract

This confirms the removed shape adapters were not the cause of the model failure.