Drop unused GM workspace from rms_norm_binary

[wangyuascend-spec]

Summary

• Keep the binary-accumulation cache in a Local tile array (updateTile[]) instead of GM workspace.
• Remove the unused workspace argument from both dyn and static rms_norm_binary entries, and drop the matching host buffers in the test mains.

Test plan

• make TESTCASE=rms_norm_binary DType=__half compiles
• make TESTCASE=rms_norm_binary_static DType=__half G_A=1 G_R=8192 TILE_A=1 TILE_R=1024 POW_R=4096 compiles
• gfrun + precision compare for dyn/static [1,8192] if CI/reviewer wants a functional re-check

Made with Cursor

[chatgpt-codex-connector]

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: 465fd5dd7f

ℹ️ 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]

Keep float vector tiles at the 512-byte minimum

With kWsCols = 32, every tile_v used by TROWSUM, TADD, TMULS, TRECIP, and TROWEXPANDMUL has only 128 bytes of physical capacity. The one-level TileOP normalization kernels require a float scalar/broadcast strip of at least 128 columns (512 bytes), as documented and implemented in rms_norm/rms_norm_pto.hpp and both group-norm-grad kernels; therefore the dynamic and static rms_norm_binary paths can compile but fail during functional execution with an invalid active size/Match Instruction error. Retain 128 physical columns while keeping ValidCol=1.

Useful? React with 👍 / 👎.