test(miopen): add BatchNorm numel>INT_MAX scale coverage (ALMIOPEN-2144) - #10259
Open
peymanr wants to merge 5 commits into
Open
test(miopen): add BatchNorm numel>INT_MAX scale coverage (ALMIOPEN-2144)#10259peymanr wants to merge 5 commits into
peymanr wants to merge 5 commits into
Conversation
JIRA ID: ALMIOPEN-2144
Add projects/miopen/test/gtest/bn_fwd_train_large.cpp with a single
Standard-tier test that runs MIOpen spatial BN forward training on a
shape where N*C*H*W = 2*2048*1024*512 = 2^31 = 2,147,483,648 > INT_MAX
(2^31-1 = 2,147,483,647).
All existing BN gtest shapes stay below INT_MAX. The largest 2-D case
tops out at ~6.7 M elements; the largest 3-D case is {2,2048,16,128,128}
(~1.07e9, currently DISABLED via WORKAROUND_SWDEV_549725). No shape
exercises the int32 index-overflow path at numel > INT_MAX.
The new test:
- Skips gracefully via GTEST_SKIP() when free GPU memory < 20 GiB
(input + output require ~16 GiB; 20 GiB includes a safety margin)
- Allocates all tensors directly on GPU via hipMalloc to avoid an
equivalent host allocation
- Initialises input with hipMemset(0x3F) — a non-zero, finite pattern
- Asserts every sampled output element is finite; NaN/Inf flags the
int32 index-overflow failure mode
- No CPU reference comparison (host 8 GiB buffer is prohibitive)
MIOpenDriver equivalent for manual verification on TW325 (gfx942):
MIOpenDriver bnorm -n 2 -c 2048 -H 1024 -W 512 -m 1 --forw 1 -V 0 -t 1
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
Pre-commit check failed⛔ pre-commit failed Please run locally:
This repo uses |
- clang-format: collapsed short one-liner bodies, aligned function call args per MIOpen .clang-format rules (AlignAfterOpenBracket, AllowShortFunctionsOnASingleLine) - gtest-format-miopen: renamed instantiation DISABLED_Standard → Standard; DISABLED_ prefix on instantiation name is not a valid TEST_TYPE per the naming convention checker (must match Smoke|Standard|Full|Perf|Unit) - no trailing-whitespace violations; end-of-file newline already present
|
🎉 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
1 participant
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: ALMIOPEN-2144
What
Adds
projects/miopen/test/gtest/bn_fwd_train_large.cpp— a gtest for spatial BatchNorm forward training at numel > INT_MAX.Shape:
{N=2, C=2048, H=1024, W=512}→ N×C×H×W = 2^31 = 2,147,483,648 > INT_MAX.Why
All existing BN test shapes stay below INT_MAX. This PR documents the gap and adds a test as
DISABLED_Standard.Test result
Validated on AMD Instinct MI325X (gfx942, xnack-) — kernel crashes at this scale due to int32 index overflow in the BN spatial kernel, the same class of bug as ALMIOPEN-2151 / ALMIOPEN-2152. The test is added as
DISABLED_Standardto track the coverage gap; it should be re-enabled once the BN overflow is fixed.Notes
DISABLED_Standard— appears in CI test discovery but does not run until re-enabled🤖 Assisted by PR Pundit