test(miopen): add NHWC conv fallback numel>INT_MAX scale coverage (ALMIOPEN-2141) - #10256
Open
peymanr wants to merge 4 commits into
Open
test(miopen): add NHWC conv fallback numel>INT_MAX scale coverage (ALMIOPEN-2141)#10256peymanr wants to merge 4 commits into
peymanr wants to merge 4 commits into
Conversation
…t test
Adds GetNhwcScaleTestCase() with shape {1,1,46342,46342} and 1x1 kernel in
NHWC layout (numel=2,147,580,964 > INT_MAX) to the ConvDirectNaiveConvFwd
unit test. Registers a FullScale instantiation with GPU-side reference
verification (CPU ref is impractical at 4.3 GB per tensor).
Covers the conv fallback path (ConvDirectNaiveConvFwd, the solver used when
no optimised solver applies) against int32 index-wrap OOB at NHWC tensors
larger than INT_MAX elements.
JIRA ID: ALMIOPEN-2141
…tiv test Adds GetNhwcFusedScaleCase() with shape N=1,C=1,H=46342,W=46342,K=1,1x1 kernel (numel=2,147,580,964 > INT_MAX) to cba_infer.cpp. Registers a FullScale instantiation for GPU_ConvBiasActivInfer_FP16 with miopenTensorNHWC layout, covering the conv+bias+activation fusion path at NHWC tensor sizes that overflow int32 index arithmetic. JIRA ID: ALMIOPEN-2141
✅ 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 |
Fix clang-format violations in cba_infer.cpp flagged by pre-commit CI. Resolves: ALMIOPEN-2141
|
🎉 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.
What
Adds two
FullScale-tier test cases covering NHWC conv paths at tensor element counts greater than INT_MAX (2,147,483,647):1. Naive / fallback forward conv (
unit_conv_solver_ConvDirectNaiveConvFwd.cpp)GetNhwcScaleTestCase(): shape{N=1, C=1, H=46342, W=46342}× 1×1 kernel, NHWC layoutGetScaleTestParams()usingUseGpuRef()(CPU reference is impractical at 4.3 GB per tensor)FullScaleinstantiation forGPU_UnitTestConvSolverDirectNaiveFwd_FP162. Fused conv+bias+activation (
cba_infer.cpp)GetNhwcFusedScaleCase(): same shape, 1×1 kernel,miopenConvolutionFullScaleinstantiation forGPU_ConvBiasActivInfer_FP16withmiopenTensorNHWClayoutWhy
Two NHWC conv surfaces were untested at numel > INT_MAX:
ConvDirectNaiveConvFwd— the solver used when no optimised solver applies): index arithmetic that usesint32_twraps at 2 GB elements, causing silent out-of-bounds writes.ConvHipDirectFwdFused/ConvCKIgemmFwdBiasActivFusedvia fusion plan): same risk in the fused launch path.The grouped-CK int32-wrap case is already covered by PR #7663's degenerate case. This PR specifically targets the fallback and fused surfaces identified in ALMIOPEN-2141.
How
projects/miopen/test/gtest/FullScaletier (large-memory GPU required, e.g. MI300X / MI325X with ≥10 GB free VRAM)Testing
Run on gfx942/gfx950 hardware:
JIRA ID: ALMIOPEN-2141
🤖 Assisted by PR Pundit