fix(miopen): align LRN/LayerNorm grids and guard OOB threads on gfx1250 - #10251
Open
SreecharanGundaboluAMD wants to merge 4 commits into
Open
fix(miopen): align LRN/LayerNorm grids and guard OOB threads on gfx1250#10251SreecharanGundaboluAMD wants to merge 4 commits into
SreecharanGundaboluAMD wants to merge 4 commits into
Conversation
… threads The global work size for MIOpenLRNAcrossChannels4 was set to the raw map_size_4 value, which is not guaranteed to be a multiple of GROUP_SIZE_X. HIP requires global work size to be a multiple of local work size; misaligned launches produce incorrect results by silently dropping boundary threads. Round up the grid x-dimension to the nearest multiple of _grp_tile0 (== GROUP_SIZE_X) using AlignUp, then add an early-return guard in the kernel so the extra threads exit before touching global memory. Co-Authored-By: Claude <noreply@anthropic.com>
…B threads The global work size for MIOpenLRNAcrossChannelsBwd1 was set to the raw _in_df_width / _in_df_height values, which are not guaranteed to be multiples of GROUP_SIZE_X / GROUP_SIZE_Y. This could cause invalid kernel launches and dropped boundary threads. Round up both x and y grid dimensions using AlignUp, and add an early-return guard in the kernel to prevent excess threads from indexing out of bounds into top_df, scale, top, bot, and bot_df. Co-Authored-By: Claude <noreply@anthropic.com>
…threads The non-parallel path of LayernormBackward set xgridsize to the raw problem.inner_size, which is not guaranteed to be a multiple of xlocalsize. The parallel and reduce-sum sibling paths already used AlignUp correctly; this path was an omission. Apply AlignUp(problem.inner_size, xlocalsize) to match the sibling paths, and add an early-return guard in layernormbwdweightbias so excess threads do not perform out-of-bounds writes to dw and db. Co-Authored-By: Claude <noreply@anthropic.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : ROCM-28415
Motivation
MIOpen unit tests on gfx1250 reported 180 failures across two test suites:
Full/GPU*LayerNormBwdTest*{FP32,FP16,BFP16}.LayerNormTestBw— NaN gradients indwanddbfor 5D tensors (NCDHW-style, layout 8)Full/GPU*Lrn*{FP32,FP16}.TestFloat[16(CrossChannel mode only) — large forward-result mismatches well above toleranceBoth failures share the same root cause: kernel launch global work sizes were set to raw problem dimensions that are not guaranteed to be multiples of the local work size. HIP silently drops the boundary threads in misaligned launches, causing incorrect results for inputs whose spatial extent is not evenly divisible by the tile size.
Technical Details
Three commits each follow the same two-step fix pattern:
1.
fix(miopen/lrn): align forward LRN across-channels grid and guard OOB threadsprojects/miopen/src/hip/mloNorm.cpp: rounded the x-grid forMIOpenLRNAcrossChannels4up to the nearest multiple of_grp_tile0usingAlignUp(map_size_4, _grp_tile0).projects/miopen/src/kernels/MIOpenLRNFwd.cpp: added an early-return guard so excess threads exit before touching global memory.2.
fix(miopen/lrn): align backward LRN across-channels grid and guard OOB threadsprojects/miopen/src/hip/mloNorm.cpp: rounded up both x and y grid dimensions forMIOpenLRNAcrossChannelsBwd1usingAlignUpon_in_df_width / _in_df_height.projects/miopen/src/kernels/MIOpenLRNBwd.cpp: added an early-return guard to prevent out-of-bounds reads/writes intotop_df,scale,top,bot, andbot_df.3.
fix(miopen/layernorm): align backward weight/bias grid and guard OOB threadsprojects/miopen/src/solver/layernorm/backward_layernorm.cpp: appliedAlignUp(problem.inner_size, xlocalsize)toxgridsizein the non-parallel path, matching the pattern already used by the parallel and reduce-sum sibling paths.projects/miopen/src/kernels/MIOpenLayerNorm.cpp: added an early-return guard inlayernormbwdweightbiasso excess threads do not perform out-of-bounds writes todwanddb.The fixes are minimal and surgical — no algorithmic or data-layout changes. The
AlignUppattern is already used elsewhere in MIOpen for grid sizing; this change brings the affected code paths into conformance with that convention.Risk Assessment
Risk Level: 🟢 Low
Impacted Components
projects/miopen/src/hip/mloNorm.cpp— LRN kernel launch parametersprojects/miopen/src/kernels/MIOpenLRNFwd.cpp— LRN forward kernel (across-channels)projects/miopen/src/kernels/MIOpenLRNBwd.cpp— LRN backward kernel (across-channels)projects/miopen/src/kernels/MIOpenLayerNorm.cpp— LayerNorm backward weight/bias kernelprojects/miopen/src/solver/layernorm/backward_layernorm.cpp— LayerNorm backward solverPotential Side Effects
Mitigation Steps
GPU*Lrn*andGPU*LayerNormBwdTest*), which previously reproduced the failures.Test Plan
Ran the following test suites on gfx1250 (MI455X - A0, ROCm 7.15.0-dev):
Test Result
All previously failing test cases pass after the fix. WithinChannel LRN tests (which were already passing) continue to pass, confirming no regressions in the sibling code path.
Submission Checklist