Repository navigation
test(miopen): add gfx103X tolerance for ConvHipImplicitGemmV4R1WrW (ROCM-23567) - #14905
Merged
JonathanLichtnerAMD merged 2 commits intoNov 4, 2026
Merged
JonathanLichtnerAMD merged 2 commits into
JonathanLichtnerAMD merged 2 commits into
Conversation
…OCM-23567) gfx1030 (Navi21) was falling back to the default tolerance (1.0, i.e. exactly FLT_EPSILON) because supported_gpus declared gfx103X but SetTolerance only covered gfx908/gfx90A. The reported RMS error (2.2468e-07) is 1.88 ULP of accumulation-order divergence on a backward-weights reduction, not a correctness issue. 4.0f gives ~2.1x headroom over the observed value while staying well under the 40.0f already used for gfx908/gfx90A.
✅ 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. |
4.0f was derived from the ~1.88-ULP error in the original ticket, but that was a single CI sample. Verified on real gfx1030 hardware (AMD Radeon Pro V620, ctr-mlib-02) that the actual error is ~30 ULP, well over 4.0f. Matching the existing gfx908/gfx90A value (40.0f) rather than guessing another intermediate number, since this solver's static_ck numerics already carry that same documented looseness on gfx9 (see the existing "40.0f is too much" TODO above).
JonathanLichtnerAMD
deleted the
users/jlichtne/ROCM-23567-gfx103x-wrw-tolerance
branch
November 4, 2026 20:04
cycheng
pushed a commit
that referenced
this pull request
Nov 6, 2026
…OCM-23567) (#14905) JIRA ID: ROCM-23567 ## Summary `Smoke/GPU_UnitTestConvSolverHipImplicitGemmV4R1WrW_FP32.ConvHipImplicitGemmV4R1WrW/0` has been failing in TheRock's nightly gfx103X run since 2026-04 (ROCM-23567). Root cause is a test calibration gap, not a product defect: - `GetTestParams()` declares `Gpu::gfx103X` in `supported_gpus` but only calls `SetTolerance` for `gfx908`/`gfx90A` (at `40.0f`). - gfx1030 therefore falls back to the default tolerance of `1.0`, making the threshold exactly `FLT_EPSILON` = `1.1920928955078125e-07` — matching the reported threshold bit-for-bit. - The verification metric is `rms_range` (RMS relative error), so the reported `2.2468e-07` is **1.88 ULP** of accumulation-order divergence on a backward-weights reduction — not a correctness problem. This adds `p.SetTolerance(Gpu::gfx103X, miopenFloat, 40.0f);`, matching the existing `gfx908`/`gfx90A` value. An initial attempt used `4.0f` (~2.1x headroom over the ticket's reported 1.88x), but real-hardware testing (see below) showed ~30 ULP error on this device — well over `4.0f` — so the value was raised to match gfx9 rather than pick another unverified intermediate number. This solver's `static_ck` numerics already carry that same documented looseness on gfx9 (see the existing `/// \todo 40.0f is too much` comment above the gfx9 lines). Of the four tests originally reported in ROCM-23567, this is the only one still live — the two `ConvOclBwdWrW2` tests were removed with their solvers (#6859), and `MIOpenDriverRegressionBigTensorTest_FP32` is now excluded from every `test_categories.yaml` category (#6198). ## Test plan - [x] `cmake --build projects/miopen/build --target miopen_gtest` — builds clean on gfx1151. - [x] `./miopen_gtest.exe --gtest_filter='*UnitTestConvSolverHipImplicitGemmV4R1WrW*'` on local gfx1151 — GPU FP32/FP16/BFP16 variants report `SKIPPED` (gfx1151 isn't in `supported_gpus`), CPU applicability tests still `OK`. Confirms the `SetTolerance` call is inert off gfx103X. - [x] **Validated on real gfx1030 hardware** (AMD Radeon Pro V620, the same SKU named in the original defect report) via an AMD-internal Slurm cluster, since rocm-libraries pre-merge CI does not cover gfx103X (gfx942-only) and TheRock's gfx1030 CI runners are currently disabled cluster-wide (`therock-ci-config#57`, known-bad-runner exclusion, unrelated to this change). Built MIOpen from this branch inside `rocm/miopen-dev:latest` and ran `test_unit_conv_solver_ConvHipImplicitGemmV4R1WrW --gtest_filter='*ConvHipImplicitGemmV4R1WrW*'`: - At the default tolerance (no override): FP32 failed, `3.59e-06` actual vs `1.19e-07` threshold — reproduces the ticket. - At `4.0f`: FP32 still failed, `3.59e-06` actual vs `4.77e-07` threshold (~30 ULP, not the ~1.88 ULP the ticket reported) — confirms `4.0f` was insufficient and motivated the revision to `40.0f`. - At `40.0f` (final): **all 5 tests pass** — `[ PASSED ] 5 tests.` 🤖 Generated with [Claude Code](https://claude.com/claude-code)
Contributor
TheRock Submodule Bump ActivityNewest first
|
JonathanLichtnerAMD
added a commit
that referenced
this pull request
Nov 6, 2026
JIRA ID: ALMIOPEN-3158 ## Motivation Follow-up to #14905 (ROCM-23567). The stale `/// \todo 40.0f is too much. The solver needs to be checked.` in `unit_conv_solver_ConvHipImplicitGemmV4R1WrW.cpp` implied a possible solver defect. ## Technical Details Replaces the stale todo with a comment stating the tolerance reflects FP32 accumulation order over K = N*Ho*Wo = 193600 with zero-mean random data, not a solver defect. Comment-only; no functional change. ## Risk Assessment **Risk Level: 🟢 Low** Comment-only change. ## Test Plan Investigation runs on gfx90a (MI210): measure rms relative error vs the GPU reference (accumulates in double) for N=8/16/32/64, with random data and with small-integer data. Nothing to build since the change is comment-only. ## Test Result - Random data, rms relative error grows with K: N=8: 2.4 eps, N=16: 4.5, N=32: 11.6, N=64: 30.1 eps (deterministic). The N=64 value 3.59e-06 matches the gfx1030 result. - Small-integer inputs (exactly representable in FP32): N=8/16/32/64 give exactly zero error at default tolerance, so the kernel computes the correct sum.
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
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-23567
Summary
Smoke/GPU_UnitTestConvSolverHipImplicitGemmV4R1WrW_FP32.ConvHipImplicitGemmV4R1WrW/0has been failing in TheRock's nightly gfx103X run since 2026-04 (ROCM-23567). Root cause is a test calibration gap, not a product defect:GetTestParams()declaresGpu::gfx103Xinsupported_gpusbut only callsSetToleranceforgfx908/gfx90A(at40.0f).1.0, making the threshold exactlyFLT_EPSILON=1.1920928955078125e-07— matching the reported threshold bit-for-bit.rms_range(RMS relative error), so the reported2.2468e-07is 1.88 ULP of accumulation-order divergence on a backward-weights reduction — not a correctness problem.This adds
p.SetTolerance(Gpu::gfx103X, miopenFloat, 40.0f);, matching the existinggfx908/gfx90Avalue. An initial attempt used4.0f(~2.1x headroom over the ticket's reported 1.88x), but real-hardware testing (see below) showed ~30 ULP error on this device — well over4.0f— so the value was raised to match gfx9 rather than pick another unverified intermediate number. This solver'sstatic_cknumerics already carry that same documented looseness on gfx9 (see the existing/// \todo 40.0f is too muchcomment above the gfx9 lines).Of the four tests originally reported in ROCM-23567, this is the only one still live — the two
ConvOclBwdWrW2tests were removed with their solvers (#6859), andMIOpenDriverRegressionBigTensorTest_FP32is now excluded from everytest_categories.yamlcategory (#6198).Test plan
cmake --build projects/miopen/build --target miopen_gtest— builds clean on gfx1151../miopen_gtest.exe --gtest_filter='*UnitTestConvSolverHipImplicitGemmV4R1WrW*'on local gfx1151 — GPU FP32/FP16/BFP16 variants reportSKIPPED(gfx1151 isn't insupported_gpus), CPU applicability tests stillOK. Confirms theSetTolerancecall is inert off gfx103X.therock-ci-config#57, known-bad-runner exclusion, unrelated to this change). Built MIOpen from this branch insiderocm/miopen-dev:latestand rantest_unit_conv_solver_ConvHipImplicitGemmV4R1WrW --gtest_filter='*ConvHipImplicitGemmV4R1WrW*':3.59e-06actual vs1.19e-07threshold — reproduces the ticket.4.0f: FP32 still failed,3.59e-06actual vs4.77e-07threshold (~30 ULP, not the ~1.88 ULP the ticket reported) — confirms4.0fwas insufficient and motivated the revision to40.0f.40.0f(final): all 5 tests pass —[ PASSED ] 5 tests.🤖 Generated with Claude Code