test(normalization): back up dynamic 4PE validation implementation - #106
wangyuascend-spec wants to merge 12 commits into
Conversation
# Conflicts: # benchmark/one-level-arch/test/common/Makefile.common
![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: 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".
| make TESTCASE=rms_norm_binary COMPILER_DIR="$COMPILER_DIR" DType=__half \ | ||
| G_A=4 G_R=8192 TILE_A=1 TILE_R=1024 POW_R=4096 PE_NUM=4 diss | ||
| make TESTCASE=rms_norm_binary_static COMPILER_DIR="$COMPILER_DIR" DType=__half \ | ||
| G_A=16 G_R=16384 PE_NUM=4 diss |
There was a problem hiding this comment.
This G_R=16384 invocation leaves POW_R at the Makefile default of 4096, but the static kernel requires G_R <= 2 * POW_R; consequently the final command in compile.all hits the static_assert in rms_norm_binary_pto.hpp and, because the script uses set -e, makes every batch build fail. Pass POW_R=8192 (and the intended tile values) for this shape.
Useful? React with 👍 / 👎.
| make TESTCASE=rms_norm_binary COMPILER_DIR="$COMPILER_DIR" DType=__half \ | ||
| G_A=4 G_R=8192 TILE_A=1 TILE_R=1024 POW_R=4096 PE_NUM=4 diss |
There was a problem hiding this comment.
The dynamic 4PE command still overrides the shape to [4,8192], while this script's header, the generator defaults, and the comparison directory all target [16,16384]. Running the documented precision workflow therefore produces an ELF whose CHK_DIR and output size do not match the generated inputs/golden data, so the newly advertised dynamic validation case cannot actually be checked.
Useful? React with 👍 / 👎.
| kPeBarrier[pe] = phase; | ||
| for (int participant = 0; participant < peNum; ++participant) { | ||
| while (kPeBarrier[participant] < phase) { | ||
| } |
There was a problem hiding this comment.
When the same process invokes group_norm_grad more than once with multiple PEs, the first call leaves every barrier slot at phase 3. The next call's phase-1 and phase-2 waits then pass immediately because they only test < phase, allowing a PE to consume ds, db, or c2/c3 before other PEs have refreshed them and producing stale gradients. Use an invocation-specific monotonically increasing epoch or safely reset the slots between calls.
Useful? React with 👍 / 👎.
| 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.
Renaming the single-PE target to include _gA$(G_A)_gR$(G_R) breaks src/run_precision_check.py, which still builds TESTCASE=rms_norm and then looks specifically for kernel_normalization_rms_norm_rms_norm_DType__half.elf. With the new defaults, make emits the _gA512_gR512.elf name instead, so the end-to-end precision helper always stops at its ELF not found check after a successful build; update the runner and its comparison directory derivation or retain the canonical target name.
Useful? React with 👍 / 👎.
| // tiling: {N, C, G, HxW, tile_hw} | ||
| int64_t tiling_info[5] = {N_BATCH, C_CH, G_GRP, HxW_SZ, TILE_HW}; | ||
| constexpr int64_t kTileHw = group_norm_tile_hw<dtype>(HxW_SZ); | ||
| static_assert(N_BATCH > 0 && C_CH > 0 && G_GRP > 0 && HxW_SZ > 0); | ||
| static_assert(C_CH % G_GRP == 0 && kTileHw > 0); | ||
| int64_t tiling_info[5] = {N_BATCH, C_CH, G_GRP, HxW_SZ, kTileHw}; |
There was a problem hiding this comment.
The dynamic host now always derives kTileHw as min(HxW_SZ, 512) and the Makefile no longer forwards TILE_HW, so explicit configurations are silently ignored. In particular, the first dynamic command in this suite's compile.all requests HxW_SZ=16 TILE_HW=8, but the executable actually passes 16 and therefore no longer exercises the spatial R-split path that the case was intended to cover. Preserve the TILE_HW override when constructing tiling_info.
Useful? React with 👍 / 👎.
1 participant
Backup PR for the current normalization dynamic 4PE implementation from PR #84. Includes the ops-20260828 integration, runtime shape-derived tiling, enlarged validation shapes, GroupNormGrad HxW tiling support, and removal of the hard-coded PE barrier address. The four dynamic 4PE gfrun accuracy cases passed locally. Source snapshot: 2af846f.