Skip to content

feat(normalization): add 128-byte static RMS Norm tile - #139

Open
wangyuascend-spec wants to merge 1 commit into
PTO-ISA:mainfrom
wangyuascend-spec:feat/rms-norm-static-128b
Open

wangyuascend-spec wants to merge 1 commit into
PTO-ISA:mainfrom
wangyuascend-spec:feat/rms-norm-static-128b

Conversation

@wangyuascend-spec

Copy link
Copy Markdown
Collaborator

Summary

  • Add a standalone rms_norm_static_128b kernel for the fixed [512, 8192] fp16, 4PE case.
  • Use a physical and valid compute Tile shape of [1, 32]; the FP32 reduction source is therefore 128 bytes.
  • Keep the current RegBase reciprocal/sqrt, Newton, and compensated-residual formula with kernel-local epsilon 1e-6.
  • Do not add CMP/SEL handling for zero, infinity, or exceptional values.
  • Add the matching test source and register it in Makefile and compile.all.

Tile layout

  • FP32 src/squared/dst Tile: [1, 32], 128 bytes
  • FP16 load/store Tile: [1, 32], 64 bytes
  • FP32 row-reduction Tile: [1, 1], 4 bytes
  • Each row is processed in 256 reduction strips.

Validation

Environment:

  • SuperNPUBench base: ops-20260914 / 085b3bc
  • Compiler: 0908 issue100-patched toolchain
  • gfrun: SuperScalarModel 6a9b739c
  • PE count: 4

Result:

  • rms_norm_static_128b, [512, 8192], fp16: PASS
  • gfrun: R2=0
  • Total instructions: 9,713,621
  • Build parallelism: -j1

@chatgpt-codex-connector chatgpt-codex-connector Bot left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 3aabfd6cb3

ℹ️ 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".

make TESTCASE=rms_norm_static COMPILER_DIR="$COMPILER_DIR" DType=__half G_A=512 G_R=8192 PE_NUM=4 diss

# Fixed-shape static 4PE counterpart with a 128-byte FP32 compute Tile.
make TESTCASE=rms_norm_static_128b COMPILER_DIR=${COMPILER_DIR} DType=__half G_A=512 G_R=8192 PE_NUM=4 diss

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

P2 Badge Quote the compiler path in the new build command

When COMPILER_DIR contains whitespace, the unquoted expansion is split into multiple arguments, truncating the Make variable assignment and treating the remaining path component as an additional target, so this new testcase fails to build even though the preceding testcases accept the same path. Quote the expansion as COMPILER_DIR="$COMPILER_DIR", consistent with the two commands above.

Useful? React with 👍 / 👎.

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.

1 participant