fix(rms_norm): move PE partitioning into kernel - #84
wangyuascend-spec wants to merge 17 commits into
Conversation
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: d387ec0991
ℹ️ 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".
|
|
||
| // The runtime tiling is host data, so divisibility cannot be a static | ||
| // assertion. Invalid configurations do no work rather than overlap rows. | ||
| if (globalA <= 0 || globalA % peNum != 0 || peA < tile_a) { |
There was a problem hiding this comment.
When tile_a exceeds peA (for example, global A 16, peNum 4, and tile_a 8), this guard returns on every PE and leaves the entire output unwritten. Before this change, that configuration reached the sole-tail call with active_a=4, and the tail path below still supports that case, so the peA < tile_a rejection unnecessarily regresses valid dynamic tilings.
Useful? React with 👍 / 👎.
| // ValidCol=1; Cols=128 → 512B/row (TileOP IsValidActiveSize / TSize=1..7). | ||
| using tile_v = Tile<Location::Vec, float, tA, 128, BLayout::RowMajor, -1, 1>; | ||
| // Row-reduction output and row-broadcast input use physical Columns=1. | ||
| using tile_v = Tile<Location::Vec, float, tA, 1, BLayout::RowMajor, -1, 1>; |
There was a problem hiding this comment.
For any dynamic tiling with 1 < tile_a <= peA, rms_norm_tile constructs these vector tiles with active_a valid rows, but this type now has physical capacity tA * 1 = 1 float. That violates the capacity invariant stated immediately above the tile declarations and leaves TROWSUM plus the Newton intermediates without storage for rows after the first; the previous 128-column carrier provided that spare capacity. Retain enough physical capacity or make the physical row count cover the maximum active row count.
Useful? React with 👍 / 👎.
|
Fixed the 4PE res_check failure in commit 82829f9. The precision harness now uses shared static buffers, limits host file I/O to PE0, waits for input readiness, and waits for all PE completion flags before writing output. Verified with softcore.multiThreadNum=4: dynamic and static RmsNorm both finish with R2=0 and pass accuracy (8192 elements, max_abs=0.001953125, MSE=4.747562343254685e-10; atol/rtol=0.02, mse_tol=0.001). Single-PE dynamic precision regression also passes with the same metrics. Fixes the SuperNPUBench-side problem reported in #91. |
|
GroupNormGrad dynamic path has been updated to kernel-owned 4PE partitioning in commit 798f2be. Key points:
Validation (DType=__half, N=2, C=16, G=4, HxW=16, TILE_HW=8):
Compilation was run with -j1 to stay within the ~8 GiB WSL memory budget. |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 798f2be1b5
ℹ️ 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".
| const uint32_t pe = read_pe_id(); | ||
| kPeBarrier[pe] = phase; | ||
| for (int participant = 0; participant < peNum; ++participant) { | ||
| while (kPeBarrier[participant] < phase) { |
There was a problem hiding this comment.
When group_norm_grad is invoked more than once in the same process with peNum > 1, every barrier slot retains phase 2 from the preceding invocation, so both waits in the next invocation immediately pass. A faster PE can consequently read ds/db or c2_buf/c3_buf before slower PEs have refreshed them, producing nondeterministic gradients; use invocation-specific generations or reset and synchronize the slots before phase 1.
Useful? React with 👍 / 👎.
| ifneq ($(PE_NUM),1) | ||
| START_FILE = $(ROOT)/test/common/_start_multipe.s |
There was a problem hiding this comment.
When TESTCASE=group_norm_grad_static is built with PE_NUM set above 1 and run with that many PEs, this unconditional testcase-level selection starts main on every PE, but the static source calls the unpartitioned compile-time group_norm_grad overload and does not even receive a PE_NUM define. Every PE therefore computes the complete workspace and outputs concurrently, causing data races and potentially corrupt gradients; select this startup only for the dynamic multi-PE target or add equivalent partitioning to the static entry.
Useful? React with 👍 / 👎.
|
Dynamic group_norm_grad_1d has been updated to kernel-owned 4PE partitioning in commit 899fea1. Scope:
Validation (__half, N=8, C=64, G=8, TILE_D=-1):
|
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 899fea118f
ℹ️ 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".
| START_FILE ?= $(ROOT)/test/common/_start.s | ||
| CC_LINK += -nostartfiles $(START_FILE) |
There was a problem hiding this comment.
For the documented RMSNorm PE_NUM=4 builds, the RMSNorm Makefile never overrides this default, and make -n shows both dynamic and static targets still linking _start.s. The newly added _start_multipe.s explains that the model gives every PE the same SP; consequently, simultaneous calls into these stack-using kernels can overwrite one another's frames, producing corrupt or nondeterministic four-PE runs. Make the RMSNorm multi-PE targets select the per-PE stack startup as the GroupNorm targets do.
Useful? React with 👍 / 👎.
| ifneq ($(PE_NUM),1) | ||
| START_FILE = $(ROOT)/test/common/_start_multipe.s |
There was a problem hiding this comment.
When TESTCASE=group_norm_grad_1d_static is built with PE_NUM>1, this unconditional selection makes every PE enter main, but the static source invokes the unpartitioned compile-time overload and is not compiled with -DPE_NUM. Thus every PE executes the entire test independently, including concurrent RES_CHECK I/O to the same files; restrict the multi-PE startup to the dynamic testcase or add static partitioning and coordinated I/O.
Useful? React with 👍 / 👎.
| static volatile uint32_t *const kPeBarrier = | ||
| reinterpret_cast<volatile uint32_t *>(0x30000); |
There was a problem hiding this comment.
When the multi-PE overload is used in a bare-metal build, every barrier access targets 0x30000, but the repository's bare-metal memory map starts RAM at 0x60000000 and firmware at 0x70000000; no region or linker reservation covers this address. The first kPeBarrier[pe] store can therefore fault or access device/undefined space instead of synchronizing, so the barrier storage should come from allocated workspace or a linker-defined shared-memory symbol.
Useful? React with 👍 / 👎.
| const uint32_t tid = get_thread_idx(); | ||
| const int64_t pe_offset = static_cast<int64_t>(tid) * peA * gR; |
There was a problem hiding this comment.
If the runtime starts more PEs than the kernel's peNum template argument, an extra PE computes pe_offset from its unrestricted thread ID and advances x and out beyond the globalA rows before processing peA more rows. This causes out-of-bounds reads and writes—for example, thread 4 in rms_norm<dtype, 4> starts exactly at the end of the buffers—so return when tid >= peNum, as the new GroupNorm entry points already do.
Useful? React with 👍 / 👎.
|
Completed the missing normalization 4PE coverage in commit 7ed74eb and reran the authoritative accuracy workflow with the ops-20260828-pinned environment. Tested revisions:
Coverage added:
All six PE4 precision ELFs build successfully with -j1. Results:
Oracle: operator generators plus external compare scripts, atol=0.02, rtol=0.02, MSE tolerance=0.001. The four GroupNormGrad cases fail before output comparison, so they are MODEL_FAIL, not ACCURACY_FAIL. Logs are retained locally under /tmp/norm--build.log, /tmp/norm--gfrun.log, and /tmp/norm-*-compare.log. |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 7ed74eb38e
ℹ️ 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".
| static_cast<size_t>(C) * sizeof(dtype)); | ||
| writeBinaryFile(CHK_DIR "/dbeta.bin", (uint8_t *)dbeta, | ||
| static_cast<size_t>(C) * sizeof(dtype)); | ||
| kernel_done[tid] = 1; |
There was a problem hiding this comment.
In a RES_CHECK run started with more PEs than the compiled PE_NUM, the kernel correctly returns for tid >= PE_NUM, but the launcher then indexes the PE_NUM-element array with that same out-of-range tid. For tid == PE_NUM, this can overwrite the adjacent output_written flag, allowing workers to exit before PE 0 writes the results; larger IDs corrupt other shared state. Check tid before the test synchronization code (the same sequence also appears in the 1D and RMSNorm launchers).
Useful? React with 👍 / 👎.
| if (tile_hw <= 0 || tile_hw > tCap || HxW > tCap || D > tCap || | ||
| tid >= static_cast<uint32_t>(peNum)) { | ||
| return; |
There was a problem hiding this comment.
When either GroupNorm entry point is instantiated with PE_NUM <= 0, this unsigned comparison does not safely validate the template argument: zero silently makes every PE return without producing gradients, while a negative value converts to a large unsigned bound and lets the subsequent nc += peNum loops walk into negative indices and access memory before the tensors. Add a positive peNum static assertion, as the RMSNorm entry point already does; the same missing invariant affects the static overloads and group_norm_grad_1d.
Useful? React with 👍 / 👎.
|
Added rms_norm_binary internal 4PE partitioning in cf76ae8. Dynamic and static [4,8192] PE4 ELFs compile successfully with -j1 and use _start_multipe.s; the kernel partitions A by get_thread_idx(), while the test invokes the kernel once and synchronizes result I/O. Both PE4 gfrun runs currently stop at the existing SuperScalarModel ValidateBasicBinaryTepl dtype/shape/physical-stride assertion, so output.bin is not produced and precision comparison remains blocked. 1PE dynamic/static compatibility builds also pass. |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: cf76ae8adf
ℹ️ 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".
| if (globalA <= 0 || globalA % peNum != 0 || gA < tA) { | ||
| return; |
There was a problem hiding this comment.
When the runtime starts more PEs than the compiled peNum, this guard lets every excess PE continue; for tid == peNum, the offsets on lines 86–89 point exactly past x, out, and the per-row workspace, after which the kernel reads and writes out of bounds. Add a tid >= peNum check here (and the equivalent tid >= gA / peA check in the static entry) before applying the offsets.
Useful? React with 👍 / 👎.
|
Fixed the rms_norm_binary gfrun blocker in 8c71741 by matching the confirmed row-reduction hardware contract: reduction/cache tiles now use physical Columns=1, so TROWSUM output, workspace TSTORE/TLOAD, and TADD operands share the same layout. Validation with -j1: 1PE dynamic/static build + gfrun + precision PASS (max_abs=0, mse=0); 4PE [4,8192] dynamic/static build + gfrun + precision PASS (max_abs=0.0009765625, mse=6.0254e-11; atol/rtol=0.02, mse_tol=0.001). The previous ValidateBasicBinaryTepl assertion is gone. |
|
Fixed GroupNormGrad and GroupNormGrad1D physical-column adaptation in commit 9ca860e. Scalar reduction/broadcast tile_v now uses physical Columns=1 in dynamic and static kernels, matching TROWSUM and row-expand hardware contracts. Validation (-j1, res_check=on): GroupNormGrad 1PE/4PE dynamic+static all build, gfrun R2=0, and precision PASS with dx/dgamma/dbeta max_abs=0. GroupNormGrad1D 1PE/4PE dynamic+static all build, gfrun R2=0, and precision PASS with dx max_abs=1.52587890625e-05, dgamma/dbeta max_abs=0. |
# Conflicts: # benchmark/one-level-arch/test/common/Makefile.common
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: e1cbb6d3ac
ℹ️ 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".
| ifeq ($(PE_NUM),1) | ||
| TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType).elf | ||
| else | ||
| TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R)_PE$(PE_NUM).elf |
There was a problem hiding this comment.
For PE_NUM=4 res_check=on, Makefile.common derives CHK_DIR from this newly suffixed target (for example, ..._gA4_gR8192_PE4), while the generator and comparison scripts still default to the unsuffixed single-PE directory. Consequently the launcher opens a missing input file in a new directory, ignores the short-read failure, runs on zero-initialized buffers, and writes output somewhere the comparator does not inspect; a stale single-PE output can therefore appear to validate the new multi-PE experiment. The same target-directory mismatch is introduced in the RMSNorm and both GroupNorm Makefiles.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 2af846f34f
ℹ️ 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".
| DEFINES += -DDType=$(DType) -DPE_NUM=$(PE_NUM) -DG_A=$(G_A) -DG_R=$(G_R) | ||
| ifeq ($(PE_NUM),1) | ||
| TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType).elf | ||
| TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R).elf |
There was a problem hiding this comment.
Running src/run_precision_check.py with its defaults now builds this shape-suffixed target, but the runner still hardcodes kernel_normalization_rms_norm_rms_norm_DType__half.elf at lines 21–23 and checks that old path after compilation. The documented end-to-end precision check therefore either exits with ELF not found or, if an old artifact remains, simulates a stale binary instead of the one just built; update the runner to derive the new target name (and matching compare directory).
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: a6582b9a9d
ℹ️ 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".
| TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R)_tA$(TILE_A)_tR$(TILE_R).elf | ||
| else | ||
| TARGET = $(ELF_HEAD)_$(TESTCASE)_DType$(DType)_gA$(G_A)_gR$(G_R)_tA$(TILE_A)_tR$(TILE_R)_PE$(PE_NUM).elf | ||
| PE_NUM ?= 4 |
There was a problem hiding this comment.
When using the shipped src/run_precision_check.py after its ELF-name mismatch is corrected, this target can only build PE_NUM=4, but the runner's gfrun command at line 93 supplies only -t 1; benchmark/one-level-arch/test/kernel/matmul/src/gfrun_matmul.py:67-72 documents that -t controls logging rather than thread count, while README.md:317 requires -s softcore.multiThreadNum=4 for multi-PE workloads. With the default single PE, the RMSNorm launcher processes only one quarter of the rows and then waits forever for kernel_done[1..3], so the precision workflow must pass the four-PE simulator setting.
Useful? React with 👍 / 👎.
| if (tid == 0) { | ||
| readBinaryFile(CHK_DIR "/input.bin", (uint8_t *)input, | ||
| static_cast<size_t>(g_a) * g_r * sizeof(dtype)); | ||
| input_ready = 1; |
There was a problem hiding this comment.
In a 4-PE RES_CHECK run, setting this volatile flag does not provide release/acquire ordering for the ordinary shared input buffers, so optimized worker code may observe input_ready == 1 without the preceding file-loaded data being ordered before its kernel reads. The analogous multi-PE harness in benchmark/one-level-arch/test/kernel/multi_thread/matmul/src/matmul_shared.cpp:67-79 places compiler memory clobbers both before publishing leader_ready and after the worker wait; add equivalent synchronization here and in the three launchers that duplicate this flag protocol to avoid validating with partially visible inputs.
Useful? React with 👍 / 👎.
- 基线 ops-20260904(a0ddcc3) + env_test 现构建工具链(model 49547742/llvm 67d3ac98/TileOP f8fb8943) - 新tag含kernels/重构,5个PR(PTO-ISA#82/PTO-ISA#84/PTO-ISA#39/PTO-ISA#78/PTO-ISA#75)merge冲突=降级不忠实;干净merge PTO-ISA#73/PTO-ISA#79/PTO-ISA#83/#74可信 - gfsim 0->8 新模型真实改善;#74本轮首次可验(moe_dispatch/combine全PASS) Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
There was a problem hiding this comment.
💡 Codex Review
Fold all live cache levels for non-power-of-two splits
When the newly configurable G_R produces a non-power-of-two number of cache updates—for example, G_R=10000 derives powR=8192 and tile_r=1024, resulting in nine updates—GetCacheId(r - 1) returns level 0 and this load discards the partial sums still held at other levels. The normalization is consequently computed from only the final partial tile; fold every live level after the loop or reject shapes whose update count is unsupported.
ℹ️ 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".
| inline void rsqrt_newton(TileVec &out, TileVec &a) { | ||
| TRSQRT(out, a); | ||
| auto body = [&](auto &x, auto &t1, auto &t2) { | ||
| TRECIP(x, a); |
There was a problem hiding this comment.
When a row's mean(x²) + eps is below roughly one third (for example, a row filled with 0.1), TRECIP initializes the iteration with 1/a, outside the convergence basin for this inverse-square-root update, so the value diverges and RMSNorm emits overflow or NaNs instead of normalized values near one. benchmark/one-level-arch/kernels/multi_thread/README.md:255-258 documents this exact sequence failing and TRSQRT fixing it; use an inverse-square-root seed here and in the duplicate rms_norm_binary helper.
Useful? React with 👍 / 👎.
| DEFINES += -DRES_CHECK -DENABLE_BINARY_OUTPUT | ||
| DEFINES += -DCHK_DIR=\"$(ROOT)/compare/$(notdir $(basename $(TARGET)))\" | ||
| ifeq ($(PLAT), linx) | ||
| CC_LINK = -nostartfiles $(START_FILE) |
There was a problem hiding this comment.
For Linx res_check=on builds outside the four normalization directories, START_FILE remains the default _start.s, and this reassignment also stops linking group_worker_runtime.o. Fresh repo-wide inspection shows that test/kernel/multi_thread/res_check_all.py still launches these existing targets with four PEs while, for example, test/kernel/multi_thread/normalization/rms_norm/Makefile does not override START_FILE; all PEs therefore direct-boot with the model's shared initial SP instead of workers entering __linx_group_worker_start on independent stacks and parking after main. These checks can overwrite one another's stack frames or stop before PE0 exports the result, so scope the direct-boot override to the new normalization targets or retain the worker entry elsewhere.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 6002cd81fd
ℹ️ 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".
| static_assert(peNum > 0 && peNum <= kMaxPeCount); | ||
| if constexpr (peNum > 1) { | ||
| const uint32_t pe = read_pe_id(); | ||
| kPeBarrier[pe] = phase; |
There was a problem hiding this comment.
Even on the first 4-PE invocation, assigning and polling these volatile slots provides no release/acquire ordering for the ordinary workspace accesses: a PE can observe phase 1 before another PE's ds/db stores are visible, or phase 2 before its c2_buf/c3_buf stores are visible, and then consume stale values in the following stage. Use the platform's memory fence/barrier primitive (or release/acquire atomics) around publishing and observing each phase; this is independent of the separate stale-generation problem on later invocations.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: bed1e4ec3c
ℹ️ 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".
| constexpr int64_t kTileR = binary_tile_r(G_R); | ||
| constexpr int64_t kPowR = floor_power_of_two(G_R - 1); |
There was a problem hiding this comment.
When G_R is overridden with a non-aligned size, these automatic settings can produce a non-power-of-two number of cache updates even though the binary accumulator retrieves only cache[GetCacheId(r - 1)]. For example, G_R=20000 passes the assertions, but kTileR=8192 and kPowR=16384 produce three updates; the final lookup selects level 0, which contains only the last 4576-column chunk, so every row is normalized using an incomplete sum. Choose tiling that guarantees a power-of-two update count or combine all populated cache levels.
Useful? React with 👍 / 👎.
| template <typename dtype, int peNum> | ||
| void rms_norm(dtype *x, const int64_t *tiling, dtype *out, float eps = 1e-6f) { | ||
| static_assert(peNum == 4, "normalization kernels support only 4PE"); |
There was a problem hiding this comment.
A repo-wide search at this commit shows that the existing multi-thread RMSNorm header still includes single_thread/normalization/rms_norm/rms_norm.hpp and invokes the removed compile-time rms_norm<dtype, peA, gA, gR, tA, tR> overload; the binary RMSNorm wrapper has the equivalent stale include and call. Since benchmark/one-level-arch/compile_all.sh still builds both multi-thread targets, they now fail to compile. Keep compatibility headers and overloads or update both consumers to the replacement API.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 88bace0106
ℹ️ 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".
| addi zero, 0x5e, ->x1 | ||
| acrc 1 |
There was a problem hiding this comment.
When the new RMSNorm, binary RMSNorm, or 1D GroupNorm target is built without RES_CHECK for performance measurement, each PE returns from main directly into this 0x5e (SYS_exit_group) path; test/common/src/group_worker_runtime.c:8-11 confirms that any worker issuing it terminates the entire group. Fresh evidence beyond the existing Makefile.common comment is that the new launchers' completion waits are compiled only under RES_CHECK. With a supported uneven partition such as G_A=333, PE3 handles 81 rows while the others handle 84 and can therefore terminate the run before their computation finishes, truncating both results and cycle measurements; synchronize all PEs and let only the leader exit while workers park.
Useful? React with 👍 / 👎.
1 participant
Summary
get_thread_idx()inside the kernel to select each PE's row range1peNum=1Implementation
For dynamic tiling, the kernel now interprets
tiling[0]as the global A dimension, derivespeA = globalA / peNum, and offsets the input/output pointers using the current PE ID. Invalid runtime partitions return without overlapping rows.The static implementation already owns its PE partitioning; this change updates its reduction/broadcast carrier Tiles to physical Columns=1 as well.
The one-level dynamic test no longer calls a per-PE slice of the kernel. Four PEs are represented by the kernel implementation, not by four test-level kernel calls.
Validation
Environment:
Built successfully:
G_A=16, G_R=512, TILE_A=2, TILE_R=512), 1 PESingle-PE dynamic and static
gfrunruns both reached the benchmark end withR2=0and passed the external FP16 accuracy oracle:The 4-PE accuracy run is currently blocked after successful ELF generation by the existing hosted
res_check=onruntime issue:Both dynamic and static 4-PE runs reproduce
SoftCore.cpp:1038: Block BARG targetduring the hosted musl exit path. RMSNorm reproduction evidence was added to SuperScalarModel#332.Toolchain note
The validation installation required the current TileOP/LLVM TEPL selector compatibility discussed in Linx-TileOP-API#29 and #32. This PR contains only the SuperNPUBench RMSNorm changes and does not include toolchain generated-header changes.