Repository navigation
feat(integration-tests): per-tensor validator selection from the engine TOML - #11716
Merged
Merged
Conversation
…ensor from the engine TOML ## Motivation The integration-test harnesses grade every output tensor with one comparison: per-element allclose, `|ref - impl| <= atol + rtol*|ref|`. That is the wrong question for a reduction output whose elements can land arbitrarily close to zero through cancellation, because the absolute error tracks the tensor's scale while `rtol*|ref|` collapses to nothing. Layernorm backward's `dscale`/`dbias` are exactly that: sums over up to ~9e5 zero-mean terms. Measured on gfx1151 / HIP_MLOPS_ENGINE, 22 of 132 fp32/upcast Smoke cases fail under allclose. The `atol` that would rescue them ranges 0.9x-52x the default depending on reduction length, swings ~2x with the input seed, and — since a tolerance override is per test, not per tensor — would also apply to `dx`, whose real error is four to five decades smaller. Relative RMS error stays in 1e-6..8e-6 across the whole set. The C++ test worked around this by calling `registerRmsValidator()` directly, which the bundle harness had no equivalent of. That blocked migrating the test to bundles (PR #11306) and left RMSNorm backward's equivalent diagnosed as flakiness and skipped. ## Technical Details Comparison mechanism: - `ComparisonTolerance` carries a `ValidatorKind` (`ALLCLOSE` or `RMS`) plus an RMS threshold, with `allClose()`/`rms()` factories. - The tolerance lookup takes `(uid, label, dataType)` instead of `dataType`. `label` is the tensor's name, or `uid=N` when unnamed — uids are not stable between a C++ graph test and the bundle captured from it (`DSCALE` is uid 8 in the test, uid 1 in the bundle), names are. - `compareTensor()` dispatches to `createRmsValidator()`, and an RMS failure reports its own threshold via the new `ComparisonContext::toleranceSummary` rather than printing atol/rtol that did not decide it. Selection: - New `[[validator_overrides]]` section in each engine's test-config TOML: `filters` (gtest-name globs) + `tensors` (tensor-label globs) + `validator` + `rms_threshold`. Both globs must match; later entries win. The parser is strict: a missing `tensors`, an unknown `validator`, or a missing or non-positive `rms_threshold` is a load error, never a silent fall-back. - allclose remains the default everywhere and this section is the only thing that changes it. `registerRmsValidator()` is removed from the shared harness so a test cannot select a validator from C++ — that selection was invisible to config and applied to every engine, while the reason for RMS is a per-engine numerical property. (miopen-provider's identically-named function belongs to its own local harness class and is untouched.) - `BundleReferenceValidationHarness` is unchanged in behaviour and now says so explicitly: reference-vs-golden always uses allclose at the default tolerance. An engine's config describes how far that engine may drift; it never relaxes the gate on our own committed data. Net coverage is unchanged: layernorm backward's `dscale`/`dbias` were RMS-graded before via the C++ call and are RMS-graded after via `HIP_MLOPS_ENGINE.toml`. Landing the layernorm-backward bundles and retiring the C++ test, and the same treatment for RMSNorm backward, are follow-ups. ## Test Plan - New unit tests for the comparison and the TOML parser. - A/B the layernorm backward Smoke tier with and without the TOML entry. - Prove per-tensor selection reaches the bundle path, not just the C++ path. - Regression the quick bundle tier and the other C++ graph suites. ## Test Result - `TestOutputComparison`: 12/12 (4 new — RMS accepts a near-zero element allclose rejects; RMS still rejects real drift; the report names the RMS threshold and not atol; the kind is chosen per tensor from the label). - `TestTestSettings`: 8 new covering both-globs matching, other tensors in the same test unaffected, other tests with the same tensor name unaffected, precedence, and all four parser rejections. - `hipdnn_integration_tests_unit_tests`: 543/543. - `hipdnn_test_sdk_tests`: pass (1 new for `toleranceSummary`). - Smoke layernorm backward fp32+upcast, with the TOML entry: 132 passed. With the entry removed: 110 passed, 22 failed. - Bundle path proven independently on `quick_RMSNorm_Variant3` (no layernorm bundles exist yet): forcing `atol=rtol=0` fails `RMSNorm_0::Y`; adding an RMS override for `*::Y` at 1e-2 turns 43/43 green; at 1e-30 it fails and reports `Tolerance: relative RMS <= 1e-30`. - `Smoke/*RMSNorm*:Smoke/*Layernorm*`: 1508 ran, 1388 passed, 120 pre-existing TOML skips, 0 failed. - `quick_*` bundle tier: 3051 ran, 1663 passed, 1388 skipped, 0 failed. - clang-format clean; clang-tidy clean on all 7 affected translation units (the one remaining finding, `kUserProfileToken` in `PlatformUtils.windows.hpp`, is pre-existing and reproduces on a clean tree). Tracking: ALMIOPEN-2561 ## Submission Checklist - [X] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
✅ 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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #11716 +/- ##
===========================================
- Coverage 70.55% 63.79% -6.76%
===========================================
Files 2810 2825 +15
Lines 462535 469166 +6631
Branches 68105 69274 +1169
===========================================
- Hits 326312 299281 -27031
- Misses 112680 148839 +36159
+ Partials 23543 21046 -2497
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
… keyword test_name_validator.py reserves the 'Test' keyword for the suite name: a case name containing it trips both the positional-keyword rule and the duplicate-keyword rule, which is what failed hipdnn-integration-tests_test_name_validation in the superbuild CI. Tracking: ALMIOPEN-2561
adickin-amd
commented
Sep 8, 2026
adickin-amd
commented
Sep 8, 2026
adickin-amd
commented
Sep 8, 2026
adickin-amd
commented
Sep 8, 2026
adickin-amd
commented
Sep 8, 2026
…e report between harnesses Review follow-up on the per-tensor validator work. The two harnesses each had their own copy of the [[validator_overrides]] lookup and their own report formatting, and the copies had already drifted: the C++ graph path passed a raw tensor name where the bundle path passed the "uid=N"-capable label, and the RMS tolerance summary only ever reached the bundle path. Layernorm backward is a C++ graph test, so the one override this branch ships ran on the path that printed "atol=0 rtol=0" for a check atol/rtol never decided. - bundle::tensorLabel gains a (uid, name) overload; both harnesses resolve a label through it, so a TOML glob and a failure report always agree on what a tensor is called. - bundle::formatMismatchReport owns the header and drift profile for both paths. TensorValidationEntry carries the ComparisonTolerance rather than a bare atol/rtol pair, which is what let the RMS summary go missing. - findTomlRmsThreshold in TomlGuards.hpp owns the config lookup, the kind == RMS test and the log line, collapsing the duplicated block at both call sites. - bundle::makeValidator turns createRmsValidator's unsupported-dtype throw into a failure naming the tensor, its dtype and the config section to narrow. A 'tensors' glob one wildcard too wide can select RMS for an integer output, and the parser cannot catch it: dtype is not known until the graph is read. - The parser rejects an allclose entry that also sets rms_threshold, matching the strictness of the surrounding rejections. - Drop a stray empty .tmp_names.txt committed by accident. JIRA ID : ALMIOPEN-2561
…ernorm_rms_validation
…ernorm_rms_validation
adickin-amd
marked this pull request as ready for review
September 10, 2026 14:30
…ernorm_rms_validation
adickin-amd
commented
Sep 10, 2026
adickin-amd
commented
Sep 10, 2026
adickin-amd
commented
Sep 10, 2026
…port the threshold Second review follow-up. Both changes are shape, not semantics — the rendered failure header is byte-identical before and after. makeValidator tested `kind != RMS` and treated everything else as allclose, so a future ValidatorKind would have been graded by a validator nobody chose. It now switches on the kind with a throwing default, which is also what the project already asks for: -Wswitch-default is in the integration-tests warning set and bugprone-switch-missing-default-case is enabled in .clang-tidy. The try/catch stays on the RMS arm — RMS is the one kind a [[validator_overrides]] glob can select, so an unsupported data type there is an operator's config mistake worth naming, whereas allclose is the default that nothing selects and has no glob to blame. ComparisonContext took a pre-rendered `toleranceSummary` string composed in dnn-providers/integration-tests and printed verbatim by a projects/hipdnn/test_sdk formatter — presentation logic living a component away from the thing that presents it, which is the same split that let the RMS summary go missing on one path in the first place. It now carries `std::optional<float> rmsThreshold` and formatComparisonHeader owns the wording, including the parenthetical about the element counts that appendComparisonDiffByDataType emits two lines later. Tests: UnhandledValidatorKindIsRefused guards the new default case; NoRmsThresholdPrintsAtolRtol guards the else branch; the renamed RmsThresholdReplacesAtolRtol now also asserts the parenthetical the SDK owns. JIRA ID : ALMIOPEN-2561
brentmaas
approved these changes
Sep 11, 2026
brentmaas
left a comment
brentmaas
left a comment
Contributor
There was a problem hiding this comment.
LGTM! One comment regarding the RMS threshold in the validator override which could help resolve the test failure.
…olds Use the existing layernorm thresholds for pure FP16 and BF16 dscale/dbias while preserving FP32 thresholds for mixed and upcast fixtures. Document output dtype matching and last-match precedence. Validation: 17 targeted integration tests passed on Windows gfx1151. A temporary tighter FP32 threshold produced 10 expected failures while all 7 pure FP16/BF16 cases passed; restoring the threshold returned 17 passes. JIRA ID : ALMIOPEN-2561
…ernorm_rms_validation
mousdahl-amd
approved these changes
Sep 14, 2026
mousdahl-amd
left a comment
mousdahl-amd
left a comment
Contributor
There was a problem hiding this comment.
Looks good, just scrub the JIRA ids.
NolanHannaAMD
approved these changes
Sep 14, 2026
NolanHannaAMD
left a comment
NolanHannaAMD
left a comment
Contributor
There was a problem hiding this comment.
Just a handful of non-critical things. Nothing blocking.
bghimireamd
reviewed
Sep 14, 2026
Contributor
|
LGTM! |
bghimireamd
approved these changes
Sep 14, 2026
…nsor Review follow-ups on the per-tensor validator selection. Tolerance logging moved out of the resolver and into one shared decision site. applyTomlToleranceOverride() is handed a test, not a tensor, so it could not know whether a [[validator_overrides]] entry outranked the atol/rtol it announced; a reader chasing a layernorm-backward result was told the wrong check ran up to four times per case. gradingForTensor() now reads the validator override first and logs exactly one line per output tensor, naming the comparison that actually graded it. Both harnesses resolve every output through it, so neither can report a check the other would not have run. ToleranceLookup drops its uid parameter. No callback read it, and the reason this PR keys on labels is that uids are not stable between a C++ graph test and the bundle captured from it. Also: two more parser rejection paths under test (an entry with no 'filters', and a non-table entry — a 'filters'-less entry reads as "apply everywhere" and is a likely hand-edit), the layernorm-backward call site says that the tolerance it passes is the fallback for engines whose config selects no validator, and JIRA ids are out of the comments this PR added.
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
brentmaas
added a commit
that referenced
this pull request
Sep 30, 2026
… layernorm bwd (#12639) ## Motivation A follow-up PR for #9500 and #11306 after the merges of #11716 and #12221 enable the use of RMS validators in test bundles and reorganised the layernorm backward tests. This PR generates those test bundles and support claims. Closes #10168. **As with #12107, generating these test bundles relied on the tooling hotfix for `place_bundles.py` (and extended to `verify_migration.py`) described in #11306.** ## Technical Details - Generated layernorm bwd test bundles. - Generated support claims for the new layernorm bwd test bundles. - Update test TOML to include the new test bundles. - Deregister layernorm bwd from always building. ## Test Plan Run the tests using the new test bundles, verify support claims. ## Test Result All tests pass, new support claims verify. ## Submission Checklist - [X] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. Co-authored-by: Adam Dickin <adam.dickin@amd.com>
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.
Summary
The integration-test harnesses grade every output tensor with one comparison: per-element allclose. That is the wrong question for a reduction output whose elements can land arbitrarily near zero through cancellation, because the absolute error tracks the tensor's scale while
rtol*|ref|collapses. This adds a[[validator_overrides]]section to each engine's test-config TOML so a named output tensor can be graded by MIOpen's aggregate relative-RMS check instead, unblocking the migration of layernorm-backward coverage to golden bundles. allclose remains the default everywhere, and the engine TOML is the only thing that changes it.JIRA ID : ALMIOPEN-2561
Follow-up to the layernorm-backward bundle work in PR #11306, whose Copilot review identified this gap. Implements the tolerance-chain intent of RFC 0011 for validator kind.
Risk Assessment
Medium risk. No product code changes — this is test infrastructure — but it touches the comparison path every integration test on every ASIC flows through, and adds a config schema surface. Mitigating: net coverage is unchanged by construction (layernorm backward's
dscale/dbiaswere RMS-graded before via a C++ call and after via the TOML), the new path is inert unless a[[validator_overrides]]entry matches, and only one entry exists, scoped to*LayernormBackward*onHIP_MLOPS_ENGINE. Residual risk is that local validation covered one ASIC and one provider preset, so the multi-arch run below is the outstanding gate.ASIC Coverage
Multi-arch CI run required; a full sweep across every supported family is not. The change adds no op and no support surface, and does not touch kernel selection or dispatch — but it does alter the comparison code path that existing test cases on every architecture already exercise, so an arch-specific regression would surface as test failures rather than build failures. The routine multi-arch lanes (
gfx94X,gfx950on Linux;gfx110Xon Windows) are the right gate.Two coverage notes for reviewers. Local validation ran on gfx1151 (Windows) only, which is outside those lanes. Local builds used the
hip-kernel-providerpreset only, somiopen-providerandhipblaslt-providerwere not compiled against these headers; a repo-wide search found no consumer of the changed headers outsidednn-providers/integration-testsandprojects/hipdnn/test_sdk(miopen-provider's identically-named harness is a separate local class), so this is expected to be covered by PR CI.The RMS threshold committed for layernorm backward is the existing per-op layernorm tolerance and is unchanged from what the C++ test already passed to the RMS validator, so it introduces no new arch-dependent constant.
Testing Summary
Testing Checklist
git commit(14 pre-commit hooks incl. clang-format) - Status: Passedhipdnn_integration_tests_unit_tests- 543/543, 12 new - Status: Passedhipdnn_test_sdk_tests- 1 new forComparisonContext::toleranceSummary- Status: Passedhipdnn_integration_tests --test-engine HIP_MLOPS_ENGINE --gtest_filter='Smoke/*LayernormBackward*Fp32*:Smoke/*LayernormBackward*Upcast*'- 132 passed - ASICs: gfx1151 - Status: Passed[[validator_overrides]]stripped - 110 passed, 22 failed as expected - ASICs: gfx1151 - Status: Passed--gtest_filter=quick_RMSNorm_Variant3.*withatol=rtol=0plus an RMS override on*::Y; 43/43 at threshold 1e-2, correct RMS-formatted failure at 1e-30 - ASICs: gfx1151 - Status: Passed--gtest_filter='Smoke/*RMSNorm*:Smoke/*Layernorm*'- 1508 ran, 1388 passed, 120 pre-existing TOML skips, 0 failed - ASICs: gfx1151 - Status: Passed--gtest_filter=quick_*- 3051 ran, 1663 passed, 1388 skipped, 0 failed - ASICs: gfx1151 - Status: Passedtidy_changed.py --base origin/develop- clean on all 7 affected translation units - Status: PassedTechnical Changes
ComparisonTolerancegains aValidatorKind(ALLCLOSEorRMS) and an RMS threshold, withallClose()/rms()factories;compareTensor()dispatches tocreateRmsValidator()accordingly.(uid, label, dataType)rather thandataTypealone.labelis the tensor's name, oruid=Nwhen unnamed — uids are not stable between a C++ graph test and the bundle captured from it (DSCALEis uid 8 in the test, uid 1 in the bundle), names are.ComparisonContext::toleranceSummary, instead of printing atol/rtol that did not decide the verdict.[[validator_overrides]]TOML section:filters(gtest-name globs) plustensors(tensor-label globs),validator, andrms_threshold. Both globs must match; later entries win. The parser rejects a missingtensors, an unknownvalidator, and a missing or non-positiverms_thresholdat load rather than silently falling back to allclose.registerRmsValidator()is removed from the sharedIntegrationGraphVerificationHarness, so a test can no longer select a validator from C++ — that selection was invisible to config and applied to every engine, whereas the reason for RMS is a per-engine numerical property. Verified no remaining callers; miopen-provider's identically-named function belongs to its own local harness class and is untouched.BundleReferenceValidationHarnesskeeps its existing behaviour and now documents it: reference-vs-golden comparison always uses allclose at the default tolerance, because an engine's config describes how far that engine may drift and must never relax the gate on committed golden data.HIP_MLOPS_ENGINE.tomlselects RMS for layernorm backward'sdscale/dbias, with the measured justification recorded alongside it; the integration-tests README documents the new section.