From 31dc74c4990104a5d6f0d513a9922d098954be41 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Fri, 4 Sep 2026 12:32:55 -0600 Subject: [PATCH 1/6] feat(integration-tests): select the output comparison validator per tensor from the engine TOML MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## 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. --- .../config/HIP_MLOPS_ENGINE.toml | 26 +++ dnn-providers/integration-tests/README.md | 29 ++- .../IntegrationGraphVerificationHarness.hpp | 57 +++--- .../src/harness/TestConfig.hpp | 14 ++ .../src/harness/TestSettings.hpp | 143 +++++++++++++++ .../BundleReferenceValidationHarness.cpp | 9 +- .../IntegrationBundleVerificationHarness.cpp | 16 +- .../src/harness/bundle/OutputComparison.cpp | 56 ++++-- .../src/harness/bundle/OutputComparison.hpp | 60 +++++-- .../IntegrationGpuLayernormBackward.cpp | 9 +- .../tests/TestOutputComparison.cpp | 136 +++++++++++++- .../tests/TestTestSettings.cpp | 170 ++++++++++++++++++ .../utilities/ComparisonReport.hpp | 14 +- .../tests/utilities/TestComparisonReport.cpp | 19 +- 14 files changed, 684 insertions(+), 74 deletions(-) diff --git a/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml b/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml index 0ffaf30696be..697c57d4f45f 100644 --- a/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml +++ b/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml @@ -22,6 +22,32 @@ version = 1 # atol: absolute tolerance # rtol: relative tolerance # +# -- Validator Overrides -- +# Change which comparison an output tensor is graded by. allclose is the default +# everywhere and this is the only thing that can change it. Each entry requires: +# filters: array of GTest globs matched against the full GTest name +# tensors: array of globs matched against the output tensor's label — its +# name, or "uid=N" when the graph did not name it +# validator: "allclose" or "rms" +# rms_threshold: required when validator = "rms" +# Later entries take precedence. +# +# Layernorm backward's dscale/dbias are reductions over the batch axes: +# dscale_j = sum_i dy*xhat, dbias_j = sum_i dy +# over up to ~9e5 zero-mean terms. The GPU's reduction order differs from the CPU +# reference's, and the resulting absolute error tracks the *tensor's* scale, not each +# element's. Elements that land near zero through cancellation then blow past +# atol + rtol*|ref| while the tensor as a whole is accurate to ~1e-6 relative. +# Measured on gfx1151: 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 and +# swings ~2x with the input seed, and applies to dx as well, whose real error is 4-5 +# decades smaller. Relative RMS stays 1e-6..8e-6 across the whole set. ALMIOPEN-2561. +[[validator_overrides]] +filters = ["*LayernormBackward*"] +tensors = ["*::DSCALE", "*::DBIAS"] +validator = "rms" +rms_threshold = 1e-4 +# # RMSNormBackward Pure-Bfp16 (dy/x/dx all bf16) fails CPU-reference tolerance # at the suite's large-reduction shapes (x:[4096,128,1,1], x:[5,256,14,14], # x:[16,3,8,14,14]) — smaller shapes in the same suites pass reliably. The diff --git a/dnn-providers/integration-tests/README.md b/dnn-providers/integration-tests/README.md index 096e5fc74d0e..10596b76ff98 100644 --- a/dnn-providers/integration-tests/README.md +++ b/dnn-providers/integration-tests/README.md @@ -388,7 +388,7 @@ Both the superbuild (target already present) and standalone provider builds (`find_package`) are supported; if the package is not found the target is skipped with a status message. -### Per-provider TOML config (tolerance overrides & skips) +### Per-provider TOML config (tolerance overrides, validator overrides & skips) Each provider owns one `--test-config` TOML file (e.g. `miopen-provider/config/MIOPEN_ENGINE.toml`, @@ -399,6 +399,9 @@ recompiling or touching test source: - **Override tolerances** for specific tests/groups, when that engine's numerics legitimately differ from the default atol/rtol (e.g. reduced precision from split-k accumulation). +- **Override the validator** for specific output tensors, when per-element + allclose is the wrong *question* for that tensor rather than merely too + tight. This is the only place a validator can be changed. - **Skip tests** on specific architectures (or globally), when that engine has no applicable kernel/solution for a case. @@ -411,6 +414,12 @@ filters = ["Smoke/IntegrationGpuConvWrw3dBfp16.Correctness/14"] atol = 1.19 rtol = 0.2 +[[validator_overrides]] +filters = ["*LayernormBackward*"] +tensors = ["*::DSCALE", "*::DBIAS"] +validator = "rms" +rms_threshold = 1e-4 + [[test_skips]] archs = ["gfx90a", "gfx10", "gfx11", "gfx12"] # optional; omit to skip everywhere filters = ["*ConvFwdBiasActiv*"] @@ -421,12 +430,30 @@ reason = "ROCm/rocm-libraries#6979 — no engine has an applicable solution for GTest name — same string a `--gtest_filter` would match. - `tolerance_overrides`: later entries take precedence when multiple filters match. Both `atol` and `rtol` are required. +- `validator_overrides`: later entries take precedence. An entry applies only + when a `filters` glob matches the test name **and** a `tensors` glob matches + the output tensor's label — its name (e.g. `LayernormBackward_0::DSCALE`), or + `uid=N` when the graph did not name it. Match on the tensor label rather than + the uid: uids differ between a C++ graph test and the bundle captured from it, + names do not. `validator` is `"allclose"` or `"rms"`; `rms_threshold` is + required and must be positive when the validator is `"rms"`. Absent any match + the comparison is allclose — **allclose is the default everywhere, and this + section is the only thing that changes it.** Use it when a per-element check + is the wrong question, not to buy slack: an output that is a long reduction + (layernorm/RMSNorm backward `dscale`/`dbias`) has elements that land + arbitrarily near zero through cancellation, so per-element relative error is + unbounded while the aggregate relative-RMS error is not. Prefer + `tolerance_overrides` for everything else. See ALMIOPEN-2561. - `test_skips`: the first matching entry wins; `reason` is surfaced in the `GTEST_SKIP` message. `archs` (substring match against the raw `gcnArchName`) and `platforms` (`"windows"`/`"linux"`) are both optional — omit either to match any. - Applies to **both** bundle/sweep tests and C++ graph tests; the lookup runs in the shared harness (`TestConfig`/`TestSettings`), not per test type. + One caveat for `validator_overrides`: golden data validated against a *reference* + executor (`BundleReferenceValidationHarness`) is always compared with 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. - `[meta] version = 1` is required; the file is rejected on parse if missing or on an unsupported version. diff --git a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp index 9d10c3fb6ab0..c1414b8036fa 100644 --- a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp +++ b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp @@ -265,6 +265,10 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam attr, float absoluteTolerance, float relativeTolerance) @@ -273,40 +277,33 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParamget_data_type()); - auto [it, inserted] = _tensorValidationMap.insert( - {attr->get_uid(), - TensorValidationEntry{hipdnn_test_sdk::utilities::createAllCloseValidator( - sdkDataType, finalAtol, finalRtol), - attr->get_name(), - finalAtol, - finalRtol, - sdkDataType}}); - if(!inserted) + + auto validator = hipdnn_test_sdk::utilities::createAllCloseValidator( + sdkDataType, finalAtol, finalRtol); + float entryAtol = finalAtol; + float entryRtol = finalRtol; + + const auto validatorOverride + = TestConfig::get().findValidatorOverride(testName, attr->get_name()); + if(validatorOverride && validatorOverride->kind == ValidatorOverrideKind::RMS) { - ADD_FAILURE() << "Duplicate validator for tensor " << attr->get_uid() << " (" - << attr->get_name() << "); keeping first registration"; + HIPDNN_PLUGIN_LOG_INFO("Validator override applied for " + << testName << " tensor " << attr->get_name() + << ": rms, threshold=" << validatorOverride->rmsThreshold); + validator = hipdnn_test_sdk::utilities::createRmsValidator( + sdkDataType, validatorOverride->rmsThreshold); + entryAtol = 0.0f; + entryRtol = 0.0f; } - _tensorIdToNameMap.insert({attr->get_uid(), attr->get_name()}); - }); - } - void registerRmsValidator(const std::shared_ptr attr, - float rmsThreshold) - { - _deferredValidators.emplace_back([this, attr, rmsThreshold]() { - auto sdkDataType - = hipdnn_test_sdk::utilities::frontendToSdkDataType(attr->get_data_type()); auto [it, inserted] = _tensorValidationMap.insert( {attr->get_uid(), TensorValidationEntry{ - hipdnn_test_sdk::utilities::createRmsValidator(sdkDataType, rmsThreshold), - attr->get_name(), - 0.0f, - 0.0f, - sdkDataType}}); + std::move(validator), attr->get_name(), entryAtol, entryRtol, sdkDataType}}); if(!inserted) { ADD_FAILURE() << "Duplicate validator for tensor " << attr->get_uid() << " (" @@ -434,8 +431,12 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParamfindToleranceOverride(testName); } + // Find a validator override for one output tensor of the given test. + // Returns std::nullopt if no config loaded or nothing matches, which means the + // default allclose comparison. + std::optional findValidatorOverride(std::string_view testName, + std::string_view tensorLabel) const + { + throwIfNotInitialized(); + if(!_testSettings.has_value()) + { + return std::nullopt; + } + return _testSettings->findValidatorOverride(testName, tensorLabel); + } + // Raw gcnArchName for device 0 detected at init time (e.g. // "gfx942:sramecc+:xnack-"). Empty if detection failed. const std::string& getCurrentArch() const diff --git a/dnn-providers/integration-tests/src/harness/TestSettings.hpp b/dnn-providers/integration-tests/src/harness/TestSettings.hpp index f87bfe76d0d5..0f2acdd7d2ac 100644 --- a/dnn-providers/integration-tests/src/harness/TestSettings.hpp +++ b/dnn-providers/integration-tests/src/harness/TestSettings.hpp @@ -38,6 +38,22 @@ struct ToleranceOverride float rtol; }; +// Which comparison an output tensor is graded by. Mirrors +// bundle::ValidatorKind, kept separate so TestSettings stays free of the +// comparison headers. +enum class ValidatorOverrideKind +{ + ALLCLOSE, + RMS, +}; + +struct ValidatorOverride +{ + ValidatorOverrideKind kind; + // Relative-RMS threshold. Meaningful only when kind == RMS. + float rmsThreshold; +}; + // Loads a per-engine TOML settings file for integration tests. // Currently supports tolerance overrides and arch-scoped test skips; // additional settings (knobs, support matrix, etc.) will be added in @@ -52,6 +68,12 @@ struct ToleranceOverride // atol = 1e-3 // rtol = 1e-2 // +// [[validator_overrides]] +// filters = ["*LayernormBackward*"] +// tensors = ["*::DSCALE", "*::DBIAS"] +// validator = "rms" +// rms_threshold = 1e-4 +// // [[test_skips]] // archs = ["gfx1100", "gfx1101"] // platforms = ["windows"] # optional @@ -66,6 +88,11 @@ struct ToleranceOverride // Filters use GTest-style globs (globMatch with * wildcards). // For tolerance_overrides, later entries take precedence over earlier // ones when multiple filters match. +// For validator_overrides, an entry applies when a 'filters' glob matches the +// gtest-formatted test name AND a 'tensors' glob matches the output tensor's +// label (its name, or "uid=N" when the graph did not name it). Later entries +// take precedence. Absent any match, allclose is used — allclose is the default +// and nothing else can select a validator. See ALMIOPEN-2561. // For test_skips, an entry matches when ALL of: // - 'archs' is omitted/empty (any arch), OR any 'archs' value is a // substring of the device's raw gcnArchName. @@ -90,6 +117,14 @@ class TestSettings } } + if(auto* arr = table["validator_overrides"].as_array()) + { + for(const auto& node : *arr) + { + _validatorOverrides.push_back(parseValidatorOverride(node)); + } + } + if(auto* arr = table["test_skips"].as_array()) { for(const auto& node : *arr) @@ -127,6 +162,44 @@ class TestSettings return _overrides.size(); } + // Find a validator override for one output tensor of one test. + // Returns the last matching entry (later entries take precedence), or + // std::nullopt when nothing matches — which means allclose. + std::optional findValidatorOverride(std::string_view testName, + std::string_view tensorLabel) const + { + std::optional result; + const std::string testNameStr(testName); + const std::string tensorLabelStr(tensorLabel); + + for(const auto& entry : _validatorOverrides) + { + const bool testMatches + = std::any_of(entry.filters.begin(), entry.filters.end(), [&](const auto& f) { + return globMatch(f, testNameStr); + }); + if(!testMatches) + { + continue; + } + const bool tensorMatches + = std::any_of(entry.tensors.begin(), entry.tensors.end(), [&](const auto& t) { + return globMatch(t, tensorLabelStr); + }); + if(tensorMatches) + { + result = ValidatorOverride{entry.kind, entry.rmsThreshold}; + } + } + + return result; + } + + size_t validatorOverrideCount() const + { + return _validatorOverrides.size(); + } + // Find a skip rule matching the given test name, device, and platform. // deviceArchRaw is the unparsed gcnArchName string (e.g. "gfx942:sramecc+:xnack-"). // platform is the lowercase platform name ("windows" or "linux"). @@ -201,6 +274,14 @@ class TestSettings float rtol; }; + struct ValidatorEntry + { + std::vector filters; + std::vector tensors; + ValidatorOverrideKind kind; + float rmsThreshold; + }; + struct SkipEntry { std::vector archs; @@ -285,6 +366,67 @@ class TestSettings return parsed; } + static ValidatorEntry parseValidatorOverride(const toml::node& node) + { + static constexpr const char* K_SECTION = "TestSettings: [[validator_overrides]]"; + + const auto* table = node.as_table(); + if(table == nullptr) + { + throw std::runtime_error(std::string(K_SECTION) + " entry is not a table"); + } + + ValidatorEntry parsed; + parsed.filters = parseStringArray(*table, "filters", K_SECTION); + if(parsed.filters.empty()) + { + throw std::runtime_error(std::string(K_SECTION) + " entry missing 'filters' array"); + } + + parsed.tensors = parseStringArray(*table, "tensors", K_SECTION); + if(parsed.tensors.empty()) + { + throw std::runtime_error(std::string(K_SECTION) + + " entry missing 'tensors' array — a validator override " + "names the tensors it applies to, never a whole test"); + } + + auto validator = (*table)["validator"].value(); + if(!validator.has_value()) + { + throw std::runtime_error(std::string(K_SECTION) + " entry missing 'validator'"); + } + + if(*validator == "allclose") + { + parsed.kind = ValidatorOverrideKind::ALLCLOSE; + parsed.rmsThreshold = 0.0f; + return parsed; + } + if(*validator != "rms") + { + throw std::runtime_error(std::string(K_SECTION) + " entry has unknown validator '" + + *validator + R"~(' (expected "allclose" or "rms"))~"); + } + + parsed.kind = ValidatorOverrideKind::RMS; + auto threshold = (*table)["rms_threshold"].value(); + if(!threshold.has_value()) + { + throw std::runtime_error( + std::string(K_SECTION) + + R"~( entry with validator = "rms" missing 'rms_threshold')~"); + } + if(!(*threshold > 0.0)) + { + throw std::runtime_error(std::string(K_SECTION) + " entry has non-positive " + + "'rms_threshold'"); + } + parsed.rmsThreshold = static_cast(*threshold); + + return parsed; + } + static SkipEntry parseTestSkip(const toml::node& node) { static constexpr const char* K_SECTION = "TestSettings: [[test_skips]]"; @@ -316,6 +458,7 @@ class TestSettings } std::vector _overrides; + std::vector _validatorOverrides; std::vector _skips; }; diff --git a/dnn-providers/integration-tests/src/harness/bundle/BundleReferenceValidationHarness.cpp b/dnn-providers/integration-tests/src/harness/bundle/BundleReferenceValidationHarness.cpp index 54d7238cc11e..cfc829162cfa 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/BundleReferenceValidationHarness.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/BundleReferenceValidationHarness.cpp @@ -127,10 +127,13 @@ void BundleReferenceValidationHarness::TestBody() }; // defaultTolerance(), never resolveTolerance(): a TOML override belongs to an - // engine and must not loosen the gate on our own data. - const auto toleranceFor = [&wrapper](hipdnn_flatbuffers_sdk::data_objects::DataType dataType) { + // engine and must not loosen the gate on our own data. For the same reason this + // path never selects a non-default validator — allclose always (ALMIOPEN-2561). + const auto toleranceFor = [&wrapper](int64_t /*uid*/, + const std::string& /*label*/, + hipdnn_flatbuffers_sdk::data_objects::DataType dataType) { const float value = tolerance::defaultTolerance(wrapper, dataType); - return ComparisonTolerance{value, value}; + return ComparisonTolerance::allClose(value, value); }; const std::string contextLine = "Golden data validation (" diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp index 2e51648109c9..5ab94dd1fb1a 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp @@ -578,7 +578,21 @@ VerificationOutcome << " rtol=" << tomlOverride->rtol); } - const auto toleranceFor = [&](hipdnn_flatbuffers_sdk::data_objects::DataType dataType) { + const auto toleranceFor = [&](int64_t /*uid*/, + const std::string& label, + hipdnn_flatbuffers_sdk::data_objects::DataType dataType) { + // allclose unless this engine's TOML names the tensor. Nothing else selects a + // validator — see ALMIOPEN-2561. + const auto validatorOverride + = TestConfig::get().findValidatorOverride(currentTestName(), label); + if(validatorOverride && validatorOverride->kind == ValidatorOverrideKind::RMS) + { + HIPDNN_PLUGIN_LOG_INFO("Validator override applied for " + << currentTestName() << " tensor " << label + << ": rms, threshold=" << validatorOverride->rmsThreshold); + return ComparisonTolerance::rms(validatorOverride->rmsThreshold); + } + ComparisonTolerance tolerance; tolerance::resolveTolerance( wrapper, dataType, currentTestName(), tolerance.atol, tolerance.rtol); diff --git a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp index 4f870fd88113..2ab2259c454d 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp @@ -6,6 +6,7 @@ #include #include +#include #include namespace hipdnn_integration_tests::bundle @@ -31,9 +32,12 @@ std::optional const std::string& contextLine) { const auto dataType = attrs.data_type(); + const bool useRms = tolerance.kind == ValidatorKind::RMS; - auto validator = hipdnn_test_sdk::utilities::createAllCloseValidator( - dataType, tolerance.atol, tolerance.rtol); + auto validator + = useRms ? hipdnn_test_sdk::utilities::createRmsValidator(dataType, tolerance.rmsThreshold) + : hipdnn_test_sdk::utilities::createAllCloseValidator( + dataType, tolerance.atol, tolerance.rtol); if(validator->allClose(expected, actual)) { return std::nullopt; @@ -47,23 +51,38 @@ std::optional ctx.dtypeName = hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType); ctx.atol = tolerance.atol; ctx.rtol = tolerance.rtol; + if(useRms) + { + // atol/rtol did not decide this failure, so do not print them as if they had. + std::ostringstream summary; + summary << "relative RMS <= " << tolerance.rmsThreshold + << " (aggregate check — the element counts below are elements that " + "differ at all, not elements that failed)"; + ctx.toleranceSummary = summary.str(); + } std::ostringstream report; report << hipdnn_test_sdk::utilities::formatComparisonHeader(ctx, expected); - hipdnn_test_sdk::utilities::appendComparisonDiffByDataType( - report, dataType, label, expected, actual, tolerance.atol, tolerance.rtol); + // Zero tolerances under RMS: the per-element budget is not what was checked, and a + // full drift profile (max/mean abs diff, worst elements) is the useful diagnostic. + hipdnn_test_sdk::utilities::appendComparisonDiffByDataType(report, + dataType, + label, + expected, + actual, + useRms ? 0.0f : tolerance.atol, + useRms ? 0.0f : tolerance.rtol); return TensorMismatch{uid, label, report.str()}; } -std::vector compareOutputs( - const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper& wrapper, - const std::vector& outputUids, - OutputTensors& actual, - const ExpectedTensorLookup& expectedFor, - const std::function& - toleranceFor, - const std::string& contextLine) +std::vector + compareOutputs(const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper& wrapper, + const std::vector& outputUids, + OutputTensors& actual, + const ExpectedTensorLookup& expectedFor, + const ToleranceLookup& toleranceFor, + const std::string& contextLine) { const auto& tensorAttrMap = wrapper.getTensorMap(); @@ -71,12 +90,13 @@ std::vector compareOutputs( for(const int64_t uid : outputUids) { const auto* attrs = tensorAttrMap.at(uid); - auto mismatch = compareTensor(uid, - *attrs, - expectedFor(uid), - *actual.at(uid), - toleranceFor(attrs->data_type()), - contextLine); + auto mismatch + = compareTensor(uid, + *attrs, + expectedFor(uid), + *actual.at(uid), + toleranceFor(uid, tensorLabel(uid, *attrs), attrs->data_type()), + contextLine); if(mismatch.has_value()) { mismatches.push_back(*std::move(mismatch)); diff --git a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp index a2b70fb7cd19..cf7ae2a01036 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp @@ -37,14 +37,53 @@ struct TensorMismatch /// bundle, or a reference executor's own output buffers. using ExpectedTensorLookup = std::function; -/// The tolerances one tensor is compared at. Resolved per data type, and overridable -/// per test from the TOML config. +/// How one output tensor is graded. +/// +/// ALLCLOSE is the default and stays the default: per-element +/// |ref - impl| <= atol + rtol*|ref|. RMS is MIOpen's aggregate relative-RMS check, +/// which normalises by the tensor's largest magnitude instead of each element's own. +/// It exists for reduction outputs whose elements can land arbitrarily close to zero +/// through cancellation — layernorm/RMSNorm backward dscale/dbias — where per-element +/// relative error is unbounded while the aggregate error is not. See ALMIOPEN-2561. +/// +/// Nothing selects RMS on its own: only an engine's TOML config can, via a +/// [[validator_overrides]] entry naming the tensor. +enum class ValidatorKind +{ + ALLCLOSE, + RMS, +}; + +/// How one tensor is compared. Resolved per output tensor, and overridable per test +/// from the TOML config. struct ComparisonTolerance { float atol = 0.0f; float rtol = 0.0f; + ValidatorKind kind = ValidatorKind::ALLCLOSE; + /// Relative-RMS threshold. Read only when `kind == RMS`; atol/rtol are ignored then. + float rmsThreshold = 0.0f; + + static ComparisonTolerance allClose(float atolIn, float rtolIn) + { + return ComparisonTolerance{atolIn, rtolIn, ValidatorKind::ALLCLOSE, 0.0f}; + } + + static ComparisonTolerance rms(float threshold) + { + return ComparisonTolerance{0.0f, 0.0f, ValidatorKind::RMS, threshold}; + } }; +/// Resolves how one output tensor is compared. +/// +/// `label` is the tensor's name, or "uid=N" when the graph did not give it one — the +/// same string the failure report uses, and what a TOML `tensors` glob matches against. +using ToleranceLookup + = std::function; + /// Compare one tensor. Returns nullopt when it matched. /// /// Pure: no gtest, no config lookups, no harness state. Everything it needs to @@ -60,16 +99,15 @@ std::optional /// Compare every uid in `outputUids`, and keep going after the first mismatch: one /// failing test should name every tensor that drifted, not just the lowest uid. /// -/// `toleranceFor` resolves the tolerance per data type, so the caller owns the TOML +/// `toleranceFor` resolves how each output is compared, so the caller owns the TOML /// override and this stays free of TestConfig. -std::vector compareOutputs( - const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper& wrapper, - const std::vector& outputUids, - OutputTensors& actual, - const ExpectedTensorLookup& expectedFor, - const std::function& - toleranceFor, - const std::string& contextLine); +std::vector + compareOutputs(const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper& wrapper, + const std::vector& outputUids, + OutputTensors& actual, + const ExpectedTensorLookup& expectedFor, + const ToleranceLookup& toleranceFor, + const std::string& contextLine); /// The tensor's name, or "uid=N" when the graph did not give it one. std::string tensorLabel(int64_t uid, diff --git a/dnn-providers/integration-tests/src/integration-tests/layernorm/IntegrationGpuLayernormBackward.cpp b/dnn-providers/integration-tests/src/integration-tests/layernorm/IntegrationGpuLayernormBackward.cpp index bce17bd25ed8..aa0690f6e740 100644 --- a/dnn-providers/integration-tests/src/integration-tests/layernorm/IntegrationGpuLayernormBackward.cpp +++ b/dnn-providers/integration-tests/src/integration-tests/layernorm/IntegrationGpuLayernormBackward.cpp @@ -160,9 +160,12 @@ class LayernormBackward auto [graphObj, outputs] = buildGraph(getSharedHandle(), testCase); this->registerValidator(outputs.dx, this->getTolerance(graphObj, outputs.dx)); - // RMS validator as the standard validator breaks down for resulting elements that happen to be near zero after summing hundreds of thousands of floating point values - this->registerRmsValidator(outputs.dscale, this->getTolerance(graphObj, outputs.dscale)); - this->registerRmsValidator(outputs.dbias, this->getTolerance(graphObj, outputs.dbias)); + // dscale/dbias are reductions over the batch axes: elements land arbitrarily + // near zero through cancellation, so allclose is the wrong check for them. The + // engine's TOML config selects the RMS validator for those two tensors — + // see [[validator_overrides]] in HIP_MLOPS_ENGINE.toml, ALMIOPEN-2561. + this->registerValidator(outputs.dscale, this->getTolerance(graphObj, outputs.dscale)); + this->registerValidator(outputs.dbias, this->getTolerance(graphObj, outputs.dbias)); this->inputFillRecipes().setGlobalSeed(layernormTestCase.seed); this->verifyGraph(graphObj); diff --git a/dnn-providers/integration-tests/tests/TestOutputComparison.cpp b/dnn-providers/integration-tests/tests/TestOutputComparison.cpp index 3ae7d2edd7db..ce8ff7e61da4 100644 --- a/dnn-providers/integration-tests/tests/TestOutputComparison.cpp +++ b/dnn-providers/integration-tests/tests/TestOutputComparison.cpp @@ -12,6 +12,8 @@ #include #include +#include + #include "harness/bundle/IntegrationTestBundle.hpp" #include "harness/bundle/OutputComparison.hpp" @@ -67,11 +69,30 @@ std::unique_ptr return tensor; } +// A [1, 3, 1, 1] tensor with the three values written individually, so a test can build +// the magnitude spread that separates allclose from RMS. +std::unique_ptr + floatTensor3(const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes& attrs, + float v0, + float v1, + float v2) +{ + auto tensor = hipdnn_test_sdk::detail::createTensorFromAttribute(attrs); + hipdnn_data_sdk::utilities::TensorView view(*tensor); + view.getHostValue({0, 0, 0, 0}) = v0; + view.getHostValue({0, 1, 0, 0}) = v1; + view.getHostValue({0, 2, 0, 0}) = v2; + tensor->markHostModified(); + return tensor; +} + ComparisonTolerance exact() { - return ComparisonTolerance{0.0f, 0.0f}; + return ComparisonTolerance::allClose(0.0f, 0.0f); } +constexpr float K_RMS_THRESHOLD = 1e-4f; + } // namespace // --------------------------------------------------------------------------- @@ -150,7 +171,8 @@ TEST(TestOutputComparison, ToleranceDecidesWhetherADifferenceMatters) EXPECT_TRUE(compareTensor(K_UID_A, attrs, *expected, *actual, exact(), "b").has_value()); EXPECT_FALSE( - compareTensor(K_UID_A, attrs, *expected, *actual, ComparisonTolerance{0.1f, 0.1f}, "b") + compareTensor( + K_UID_A, attrs, *expected, *actual, ComparisonTolerance::allClose(0.1f, 0.1f), "b") .has_value()); } @@ -179,7 +201,7 @@ TEST(TestOutputComparison, AllOutputsMatchingYieldsNoMismatches) {K_UID_A, K_UID_B}, actual, [&](int64_t uid) -> hipdnn_data_sdk::utilities::ITensor& { return *expected.at(uid); }, - [](auto) { return exact(); }, + [](int64_t, const std::string&, auto) { return exact(); }, "Bundle: b"); EXPECT_TRUE(mismatches.empty()); @@ -205,7 +227,7 @@ TEST(TestOutputComparison, EveryDriftedTensorIsReportedNotJustTheFirst) {K_UID_A, K_UID_B}, actual, [&](int64_t uid) -> hipdnn_data_sdk::utilities::ITensor& { return *expected.at(uid); }, - [](auto) { return exact(); }, + [](int64_t, const std::string&, auto) { return exact(); }, "Bundle: b"); ASSERT_EQ(mismatches.size(), 2u); @@ -235,10 +257,114 @@ TEST(TestOutputComparison, OnlyTheRequestedUidsAreCompared) {K_UID_A}, actual, [&](int64_t uid) -> hipdnn_data_sdk::utilities::ITensor& { return *expected.at(uid); }, - [](auto) { return exact(); }, + [](int64_t, const std::string&, auto) { return exact(); }, "Bundle: b"); EXPECT_TRUE(mismatches.empty()) << "uid 4 drifted but was not in the list"; } +// --------------------------------------------------------------------------- +// Validator kind. ALLCLOSE grades each element against its own magnitude; RMS grades +// the tensor against its largest. The two disagree exactly where reduction outputs +// live: an element near zero, on a tensor whose scale is large. ALMIOPEN-2561. +// --------------------------------------------------------------------------- + +TEST(TestOutputComparison, RmsAcceptsANearZeroElementThatAllcloseRejects) +{ + const auto buffer = makeGraphBuffer(); + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper wrapper{buffer.data(), + buffer.size()}; + const auto& attrs = *wrapper.getTensorMap().at(K_UID_B); + + // Middle element is 5 orders of magnitude below the tensor scale and drifts by 1e-4: + // negligible against the tensor, unbounded against itself. + auto expected = floatTensor3(attrs, 1000.0f, 0.001f, -1000.0f); + auto actual = floatTensor3(attrs, 1000.0f, 0.0011f, -1000.0f); + + EXPECT_TRUE( + compareTensor( + K_UID_B, attrs, *expected, *actual, ComparisonTolerance::allClose(0.0f, 1e-4f), "b") + .has_value()) + << "allclose should reject: rtol*|ref| is 1e-7 against a 1e-4 drift"; + + EXPECT_FALSE( + compareTensor( + K_UID_B, attrs, *expected, *actual, ComparisonTolerance::rms(K_RMS_THRESHOLD), "b") + .has_value()) + << "relative RMS is ~6e-8 against a 1e-4 threshold"; +} + +TEST(TestOutputComparison, RmsStillRejectsDriftLargeAgainstTheTensorScale) +{ + const auto buffer = makeGraphBuffer(); + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper wrapper{buffer.data(), + buffer.size()}; + const auto& attrs = *wrapper.getTensorMap().at(K_UID_B); + + auto expected = floatTensor3(attrs, 1000.0f, 0.001f, -1000.0f); + auto actual = floatTensor3(attrs, 1000.0f, 0.001f, -900.0f); + + EXPECT_TRUE( + compareTensor( + K_UID_B, attrs, *expected, *actual, ComparisonTolerance::rms(K_RMS_THRESHOLD), "b") + .has_value()) + << "RMS is a real check, not a pass-through"; +} + +TEST(TestOutputComparison, RmsFailureReportsItsThresholdNotAtolRtol) +{ + const auto buffer = makeGraphBuffer(); + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper wrapper{buffer.data(), + buffer.size()}; + const auto& attrs = *wrapper.getTensorMap().at(K_UID_B); + + auto expected = floatTensor3(attrs, 1000.0f, 0.001f, -1000.0f); + auto actual = floatTensor3(attrs, 1000.0f, 0.001f, -900.0f); + + const auto mismatch = compareTensor( + K_UID_B, attrs, *expected, *actual, ComparisonTolerance::rms(K_RMS_THRESHOLD), "b"); + + ASSERT_TRUE(mismatch.has_value()); + EXPECT_NE(mismatch->report.find("relative RMS"), std::string::npos); + EXPECT_EQ(mismatch->report.find("atol="), std::string::npos) + << "atol/rtol did not decide this failure and must not be printed as if they had"; +} + +// The point of the lookup taking a uid and a label: one graph, two outputs, two +// validators. This is what an engine TOML [[validator_overrides]] entry drives. +TEST(TestOutputComparison, ValidatorKindIsChosenPerTensor) +{ + const auto buffer = makeGraphBuffer(); + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper wrapper{buffer.data(), + buffer.size()}; + const auto& map = wrapper.getTensorMap(); + + // uid 4 is the near-zero reduction-like output; uid 5 ("y_out") drifts outright. + OutputTensors expected; + expected[K_UID_A] = floatTensor(*map.at(K_UID_A), 1.0f); + expected[K_UID_B] = floatTensor3(*map.at(K_UID_B), 1000.0f, 0.001f, -1000.0f); + + OutputTensors actual; + actual[K_UID_A] = floatTensor(*map.at(K_UID_A), 99.0f); + actual[K_UID_B] = floatTensor3(*map.at(K_UID_B), 1000.0f, 0.0011f, -1000.0f); + + std::vector labelsSeen; + const auto mismatches = compareOutputs( + wrapper, + {K_UID_A, K_UID_B}, + actual, + [&](int64_t uid) -> hipdnn_data_sdk::utilities::ITensor& { return *expected.at(uid); }, + [&](int64_t, const std::string& label, auto) { + labelsSeen.push_back(label); + return label == "uid=4" ? ComparisonTolerance::rms(K_RMS_THRESHOLD) + : ComparisonTolerance::allClose(0.0f, 1e-4f); + }, + "Bundle: b"); + + EXPECT_EQ(labelsSeen, (std::vector{"y_out", "uid=4"})) + << "the lookup must be given the label a TOML glob would match on"; + ASSERT_EQ(mismatches.size(), 1u) << "uid 4 passes under RMS; uid 5 fails under allclose"; + EXPECT_EQ(mismatches[0].uid, K_UID_A); +} + // NOLINTEND(readability-identifier-naming) diff --git a/dnn-providers/integration-tests/tests/TestTestSettings.cpp b/dnn-providers/integration-tests/tests/TestTestSettings.cpp index 8837d2d1c6b7..c8e8b1ddc9cd 100644 --- a/dnn-providers/integration-tests/tests/TestTestSettings.cpp +++ b/dnn-providers/integration-tests/tests/TestTestSettings.cpp @@ -295,6 +295,176 @@ rtol = 5e-2 EXPECT_FLOAT_EQ(result->rtol, 1e-2F); } +// --------------------------------------------------------------------------- +// [[validator_overrides]] +// +// allclose is the default everywhere and this section is the only thing that can +// change it, so the parser is strict: an entry that does not say exactly what it +// means is a load error, not a silent fall-back to allclose. ALMIOPEN-2561. +// --------------------------------------------------------------------------- + +namespace +{ + +const char* const K_RMS_CONFIG = R"( +[meta] +version = 1 + +[[validator_overrides]] +filters = ["*LayernormBackward*"] +tensors = ["*::DSCALE", "*::DBIAS"] +validator = "rms" +rms_threshold = 1e-4 +)"; + +} // namespace + +TEST(TestSettingsValidatorOverrides, NoOverridesMeansNoOverride) +{ + const TempTomlFile file(R"( +[meta] +version = 1 +)"); + + const TestSettings settings(file.path()); + EXPECT_EQ(settings.validatorOverrideCount(), 0U); + EXPECT_FALSE(settings.findValidatorOverride("AnyTest", "AnyTensor").has_value()); +} + +TEST(TestSettingsValidatorOverrides, MatchesOnBothTestNameAndTensorLabel) +{ + const TempTomlFile file(K_RMS_CONFIG); + const TestSettings settings(file.path()); + + ASSERT_EQ(settings.validatorOverrideCount(), 1U); + + const auto hit = settings.findValidatorOverride( + "Smoke/IntegrationGpuLayernormBackwardPure5DFp32.Correctness/9", + "LayernormBackward_0::DSCALE"); + ASSERT_TRUE(hit.has_value()); + EXPECT_EQ(hit->kind, hipdnn_integration_tests::ValidatorOverrideKind::RMS); + EXPECT_FLOAT_EQ(hit->rmsThreshold, 1e-4F); +} + +// The tensor glob is what keeps dx on allclose while dscale moves to RMS. Matching +// on the test name alone would take the whole graph with it. +TEST(TestSettingsValidatorOverrides, OtherTensorsInTheSameTestAreUnaffected) +{ + const TempTomlFile file(K_RMS_CONFIG); + const TestSettings settings(file.path()); + + EXPECT_FALSE( + settings + .findValidatorOverride("Smoke/IntegrationGpuLayernormBackwardPure5DFp32.Correctness/9", + "LayernormBackward_0::DX") + .has_value()); +} + +TEST(TestSettingsValidatorOverrides, OtherTestsWithTheSameTensorNameAreUnaffected) +{ + const TempTomlFile file(K_RMS_CONFIG); + const TestSettings settings(file.path()); + + EXPECT_FALSE( + settings + .findValidatorOverride("Smoke/IntegrationGpuRMSNormBackwardPureFp32.Correctness/0", + "RMSNormBackward_0::DSCALE") + .has_value()); +} + +TEST(TestSettingsValidatorOverrides, LaterEntriesTakePrecedence) +{ + const TempTomlFile file(R"( +[meta] +version = 1 + +[[validator_overrides]] +filters = ["*LayernormBackward*"] +tensors = ["*::DSCALE"] +validator = "rms" +rms_threshold = 1e-4 + +[[validator_overrides]] +filters = ["*LayernormBackward*Bfp16*"] +tensors = ["*::DSCALE"] +validator = "allclose" +)"); + + const TestSettings settings(file.path()); + + const auto fp32 + = settings.findValidatorOverride("Smoke/LayernormBackwardFp32.C/0", "x::DSCALE"); + ASSERT_TRUE(fp32.has_value()); + EXPECT_EQ(fp32->kind, hipdnn_integration_tests::ValidatorOverrideKind::RMS); + + const auto bf16 + = settings.findValidatorOverride("Smoke/LayernormBackwardBfp16.C/0", "x::DSCALE"); + ASSERT_TRUE(bf16.has_value()); + EXPECT_EQ(bf16->kind, hipdnn_integration_tests::ValidatorOverrideKind::ALLCLOSE); +} + +TEST(TestSettingsValidatorOverrides, ThrowsOnMissingTensors) +{ + const TempTomlFile file(R"( +[meta] +version = 1 + +[[validator_overrides]] +filters = ["*LayernormBackward*"] +validator = "rms" +rms_threshold = 1e-4 +)"); + + EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); +} + +TEST(TestSettingsValidatorOverrides, ThrowsOnUnknownValidatorKind) +{ + const TempTomlFile file(R"( +[meta] +version = 1 + +[[validator_overrides]] +filters = ["*LayernormBackward*"] +tensors = ["*::DSCALE"] +validator = "relative-rms" +rms_threshold = 1e-4 +)"); + + EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); +} + +TEST(TestSettingsValidatorOverrides, ThrowsOnRmsWithoutThreshold) +{ + const TempTomlFile file(R"( +[meta] +version = 1 + +[[validator_overrides]] +filters = ["*LayernormBackward*"] +tensors = ["*::DSCALE"] +validator = "rms" +)"); + + EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); +} + +TEST(TestSettingsValidatorOverrides, ThrowsOnNonPositiveRmsThreshold) +{ + const TempTomlFile file(R"( +[meta] +version = 1 + +[[validator_overrides]] +filters = ["*LayernormBackward*"] +tensors = ["*::DSCALE"] +validator = "rms" +rms_threshold = 0.0 +)"); + + EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); +} + // --------------------------------------------------------------------------- // [[test_skips]] parsing // --------------------------------------------------------------------------- diff --git a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/ComparisonReport.hpp b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/ComparisonReport.hpp index 178cdbcb7ffd..f79ce42a9368 100644 --- a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/ComparisonReport.hpp +++ b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/ComparisonReport.hpp @@ -22,6 +22,9 @@ struct ComparisonContext std::string dtypeName; float atol = 0.0f; float rtol = 0.0f; + /// Replaces the "atol=… rtol=…" line when set, for checks that atol/rtol do not + /// describe (e.g. an aggregate relative-RMS comparison). + std::string toleranceSummary; }; inline std::string formatComparisonHeader(const ComparisonContext& ctx, @@ -31,8 +34,15 @@ inline std::string formatComparisonHeader(const ComparisonContext& ctx, os << "\nComparison FAILED\n" << " " << ctx.contextLine << "\n" << " Tensor: " << ctx.tensorLabel << "\n" - << " Shape: " << StreamVec(tensor.dims()) << " " << ctx.dtypeName << "\n" - << " Tolerance: atol=" << ctx.atol << " rtol=" << ctx.rtol << "\n"; + << " Shape: " << StreamVec(tensor.dims()) << " " << ctx.dtypeName << "\n"; + if(ctx.toleranceSummary.empty()) + { + os << " Tolerance: atol=" << ctx.atol << " rtol=" << ctx.rtol << "\n"; + } + else + { + os << " Tolerance: " << ctx.toleranceSummary << "\n"; + } return os.str(); } diff --git a/projects/hipdnn/test_sdk/tests/utilities/TestComparisonReport.cpp b/projects/hipdnn/test_sdk/tests/utilities/TestComparisonReport.cpp index 806a7833c136..1bc0539cc8a9 100644 --- a/projects/hipdnn/test_sdk/tests/utilities/TestComparisonReport.cpp +++ b/projects/hipdnn/test_sdk/tests/utilities/TestComparisonReport.cpp @@ -19,7 +19,7 @@ using DT = hipdnn_flatbuffers_sdk::data_objects::DataType; TEST(TestFormatComparisonHeader, ContainsAllFields) { const ComparisonContext ctx{ - "Bundle: /path/to/bundle", "output_tensor (UID 42, output)", "FLOAT", 1e-5f, 1e-4f}; + "Bundle: /path/to/bundle", "output_tensor (UID 42, output)", "FLOAT", 1e-5f, 1e-4f, ""}; const Tensor tensor({2, 3, 4}); const std::string header = formatComparisonHeader(ctx, tensor); @@ -34,7 +34,7 @@ TEST(TestFormatComparisonHeader, ContainsAllFields) TEST(TestFormatComparisonHeader, IncludesShape) { - const ComparisonContext ctx{"Test: MyTest.Case", "x", "HALF", 0.0f, 0.0f}; + const ComparisonContext ctx{"Test: MyTest.Case", "x", "HALF", 0.0f, 0.0f, ""}; const Tensor tensor({8, 16}); const std::string header = formatComparisonHeader(ctx, tensor); @@ -43,6 +43,21 @@ TEST(TestFormatComparisonHeader, IncludesShape) EXPECT_NE(header.find("8, 16"), std::string::npos); } +// A check that atol/rtol do not describe says so, instead of printing two numbers that +// had nothing to do with the verdict. +TEST(TestFormatComparisonHeader, ToleranceSummaryReplacesAtolRtol) +{ + const ComparisonContext ctx{ + "Test: MyTest.Case", "dscale", "FLOAT", 1e-5f, 1e-4f, "relative RMS <= 0.0001"}; + + const Tensor tensor({8, 16}); + const std::string header = formatComparisonHeader(ctx, tensor); + + EXPECT_NE(header.find("relative RMS <= 0.0001"), std::string::npos); + EXPECT_EQ(header.find("atol="), std::string::npos); + EXPECT_EQ(header.find("rtol="), std::string::npos); +} + // ================================================================================================= // appendComparisonDiffByDataType — typed dispatch // ================================================================================================= From a1ff3da1e6055a1b12b31acd6e83303ada6dbdec Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Tue, 8 Sep 2026 08:31:56 -0600 Subject: [PATCH 2/6] test(integration-tests): rename validator-override cases off the Test 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 --- .tmp_names.txt | 0 dnn-providers/integration-tests/tests/TestTestSettings.cpp | 6 +++--- 2 files changed, 3 insertions(+), 3 deletions(-) create mode 100644 .tmp_names.txt diff --git a/.tmp_names.txt b/.tmp_names.txt new file mode 100644 index 000000000000..e69de29bb2d1 diff --git a/dnn-providers/integration-tests/tests/TestTestSettings.cpp b/dnn-providers/integration-tests/tests/TestTestSettings.cpp index c8e8b1ddc9cd..34b0599a8608 100644 --- a/dnn-providers/integration-tests/tests/TestTestSettings.cpp +++ b/dnn-providers/integration-tests/tests/TestTestSettings.cpp @@ -331,7 +331,7 @@ version = 1 EXPECT_FALSE(settings.findValidatorOverride("AnyTest", "AnyTensor").has_value()); } -TEST(TestSettingsValidatorOverrides, MatchesOnBothTestNameAndTensorLabel) +TEST(TestSettingsValidatorOverrides, MatchesOnBothNameAndTensorGlob) { const TempTomlFile file(K_RMS_CONFIG); const TestSettings settings(file.path()); @@ -348,7 +348,7 @@ TEST(TestSettingsValidatorOverrides, MatchesOnBothTestNameAndTensorLabel) // The tensor glob is what keeps dx on allclose while dscale moves to RMS. Matching // on the test name alone would take the whole graph with it. -TEST(TestSettingsValidatorOverrides, OtherTensorsInTheSameTestAreUnaffected) +TEST(TestSettingsValidatorOverrides, OtherTensorsInTheSameCaseAreUnaffected) { const TempTomlFile file(K_RMS_CONFIG); const TestSettings settings(file.path()); @@ -360,7 +360,7 @@ TEST(TestSettingsValidatorOverrides, OtherTensorsInTheSameTestAreUnaffected) .has_value()); } -TEST(TestSettingsValidatorOverrides, OtherTestsWithTheSameTensorNameAreUnaffected) +TEST(TestSettingsValidatorOverrides, OtherSuitesWithTheSameTensorAreUnaffected) { const TempTomlFile file(K_RMS_CONFIG); const TestSettings settings(file.path()); From d3a0d139e63f8c02e149a76fccff3628ec655bb0 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Tue, 8 Sep 2026 16:27:03 -0600 Subject: [PATCH 3/6] refactor(integration-tests): share the validator selection and failure 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 --- .tmp_names.txt | 0 dnn-providers/integration-tests/README.md | 18 ++- .../IntegrationGraphVerificationHarness.hpp | 94 +++++-------- .../src/harness/TestSettings.hpp | 13 ++ .../src/harness/TomlGuards.hpp | 26 ++++ .../IntegrationBundleVerificationHarness.cpp | 9 +- .../src/harness/bundle/OutputComparison.cpp | 127 +++++++++++++----- .../src/harness/bundle/OutputComparison.hpp | 37 +++++ .../tests/TestOutputComparison.cpp | 86 +++++++++++- .../tests/TestTestSettings.cpp | 19 +++ .../tests/TestTomlGuards.cpp | 13 ++ 11 files changed, 336 insertions(+), 106 deletions(-) delete mode 100644 .tmp_names.txt diff --git a/.tmp_names.txt b/.tmp_names.txt deleted file mode 100644 index e69de29bb2d1..000000000000 diff --git a/dnn-providers/integration-tests/README.md b/dnn-providers/integration-tests/README.md index 10596b76ff98..f348d86b9314 100644 --- a/dnn-providers/integration-tests/README.md +++ b/dnn-providers/integration-tests/README.md @@ -436,13 +436,17 @@ reason = "ROCm/rocm-libraries#6979 — no engine has an applicable solution for `uid=N` when the graph did not name it. Match on the tensor label rather than the uid: uids differ between a C++ graph test and the bundle captured from it, names do not. `validator` is `"allclose"` or `"rms"`; `rms_threshold` is - required and must be positive when the validator is `"rms"`. Absent any match - the comparison is allclose — **allclose is the default everywhere, and this - section is the only thing that changes it.** Use it when a per-element check - is the wrong question, not to buy slack: an output that is a long reduction - (layernorm/RMSNorm backward `dscale`/`dbias`) has elements that land - arbitrarily near zero through cancellation, so per-element relative error is - unbounded while the aggregate relative-RMS error is not. Prefer + required and must be positive when the validator is `"rms"`, and must be + absent when it is `"allclose"` — an entry that does not say exactly what it + means is a load error, never a silent fall-back. `"rms"` is only defined for + float, half, bfloat16 and double outputs; a glob wide enough to catch an + integer output fails that tensor with a message naming the glob to narrow. + Absent any match the comparison is allclose — **allclose is the default + everywhere, and this section is the only thing that changes it.** Use it when + a per-element check is the wrong question, not to buy slack: an output that is + a long reduction (layernorm/RMSNorm backward `dscale`/`dbias`) has elements + that land arbitrarily near zero through cancellation, so per-element relative + error is unbounded while the aggregate relative-RMS error is not. Prefer `tolerance_overrides` for everything else. See ALMIOPEN-2561. - `test_skips`: the first matching entry wins; `reason` is surfaced in the `GTEST_SKIP` message. `archs` (substring match against the raw diff --git a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp index c1414b8036fa..e33779f164b6 100644 --- a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp +++ b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp @@ -20,9 +20,7 @@ #include #include -#include -#include -#include +#include #include #include #include @@ -39,6 +37,7 @@ #include "harness/SupportMatrixCollector.hpp" #include "harness/TestConfig.hpp" #include "harness/TomlGuards.hpp" +#include "harness/bundle/OutputComparison.hpp" #include "harness/input-init/FillInputs.hpp" #include "harness/input-init/InputFillRecipes.hpp" #include "harness/tolerance/ToleranceResolver.hpp" @@ -125,9 +124,11 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam validator; - std::string name; - float atol = 0.0f; - float rtol = 0.0f; + /// Why no validator could be built; non-empty exactly when `validator` is null. + std::string validatorError; + /// The tensor's name, or "uid=N" — the same label the TOML globs match on. + std::string label; + bundle::ComparisonTolerance tolerance; hipdnn_flatbuffers_sdk::data_objects::DataType dataType = hipdnn_flatbuffers_sdk::data_objects::DataType::UNSET; }; @@ -279,35 +280,28 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParamget_data_type()); + const auto label = bundle::tensorLabel(attr->get_uid(), attr->get_name()); - auto validator = hipdnn_test_sdk::utilities::createAllCloseValidator( - sdkDataType, finalAtol, finalRtol); - float entryAtol = finalAtol; - float entryRtol = finalRtol; - - const auto validatorOverride - = TestConfig::get().findValidatorOverride(testName, attr->get_name()); - if(validatorOverride && validatorOverride->kind == ValidatorOverrideKind::RMS) + auto tolerance = bundle::ComparisonTolerance::allClose(finalAtol, finalRtol); + if(const auto rmsThreshold = findTomlRmsThreshold(testName, label)) { - HIPDNN_PLUGIN_LOG_INFO("Validator override applied for " - << testName << " tensor " << attr->get_name() - << ": rms, threshold=" << validatorOverride->rmsThreshold); - validator = hipdnn_test_sdk::utilities::createRmsValidator( - sdkDataType, validatorOverride->rmsThreshold); - entryAtol = 0.0f; - entryRtol = 0.0f; + tolerance = bundle::ComparisonTolerance::rms(*rmsThreshold); } - auto [it, inserted] = _tensorValidationMap.insert( - {attr->get_uid(), - TensorValidationEntry{ - std::move(validator), attr->get_name(), entryAtol, entryRtol, sdkDataType}}); + auto selection = bundle::makeValidator(sdkDataType, label, tolerance); + auto [it, inserted] + = _tensorValidationMap.insert({attr->get_uid(), + TensorValidationEntry{std::move(selection.validator), + std::move(selection.error), + label, + tolerance, + sdkDataType}}); if(!inserted) { ADD_FAILURE() << "Duplicate validator for tensor " << attr->get_uid() << " (" - << attr->get_name() << "); keeping first registration"; + << label << "); keeping first registration"; } _tensorIdToNameMap.insert({attr->get_uid(), attr->get_name()}); }); @@ -420,37 +414,21 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParamsecond; - bool valid = entry.validator->allClose(*refTensor, *gpuTensor); - if(!valid) + if(entry.validator == nullptr) { - using hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType; - const std::string tensorLabel - = entry.name + " (UID " + std::to_string(tensorId) + ", output)"; - const std::string dtypeName - = entry.dataType != hipdnn_flatbuffers_sdk::data_objects::DataType::UNSET - ? EnumNameDataType(entry.dataType) - : "unknown"; - - hipdnn_test_sdk::utilities::ComparisonContext ctx; - ctx.contextLine = "Test: " + currentTestName(); - ctx.tensorLabel = tensorLabel; - ctx.dtypeName = dtypeName; - ctx.atol = entry.atol; - ctx.rtol = entry.rtol; - - std::ostringstream report; - report << hipdnn_test_sdk::utilities::formatComparisonHeader(ctx, *refTensor); - if(entry.dataType != hipdnn_flatbuffers_sdk::data_objects::DataType::UNSET) - { - hipdnn_test_sdk::utilities::appendComparisonDiffByDataType(report, - entry.dataType, - tensorLabel, - *refTensor, - *gpuTensor, - entry.atol, - entry.rtol); - } - EXPECT_TRUE(false) << report.str(); + ADD_FAILURE() << entry.validatorError; + continue; + } + + if(!entry.validator->allClose(*refTensor, *gpuTensor)) + { + ADD_FAILURE() << bundle::formatMismatchReport(tensorId, + entry.label, + entry.dataType, + *refTensor, + *gpuTensor, + entry.tolerance, + "Test: " + currentTestName()); } } } @@ -601,7 +579,7 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParamsecond.name; + return valIt->second.label; } auto nameIt = _tensorIdToNameMap.find(tensorId); if(nameIt != _tensorIdToNameMap.end()) diff --git a/dnn-providers/integration-tests/src/harness/TestSettings.hpp b/dnn-providers/integration-tests/src/harness/TestSettings.hpp index 0f2acdd7d2ac..388ebe1da650 100644 --- a/dnn-providers/integration-tests/src/harness/TestSettings.hpp +++ b/dnn-providers/integration-tests/src/harness/TestSettings.hpp @@ -93,6 +93,9 @@ struct ValidatorOverride // label (its name, or "uid=N" when the graph did not name it). Later entries // take precedence. Absent any match, allclose is used — allclose is the default // and nothing else can select a validator. See ALMIOPEN-2561. +// Validator entries are parsed strictly: 'rms' requires a positive 'rms_threshold' +// and 'allclose' must not carry one, because an entry that does not say exactly +// what it means is a load error rather than a silent fall-back to allclose. // For test_skips, an entry matches when ALL of: // - 'archs' is omitted/empty (any arch), OR any 'archs' value is a // substring of the device's raw gcnArchName. @@ -399,6 +402,16 @@ class TestSettings if(*validator == "allclose") { + // A threshold on an allclose entry means the file does not say what its + // author meant — either 'validator' was edited and the threshold left + // behind, or the reverse. Ignoring it silently is how an override quietly + // stops being the check someone thinks it is. + if((*table)["rms_threshold"]) + { + throw std::runtime_error( + std::string(K_SECTION) + + R"~( entry with validator = "allclose" also sets 'rms_threshold')~"); + } parsed.kind = ValidatorOverrideKind::ALLCLOSE; parsed.rmsThreshold = 0.0f; return parsed; diff --git a/dnn-providers/integration-tests/src/harness/TomlGuards.hpp b/dnn-providers/integration-tests/src/harness/TomlGuards.hpp index b75aabfff3aa..c8232c345bf3 100644 --- a/dnn-providers/integration-tests/src/harness/TomlGuards.hpp +++ b/dnn-providers/integration-tests/src/harness/TomlGuards.hpp @@ -43,6 +43,32 @@ inline bool applyTomlToleranceOverride(const std::string& testName, float& atol, return true; } +// The relative-RMS threshold this engine's TOML selects for one output tensor of one +// test, or nullopt when no [[validator_overrides]] entry names it — which means +// allclose, the default and the only thing any other code path can produce. +// +// Shared by both verification harnesses: the selection is a per-engine numerical +// property, and two copies of it would be two chances to disagree about what the +// config said. `tensorLabel` must be the label form the TOML globs are written +// against — bundle::tensorLabel(uid, name), never a raw tensor name. +inline std::optional findTomlRmsThreshold(const std::string& testName, + const std::string& tensorLabel) +{ + if(testName.empty()) + { + return std::nullopt; + } + const auto selected = TestConfig::get().findValidatorOverride(testName, tensorLabel); + if(!selected || selected->kind != ValidatorOverrideKind::RMS) + { + return std::nullopt; + } + HIPDNN_PLUGIN_LOG_INFO("Validator override applied for " + << testName << " tensor " << tensorLabel + << ": rms, threshold=" << selected->rmsThreshold); + return selected->rmsThreshold; +} + inline std::optional checkTomlSkip(const std::string& testName) { if(testName.empty()) diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp index 5ab94dd1fb1a..573185efbb72 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp @@ -583,14 +583,9 @@ VerificationOutcome hipdnn_flatbuffers_sdk::data_objects::DataType dataType) { // allclose unless this engine's TOML names the tensor. Nothing else selects a // validator — see ALMIOPEN-2561. - const auto validatorOverride - = TestConfig::get().findValidatorOverride(currentTestName(), label); - if(validatorOverride && validatorOverride->kind == ValidatorOverrideKind::RMS) + if(const auto rmsThreshold = findTomlRmsThreshold(currentTestName(), label)) { - HIPDNN_PLUGIN_LOG_INFO("Validator override applied for " - << currentTestName() << " tensor " << label - << ": rms, threshold=" << validatorOverride->rmsThreshold); - return ComparisonTolerance::rms(validatorOverride->rmsThreshold); + return ComparisonTolerance::rms(*rmsThreshold); } ComparisonTolerance tolerance; diff --git a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp index 2ab2259c454d..127ac37a1915 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp @@ -3,7 +3,9 @@ #include "harness/bundle/OutputComparison.hpp" +#include #include +#include #include #include @@ -12,43 +14,73 @@ namespace hipdnn_integration_tests::bundle { -std::string tensorLabel(int64_t uid, - const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes& attrs) +std::string tensorLabel(int64_t uid, const std::string& name) { - const auto* name = attrs.name(); - if(name != nullptr && !name->empty()) + if(!name.empty()) { - return name->str(); + return name; } return "uid=" + std::to_string(uid); } -std::optional - compareTensor(int64_t uid, - const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes& attrs, - hipdnn_data_sdk::utilities::ITensor& expected, - hipdnn_data_sdk::utilities::ITensor& actual, - ComparisonTolerance tolerance, - const std::string& contextLine) +std::string tensorLabel(int64_t uid, + const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes& attrs) { - const auto dataType = attrs.data_type(); - const bool useRms = tolerance.kind == ValidatorKind::RMS; + const auto* name = attrs.name(); + return tensorLabel(uid, name != nullptr ? name->str() : std::string{}); +} - auto validator - = useRms ? hipdnn_test_sdk::utilities::createRmsValidator(dataType, tolerance.rmsThreshold) - : hipdnn_test_sdk::utilities::createAllCloseValidator( - dataType, tolerance.atol, tolerance.rtol); - if(validator->allClose(expected, actual)) +ValidatorSelection makeValidator(hipdnn_flatbuffers_sdk::data_objects::DataType dataType, + const std::string& label, + const ComparisonTolerance& tolerance) +{ + if(tolerance.kind != ValidatorKind::RMS) { - return std::nullopt; + return {hipdnn_test_sdk::utilities::createAllCloseValidator( + dataType, tolerance.atol, tolerance.rtol), + {}}; } - const auto label = tensorLabel(uid, attrs); + try + { + return {hipdnn_test_sdk::utilities::createRmsValidator(dataType, tolerance.rmsThreshold), + {}}; + } + catch(const std::exception& e) + { + // An over-broad [[validator_overrides]] 'tensors' glob, not a numerical failure: + // say which tensor it caught and what to narrow, rather than unwinding the test. + std::ostringstream error; + error << "\nValidator override NOT APPLICABLE\n" + << " Tensor: " << label << "\n" + << " Data type: " << hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType) + << "\n" + << " A [[validator_overrides]] entry in this engine's TOML config selected the\n" + " rms validator for this tensor, but it does not support this data type (" + << e.what() + << ").\n" + " Narrow that entry's 'tensors' glob so it no longer matches this tensor.\n"; + return {nullptr, error.str()}; + } +} + +std::string formatMismatchReport(int64_t uid, + const std::string& label, + hipdnn_flatbuffers_sdk::data_objects::DataType dataType, + hipdnn_data_sdk::utilities::ITensor& expected, + hipdnn_data_sdk::utilities::ITensor& actual, + const ComparisonTolerance& tolerance, + const std::string& contextLine) +{ + using hipdnn_flatbuffers_sdk::data_objects::DataType; + const bool useRms = tolerance.kind == ValidatorKind::RMS; hipdnn_test_sdk::utilities::ComparisonContext ctx; ctx.contextLine = contextLine; ctx.tensorLabel = label + " (UID " + std::to_string(uid) + ", output)"; - ctx.dtypeName = hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType); + ctx.dtypeName = dataType != DataType::UNSET + ? hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType) + : "unknown"; ctx.atol = tolerance.atol; ctx.rtol = tolerance.rtol; if(useRms) @@ -63,17 +95,46 @@ std::optional std::ostringstream report; report << hipdnn_test_sdk::utilities::formatComparisonHeader(ctx, expected); - // Zero tolerances under RMS: the per-element budget is not what was checked, and a - // full drift profile (max/mean abs diff, worst elements) is the useful diagnostic. - hipdnn_test_sdk::utilities::appendComparisonDiffByDataType(report, - dataType, - label, - expected, - actual, - useRms ? 0.0f : tolerance.atol, - useRms ? 0.0f : tolerance.rtol); - - return TensorMismatch{uid, label, report.str()}; + if(dataType != DataType::UNSET) + { + // Zero tolerances under RMS: the per-element budget is not what was checked, and a + // full drift profile (max/mean abs diff, worst elements) is the useful diagnostic. + hipdnn_test_sdk::utilities::appendComparisonDiffByDataType(report, + dataType, + label, + expected, + actual, + useRms ? 0.0f : tolerance.atol, + useRms ? 0.0f : tolerance.rtol); + } + return report.str(); +} + +std::optional + compareTensor(int64_t uid, + const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes& attrs, + hipdnn_data_sdk::utilities::ITensor& expected, + hipdnn_data_sdk::utilities::ITensor& actual, + ComparisonTolerance tolerance, + const std::string& contextLine) +{ + const auto dataType = attrs.data_type(); + const auto label = tensorLabel(uid, attrs); + + auto selection = makeValidator(dataType, label, tolerance); + if(selection.validator == nullptr) + { + return TensorMismatch{uid, label, std::move(selection.error)}; + } + if(selection.validator->allClose(expected, actual)) + { + return std::nullopt; + } + + return TensorMismatch{ + uid, + label, + formatMismatchReport(uid, label, dataType, expected, actual, tolerance, contextLine)}; } std::vector diff --git a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp index cf7ae2a01036..f57fe1fc7bbe 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp @@ -14,6 +14,7 @@ #include #include #include +#include namespace hipdnn_integration_tests::bundle { @@ -110,7 +111,43 @@ std::vector const std::string& contextLine); /// The tensor's name, or "uid=N" when the graph did not give it one. +/// +/// Every harness resolves a tensor's label through this. The label is both what a +/// TOML `tensors` glob is matched against and what the failure report prints, so the +/// two must not be allowed to disagree about what a tensor is called. +std::string tensorLabel(int64_t uid, const std::string& name); + std::string tensorLabel(int64_t uid, const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes& attrs); +/// The validator one tolerance selects, or the reason it could not be built. +struct ValidatorSelection +{ + std::unique_ptr validator; + std::string error; ///< non-empty exactly when `validator` is null +}; + +/// Builds the validator `tolerance` selects for one output tensor. +/// +/// RMS is implemented for FLOAT/HALF/BFLOAT16/DOUBLE only, so a `tensors` glob one +/// wildcard too wide can select it for an integer output. The TOML parser cannot catch +/// that — a tensor's data type is not known until its graph is read — so it surfaces +/// here as a named failure that says which glob over-matched, rather than as an +/// exception unwinding out of the test body. +ValidatorSelection makeValidator(hipdnn_flatbuffers_sdk::data_objects::DataType dataType, + const std::string& label, + const ComparisonTolerance& tolerance); + +/// The failure report for one tensor: header plus per-element drift profile. +/// +/// Shared by both harnesses so a tensor graded by RMS reports the threshold that +/// decided it on either path, instead of an atol/rtol pair that did not. +std::string formatMismatchReport(int64_t uid, + const std::string& label, + hipdnn_flatbuffers_sdk::data_objects::DataType dataType, + hipdnn_data_sdk::utilities::ITensor& expected, + hipdnn_data_sdk::utilities::ITensor& actual, + const ComparisonTolerance& tolerance, + const std::string& contextLine); + } // namespace hipdnn_integration_tests::bundle diff --git a/dnn-providers/integration-tests/tests/TestOutputComparison.cpp b/dnn-providers/integration-tests/tests/TestOutputComparison.cpp index ce8ff7e61da4..154d2d1fd64d 100644 --- a/dnn-providers/integration-tests/tests/TestOutputComparison.cpp +++ b/dnn-providers/integration-tests/tests/TestOutputComparison.cpp @@ -9,6 +9,7 @@ #include #include +#include #include #include @@ -30,6 +31,9 @@ namespace constexpr int64_t K_UID_A = 5; // named "y_out" constexpr int64_t K_UID_B = 4; // unnamed +// Not wired to a node: it exists so a test can ask what an integer output does when a +// 'tensors' glob is wide enough to select RMS for it. +constexpr int64_t K_UID_INT = 6; // named "counts", INT32 // The batchnorm graph the other bundle tests use, with uid 5 given a name so both // label paths are reachable: uid 5 reports its name, uid 4 falls back to its uid. @@ -49,7 +53,9 @@ const std::string K_GRAPH R"({"name": "", "uid": 4, "strides": [3, 1, 1, 1], "dims": [1, 3, 1, 1], )" R"("data_type": "float", "virtual": false}, )" R"({"name": "y_out", "uid": 5, "strides": [60, 20, 5, 1], "dims": [2, 3, 4, 5], )" - R"("data_type": "float", "virtual": false}], "io_data_type": "float", )" + R"("data_type": "float", "virtual": false}, )" + R"({"name": "counts", "uid": 6, "strides": [3, 1, 1, 1], "dims": [1, 3, 1, 1], )" + R"("data_type": "int32", "virtual": false}], "io_data_type": "float", )" R"("compute_data_type": "float", "intermediate_data_type": "float", "name": ""})"; // Built the same way the bundle loader builds one, so these tests walk a real // flatbuffer graph rather than a hand-rolled attribute map. @@ -86,6 +92,16 @@ std::unique_ptr return tensor; } +// An integer tensor, filled with one value. Integer outputs exist in this suite, and +// the RMS validator has no implementation for them. +std::unique_ptr + intTensor(const hipdnn_flatbuffers_sdk::data_objects::TensorAttributes& attrs, int32_t value) +{ + auto tensor = hipdnn_test_sdk::detail::createTensorFromAttribute(attrs); + tensor->fillTensorWithValue(static_cast(value)); + return tensor; +} + ComparisonTolerance exact() { return ComparisonTolerance::allClose(0.0f, 0.0f); @@ -367,4 +383,72 @@ TEST(TestOutputComparison, ValidatorKindIsChosenPerTensor) EXPECT_EQ(mismatches[0].uid, K_UID_A); } +// --------------------------------------------------------------------------- +// Both harnesses resolve a tensor's label through the same helper: it is what a TOML +// 'tensors' glob matches on and what the report prints, and a raw tensor name is +// neither when the graph left the tensor unnamed. +// --------------------------------------------------------------------------- + +TEST(TestOutputComparison, LabelFromANameKeepsTheName) +{ + EXPECT_EQ(tensorLabel(7, "LayernormBackward_0::DSCALE"), "LayernormBackward_0::DSCALE"); +} + +TEST(TestOutputComparison, LabelFromAnEmptyNameFallsBackToTheUid) +{ + EXPECT_EQ(tensorLabel(7, ""), "uid=7"); +} + +// --------------------------------------------------------------------------- +// RMS is defined for float, half, bfloat16 and double only. A 'tensors' glob one +// wildcard too wide can still select it for an integer output, and the TOML parser +// cannot catch that: a tensor's dtype is not known until its graph is read. It has to +// come back as a comparison failure naming the tensor, not an exception out of the +// test body. +// --------------------------------------------------------------------------- + +TEST(TestOutputComparison, RmsOnAnUnsupportedDataTypeIsReportedNotThrown) +{ + const auto buffer = makeGraphBuffer(); + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper wrapper{buffer.data(), + buffer.size()}; + const auto& attrs = *wrapper.getTensorMap().at(K_UID_INT); + + auto expected = intTensor(attrs, 7); + auto actual = intTensor(attrs, 7); + + std::optional mismatch; + ASSERT_NO_THROW(mismatch = compareTensor(K_UID_INT, + attrs, + *expected, + *actual, + ComparisonTolerance::rms(K_RMS_THRESHOLD), + "Bundle: b")); + + // Equal tensors, so this is not a numerical verdict: it reports that the override + // could not be honoured at all. + ASSERT_TRUE(mismatch.has_value()); + EXPECT_EQ(mismatch->label, "counts"); + EXPECT_NE(mismatch->report.find("counts"), std::string::npos); + EXPECT_NE(mismatch->report.find("INT32"), std::string::npos); + EXPECT_NE(mismatch->report.find("validator_overrides"), std::string::npos) + << "the operator has to be told which config section over-matched"; +} + +// The guard is scoped to RMS: integer outputs still compare normally under the default. +TEST(TestOutputComparison, AllcloseStillGradesIntegerOutputs) +{ + const auto buffer = makeGraphBuffer(); + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper wrapper{buffer.data(), + buffer.size()}; + const auto& attrs = *wrapper.getTensorMap().at(K_UID_INT); + + auto expected = intTensor(attrs, 7); + auto matching = intTensor(attrs, 7); + auto drifted = intTensor(attrs, 9); + + EXPECT_FALSE(compareTensor(K_UID_INT, attrs, *expected, *matching, exact(), "b").has_value()); + EXPECT_TRUE(compareTensor(K_UID_INT, attrs, *expected, *drifted, exact(), "b").has_value()); +} + // NOLINTEND(readability-identifier-naming) diff --git a/dnn-providers/integration-tests/tests/TestTestSettings.cpp b/dnn-providers/integration-tests/tests/TestTestSettings.cpp index 34b0599a8608..fd5bc403fecf 100644 --- a/dnn-providers/integration-tests/tests/TestTestSettings.cpp +++ b/dnn-providers/integration-tests/tests/TestTestSettings.cpp @@ -465,6 +465,25 @@ rms_threshold = 0.0 EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); } +// An allclose entry carrying a threshold means the file no longer says what its +// author meant — 'validator' was edited and 'rms_threshold' left behind, or the +// reverse. Either way, the file is wrong about which check runs. +TEST(TestSettingsValidatorOverrides, ThrowsOnAllcloseWithRmsThreshold) +{ + const TempTomlFile file(R"( +[meta] +version = 1 + +[[validator_overrides]] +filters = ["*LayernormBackward*"] +tensors = ["*::DSCALE"] +validator = "allclose" +rms_threshold = 1e-4 +)"); + + EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); +} + // --------------------------------------------------------------------------- // [[test_skips]] parsing // --------------------------------------------------------------------------- diff --git a/dnn-providers/integration-tests/tests/TestTomlGuards.cpp b/dnn-providers/integration-tests/tests/TestTomlGuards.cpp index aa7934ac532d..0c92e37ff2c7 100644 --- a/dnn-providers/integration-tests/tests/TestTomlGuards.cpp +++ b/dnn-providers/integration-tests/tests/TestTomlGuards.cpp @@ -8,6 +8,7 @@ using hipdnn_integration_tests::applyTomlToleranceOverride; using hipdnn_integration_tests::checkTomlSkip; using hipdnn_integration_tests::currentTestName; +using hipdnn_integration_tests::findTomlRmsThreshold; // NOLINTBEGIN(readability-identifier-naming) -- gtest macro-generated names @@ -45,6 +46,11 @@ TEST(TestTomlGuards, ApplyTomlToleranceOverrideReturnsFalseForEmptyName) EXPECT_FLOAT_EQ(rtol, 1.0f); } +TEST(TestTomlGuards, FindTomlRmsThresholdReturnsNulloptForEmptyName) +{ + EXPECT_EQ(findTomlRmsThreshold("", "LayernormBackward_0::DSCALE"), std::nullopt); +} + // --------------------------------------------------------------------------- // checkTomlSkip / applyTomlToleranceOverride — no TOML loaded // @@ -67,4 +73,11 @@ TEST(TestTomlGuards, ApplyTomlToleranceOverrideReturnsFalseWhenNoSettings) EXPECT_FLOAT_EQ(rtol, 1.0f); } +// The default is allclose, and it is the absence of a matching [[validator_overrides]] +// entry that expresses it — no config, no RMS threshold, on either harness. +TEST(TestTomlGuards, FindTomlRmsThresholdReturnsNulloptWhenNoSettings) +{ + EXPECT_EQ(findTomlRmsThreshold("SomeTest.Name", "LayernormBackward_0::DSCALE"), std::nullopt); +} + // NOLINTEND(readability-identifier-naming) From 665e76eebf0918c572218a2c332ae6db04f21782 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 10 Sep 2026 14:53:07 -0600 Subject: [PATCH 4/6] refactor(integration-tests): switch on validator kind and give the report the threshold MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 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 --- .../src/harness/bundle/OutputComparison.cpp | 64 +++++++++++-------- .../tests/TestOutputComparison.cpp | 17 +++++ .../utilities/ComparisonReport.hpp | 19 ++++-- .../tests/utilities/TestComparisonReport.cpp | 36 +++++++++-- 4 files changed, 95 insertions(+), 41 deletions(-) diff --git a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp index 127ac37a1915..9e8f1b8f29f5 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp @@ -5,7 +5,7 @@ #include #include -#include +#include #include #include @@ -34,33 +34,44 @@ ValidatorSelection makeValidator(hipdnn_flatbuffers_sdk::data_objects::DataType const std::string& label, const ComparisonTolerance& tolerance) { - if(tolerance.kind != ValidatorKind::RMS) + switch(tolerance.kind) { + case ValidatorKind::ALLCLOSE: return {hipdnn_test_sdk::utilities::createAllCloseValidator( dataType, tolerance.atol, tolerance.rtol), {}}; - } - try - { - return {hipdnn_test_sdk::utilities::createRmsValidator(dataType, tolerance.rmsThreshold), + case ValidatorKind::RMS: + // Only RMS is caught. It is the one kind a [[validator_overrides]] glob can + // select, so an unsupported data type here is an operator's config mistake and + // deserves a legible answer. allclose is the default that nothing selects, so + // there is no glob to blame and its own throw stays a throw. + try + { + return { + hipdnn_test_sdk::utilities::createRmsValidator(dataType, tolerance.rmsThreshold), {}}; - } - catch(const std::exception& e) - { - // An over-broad [[validator_overrides]] 'tensors' glob, not a numerical failure: - // say which tensor it caught and what to narrow, rather than unwinding the test. - std::ostringstream error; - error << "\nValidator override NOT APPLICABLE\n" - << " Tensor: " << label << "\n" - << " Data type: " << hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType) - << "\n" - << " A [[validator_overrides]] entry in this engine's TOML config selected the\n" - " rms validator for this tensor, but it does not support this data type (" - << e.what() - << ").\n" - " Narrow that entry's 'tensors' glob so it no longer matches this tensor.\n"; - return {nullptr, error.str()}; + } + catch(const std::exception& e) + { + std::ostringstream error; + error << "\nValidator override NOT APPLICABLE\n" + << " Tensor: " << label << "\n" + << " Data type: " + << hipdnn_flatbuffers_sdk::data_objects::EnumNameDataType(dataType) << "\n" + << " A [[validator_overrides]] entry in this engine's TOML config selected the\n" + " rms validator for this tensor, but it does not support this data type (" + << e.what() + << ").\n" + " Narrow that entry's 'tensors' glob so it no longer matches this tensor.\n"; + return {nullptr, error.str()}; + } + + default: + // A kind that is neither, i.e. a new ValidatorKind whose case was never written. + // Grading it as allclose by omission is exactly the silent miscompare this whole + // mechanism exists to prevent, so refuse instead. + throw std::invalid_argument("makeValidator: unhandled ValidatorKind"); } } @@ -85,12 +96,9 @@ std::string formatMismatchReport(int64_t uid, ctx.rtol = tolerance.rtol; if(useRms) { - // atol/rtol did not decide this failure, so do not print them as if they had. - std::ostringstream summary; - summary << "relative RMS <= " << tolerance.rmsThreshold - << " (aggregate check — the element counts below are elements that " - "differ at all, not elements that failed)"; - ctx.toleranceSummary = summary.str(); + // Carry the threshold, not a rendered sentence: the report module owns how a + // tolerance is worded, and atol/rtol did not decide this failure. + ctx.rmsThreshold = tolerance.rmsThreshold; } std::ostringstream report; diff --git a/dnn-providers/integration-tests/tests/TestOutputComparison.cpp b/dnn-providers/integration-tests/tests/TestOutputComparison.cpp index 154d2d1fd64d..d5bfb0556555 100644 --- a/dnn-providers/integration-tests/tests/TestOutputComparison.cpp +++ b/dnn-providers/integration-tests/tests/TestOutputComparison.cpp @@ -10,6 +10,7 @@ #include #include +#include #include #include @@ -21,8 +22,10 @@ using hipdnn_integration_tests::bundle::compareOutputs; using hipdnn_integration_tests::bundle::compareTensor; using hipdnn_integration_tests::bundle::ComparisonTolerance; +using hipdnn_integration_tests::bundle::makeValidator; using hipdnn_integration_tests::bundle::OutputTensors; using hipdnn_integration_tests::bundle::tensorLabel; +using hipdnn_integration_tests::bundle::ValidatorKind; // NOLINTBEGIN(readability-identifier-naming) @@ -451,4 +454,18 @@ TEST(TestOutputComparison, AllcloseStillGradesIntegerOutputs) EXPECT_TRUE(compareTensor(K_UID_INT, attrs, *expected, *drifted, exact(), "b").has_value()); } +// A ValidatorKind with no case in makeValidator must not be graded by whichever branch +// happens to fall through. Silently comparing a tensor with a validator nobody chose is +// the failure mode this whole mechanism exists to prevent, so the unhandled kind is +// refused outright rather than defaulted. +TEST(TestOutputComparison, UnhandledValidatorKindIsRefused) +{ + ComparisonTolerance bogus = ComparisonTolerance::allClose(0.0f, 0.0f); + bogus.kind = static_cast(99); + + EXPECT_THROW( + makeValidator(hipdnn_flatbuffers_sdk::data_objects::DataType::FLOAT, "y_out", bogus), + std::invalid_argument); +} + // NOLINTEND(readability-identifier-naming) diff --git a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/ComparisonReport.hpp b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/ComparisonReport.hpp index f79ce42a9368..1e8a0bf2e612 100644 --- a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/ComparisonReport.hpp +++ b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/ComparisonReport.hpp @@ -3,6 +3,7 @@ #pragma once +#include #include #include #include @@ -22,9 +23,10 @@ struct ComparisonContext std::string dtypeName; float atol = 0.0f; float rtol = 0.0f; - /// Replaces the "atol=… rtol=…" line when set, for checks that atol/rtol do not - /// describe (e.g. an aggregate relative-RMS comparison). - std::string toleranceSummary; + /// Set when the verdict came from an aggregate relative-RMS check rather than a + /// per-element one. atol/rtol did not decide it and are not printed; this threshold + /// is. Unset means an ordinary allclose comparison. + std::optional rmsThreshold; }; inline std::string formatComparisonHeader(const ComparisonContext& ctx, @@ -35,13 +37,18 @@ inline std::string formatComparisonHeader(const ComparisonContext& ctx, << " " << ctx.contextLine << "\n" << " Tensor: " << ctx.tensorLabel << "\n" << " Shape: " << StreamVec(tensor.dims()) << " " << ctx.dtypeName << "\n"; - if(ctx.toleranceSummary.empty()) + if(ctx.rmsThreshold.has_value()) { - os << " Tolerance: atol=" << ctx.atol << " rtol=" << ctx.rtol << "\n"; + // The element listing below counts every element that differs at all, because an + // aggregate check has no per-element budget to fail against. Say so, or the + // counts read as failures. + os << " Tolerance: relative RMS <= " << *ctx.rmsThreshold + << " (aggregate check — the element counts below are elements that differ at " + "all, not elements that failed)\n"; } else { - os << " Tolerance: " << ctx.toleranceSummary << "\n"; + os << " Tolerance: atol=" << ctx.atol << " rtol=" << ctx.rtol << "\n"; } return os.str(); } diff --git a/projects/hipdnn/test_sdk/tests/utilities/TestComparisonReport.cpp b/projects/hipdnn/test_sdk/tests/utilities/TestComparisonReport.cpp index 1bc0539cc8a9..d0ce8ea26a91 100644 --- a/projects/hipdnn/test_sdk/tests/utilities/TestComparisonReport.cpp +++ b/projects/hipdnn/test_sdk/tests/utilities/TestComparisonReport.cpp @@ -5,6 +5,7 @@ #include #include +#include #include #include @@ -18,8 +19,12 @@ using DT = hipdnn_flatbuffers_sdk::data_objects::DataType; TEST(TestFormatComparisonHeader, ContainsAllFields) { - const ComparisonContext ctx{ - "Bundle: /path/to/bundle", "output_tensor (UID 42, output)", "FLOAT", 1e-5f, 1e-4f, ""}; + const ComparisonContext ctx{"Bundle: /path/to/bundle", + "output_tensor (UID 42, output)", + "FLOAT", + 1e-5f, + 1e-4f, + std::nullopt}; const Tensor tensor({2, 3, 4}); const std::string header = formatComparisonHeader(ctx, tensor); @@ -34,7 +39,7 @@ TEST(TestFormatComparisonHeader, ContainsAllFields) TEST(TestFormatComparisonHeader, IncludesShape) { - const ComparisonContext ctx{"Test: MyTest.Case", "x", "HALF", 0.0f, 0.0f, ""}; + const ComparisonContext ctx{"Test: MyTest.Case", "x", "HALF", 0.0f, 0.0f, std::nullopt}; const Tensor tensor({8, 16}); const std::string header = formatComparisonHeader(ctx, tensor); @@ -44,11 +49,11 @@ TEST(TestFormatComparisonHeader, IncludesShape) } // A check that atol/rtol do not describe says so, instead of printing two numbers that -// had nothing to do with the verdict. -TEST(TestFormatComparisonHeader, ToleranceSummaryReplacesAtolRtol) +// had nothing to do with the verdict. The header owns the wording: callers hand over the +// threshold, never a rendered sentence. +TEST(TestFormatComparisonHeader, RmsThresholdReplacesAtolRtol) { - const ComparisonContext ctx{ - "Test: MyTest.Case", "dscale", "FLOAT", 1e-5f, 1e-4f, "relative RMS <= 0.0001"}; + const ComparisonContext ctx{"Test: MyTest.Case", "dscale", "FLOAT", 1e-5f, 1e-4f, 1e-4f}; const Tensor tensor({8, 16}); const std::string header = formatComparisonHeader(ctx, tensor); @@ -56,6 +61,23 @@ TEST(TestFormatComparisonHeader, ToleranceSummaryReplacesAtolRtol) EXPECT_NE(header.find("relative RMS <= 0.0001"), std::string::npos); EXPECT_EQ(header.find("atol="), std::string::npos); EXPECT_EQ(header.find("rtol="), std::string::npos); + // The element counts printed below an aggregate verdict are elements that differ at + // all, not elements that failed. Saying so is the header's job, not the caller's. + EXPECT_NE(header.find("aggregate check"), std::string::npos); + EXPECT_NE(header.find("not elements that failed"), std::string::npos); +} + +// The threshold is optional, and absent means an ordinary per-element verdict. +TEST(TestFormatComparisonHeader, NoRmsThresholdPrintsAtolRtol) +{ + const ComparisonContext ctx{"Test: MyTest.Case", "dx", "FLOAT", 1e-5f, 1e-4f, std::nullopt}; + + const Tensor tensor({8, 16}); + const std::string header = formatComparisonHeader(ctx, tensor); + + EXPECT_NE(header.find("atol=1e-05"), std::string::npos); + EXPECT_NE(header.find("rtol=0.0001"), std::string::npos); + EXPECT_EQ(header.find("relative RMS"), std::string::npos); } // ================================================================================================= From 5f7c208055b6756d4c5fc29f58b69493887930d4 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Fri, 11 Sep 2026 15:45:06 -0600 Subject: [PATCH 5/6] fix(hip-kernel-provider): restore dtype-specific layernorm RMS thresholds 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 --- .../config/HIP_MLOPS_ENGINE.toml | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml b/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml index 947f2c4f4657..bda0263e694d 100644 --- a/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml +++ b/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml @@ -47,6 +47,23 @@ filters = ["*LayernormBackward*"] tensors = ["*::DSCALE", "*::DBIAS"] validator = "rms" rms_threshold = 1e-4 + +# Match layernorm::getTolerance() in TestTolerances.hpp for the dtype of +# dscale/dbias. Only Pure fixtures store these outputs in FP16/BF16; Mixed and +# Upcast fixtures keep them in FP32 and retain the 1e-4 threshold above. +# These more-specific entries must follow the shared rule: later matches win. +[[validator_overrides]] +filters = ["*LayernormBackwardPure*Fp16.*"] +tensors = ["*::DSCALE", "*::DBIAS"] +validator = "rms" +rms_threshold = 1e-3 + +[[validator_overrides]] +filters = ["*LayernormBackwardPure*Bfp16.*"] +tensors = ["*::DSCALE", "*::DBIAS"] +validator = "rms" +rms_threshold = 1e-2 + # # RMSNormBackward Pure-Bfp16 (dy/x/dx all bf16) fails CPU-reference tolerance # at the suite's large-reduction shapes (x:[4096,128,1,1], x:[5,256,14,14], From f2b969b794a4889670d37f1f95e10c14d4132767 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Tue, 15 Sep 2026 17:21:05 -0600 Subject: [PATCH 6/6] refactor(integration-tests): log the check that graded each output tensor MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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. --- .../config/HIP_MLOPS_ENGINE.toml | 2 +- dnn-providers/integration-tests/README.md | 2 +- .../IntegrationGraphVerificationHarness.hpp | 58 +++++++++---------- .../src/harness/TestSettings.hpp | 2 +- .../src/harness/TomlGuards.hpp | 42 ++++++++++++-- .../BundleReferenceValidationHarness.cpp | 5 +- .../IntegrationBundleVerificationHarness.cpp | 33 +++-------- .../src/harness/bundle/OutputComparison.cpp | 13 ++--- .../src/harness/bundle/OutputComparison.hpp | 15 ++--- .../IntegrationGpuLayernormBackward.cpp | 7 ++- .../tests/TestOutputComparison.cpp | 10 ++-- .../tests/TestTestSettings.cpp | 52 ++++++++++++++++- .../tests/TestTomlGuards.cpp | 14 +++++ 13 files changed, 165 insertions(+), 90 deletions(-) diff --git a/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml b/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml index bda0263e694d..c46da5c291e1 100644 --- a/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml +++ b/dnn-providers/hip-kernel-provider/config/HIP_MLOPS_ENGINE.toml @@ -41,7 +41,7 @@ version = 1 # Measured on gfx1151: 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 and # swings ~2x with the input seed, and applies to dx as well, whose real error is 4-5 -# decades smaller. Relative RMS stays 1e-6..8e-6 across the whole set. ALMIOPEN-2561. +# decades smaller. Relative RMS stays 1e-6..8e-6 across the whole set. [[validator_overrides]] filters = ["*LayernormBackward*"] tensors = ["*::DSCALE", "*::DBIAS"] diff --git a/dnn-providers/integration-tests/README.md b/dnn-providers/integration-tests/README.md index f348d86b9314..4049684f3cc1 100644 --- a/dnn-providers/integration-tests/README.md +++ b/dnn-providers/integration-tests/README.md @@ -447,7 +447,7 @@ reason = "ROCm/rocm-libraries#6979 — no engine has an applicable solution for a long reduction (layernorm/RMSNorm backward `dscale`/`dbias`) has elements that land arbitrarily near zero through cancellation, so per-element relative error is unbounded while the aggregate relative-RMS error is not. Prefer - `tolerance_overrides` for everything else. See ALMIOPEN-2561. + `tolerance_overrides` for everything else. - `test_skips`: the first matching entry wins; `reason` is surfaced in the `GTEST_SKIP` message. `archs` (substring match against the raw `gcnArchName`) and `platforms` (`"windows"`/`"linux"`) are both optional — diff --git a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp index e33779f164b6..80bbb7059de8 100644 --- a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp +++ b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp @@ -269,42 +269,38 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam attr, float absoluteTolerance, float relativeTolerance) { - float finalAtol = absoluteTolerance; - float finalRtol = relativeTolerance; - applyTomlToleranceOverride(currentTestName(), finalAtol, finalRtol); - const auto testName = currentTestName(); - _deferredValidators.emplace_back([this, attr, testName, finalAtol, finalRtol]() { - const auto sdkDataType - = hipdnn_test_sdk::utilities::frontendToSdkDataType(attr->get_data_type()); - const auto label = bundle::tensorLabel(attr->get_uid(), attr->get_name()); - - auto tolerance = bundle::ComparisonTolerance::allClose(finalAtol, finalRtol); - if(const auto rmsThreshold = findTomlRmsThreshold(testName, label)) - { - tolerance = bundle::ComparisonTolerance::rms(*rmsThreshold); - } - - auto selection = bundle::makeValidator(sdkDataType, label, tolerance); - auto [it, inserted] - = _tensorValidationMap.insert({attr->get_uid(), - TensorValidationEntry{std::move(selection.validator), - std::move(selection.error), - label, - tolerance, - sdkDataType}}); - if(!inserted) - { - ADD_FAILURE() << "Duplicate validator for tensor " << attr->get_uid() << " (" - << label << "); keeping first registration"; - } - _tensorIdToNameMap.insert({attr->get_uid(), attr->get_name()}); - }); + _deferredValidators.emplace_back( + [this, attr, testName, absoluteTolerance, relativeTolerance]() { + const auto sdkDataType + = hipdnn_test_sdk::utilities::frontendToSdkDataType(attr->get_data_type()); + const auto label = bundle::tensorLabel(attr->get_uid(), attr->get_name()); + + // One shared decision site for both harnesses: it picks the check and + // logs the one it picked. + const auto tolerance + = gradingForTensor(testName, label, absoluteTolerance, relativeTolerance); + + auto selection = bundle::makeValidator(sdkDataType, label, tolerance); + auto [it, inserted] = _tensorValidationMap.insert( + {attr->get_uid(), + TensorValidationEntry{std::move(selection.validator), + std::move(selection.error), + label, + tolerance, + sdkDataType}}); + if(!inserted) + { + ADD_FAILURE() << "Duplicate validator for tensor " << attr->get_uid() << " (" + << label << "); keeping first registration"; + } + _tensorIdToNameMap.insert({attr->get_uid(), attr->get_name()}); + }); } virtual void generateBundles(hipdnn_frontend::graph::Graph& graph, diff --git a/dnn-providers/integration-tests/src/harness/TestSettings.hpp b/dnn-providers/integration-tests/src/harness/TestSettings.hpp index 388ebe1da650..a1ab6d4bfd12 100644 --- a/dnn-providers/integration-tests/src/harness/TestSettings.hpp +++ b/dnn-providers/integration-tests/src/harness/TestSettings.hpp @@ -92,7 +92,7 @@ struct ValidatorOverride // gtest-formatted test name AND a 'tensors' glob matches the output tensor's // label (its name, or "uid=N" when the graph did not name it). Later entries // take precedence. Absent any match, allclose is used — allclose is the default -// and nothing else can select a validator. See ALMIOPEN-2561. +// and nothing else can select a validator. // Validator entries are parsed strictly: 'rms' requires a positive 'rms_threshold' // and 'allclose' must not carry one, because an entry that does not say exactly // what it means is a load error rather than a silent fall-back to allclose. diff --git a/dnn-providers/integration-tests/src/harness/TomlGuards.hpp b/dnn-providers/integration-tests/src/harness/TomlGuards.hpp index c8232c345bf3..f1c90900d600 100644 --- a/dnn-providers/integration-tests/src/harness/TomlGuards.hpp +++ b/dnn-providers/integration-tests/src/harness/TomlGuards.hpp @@ -11,6 +11,7 @@ #include #include "harness/TestConfig.hpp" +#include "harness/bundle/OutputComparison.hpp" namespace hipdnn_integration_tests { @@ -25,6 +26,11 @@ inline std::string currentTestName() return std::string(info->test_suite_name()) + "." + info->name(); } +// Applies this engine's [[tolerance_overrides]] entry for `testName`, if one matches. +// +// Silent by design. A tolerance is not necessarily the check that grades a tensor — a +// [[validator_overrides]] entry outranks it — and this function cannot know: it is given +// a test, not a tensor. Logging belongs to gradingForTensor(), which knows both. inline bool applyTomlToleranceOverride(const std::string& testName, float& atol, float& rtol) { if(testName.empty()) @@ -38,8 +44,6 @@ inline bool applyTomlToleranceOverride(const std::string& testName, float& atol, } atol = ovr->atol; rtol = ovr->rtol; - HIPDNN_PLUGIN_LOG_INFO("Tolerance override applied for " << testName << ": atol=" << atol - << " rtol=" << rtol); return true; } @@ -63,12 +67,40 @@ inline std::optional findTomlRmsThreshold(const std::string& testName, { return std::nullopt; } - HIPDNN_PLUGIN_LOG_INFO("Validator override applied for " - << testName << " tensor " << tensorLabel - << ": rms, threshold=" << selected->rmsThreshold); return selected->rmsThreshold; } +// How one output tensor is graded, and the one place that says so out loud. +// +// The validator override is read first because it outranks atol/rtol: announcing a +// tolerance before knowing whether it survives is how a reader is told the wrong check +// ran. `atol`/`rtol` are the fallback the caller resolved; they are logged only when a +// TOML entry changed them, so a quiet run means the defaults graded the tensor. +// +// Both verification harnesses resolve every output tensor through this, so neither can +// report a check the other would not have run. +inline bundle::ComparisonTolerance gradingForTensor(const std::string& testName, + const std::string& tensorLabel, + float atol, + float rtol) +{ + if(const auto rmsThreshold = findTomlRmsThreshold(testName, tensorLabel)) + { + HIPDNN_PLUGIN_LOG_INFO("Validator override applied for " + << testName << " tensor " << tensorLabel + << ": rms, threshold=" << *rmsThreshold); + return bundle::ComparisonTolerance::rms(*rmsThreshold); + } + + if(applyTomlToleranceOverride(testName, atol, rtol)) + { + HIPDNN_PLUGIN_LOG_INFO("Tolerance override applied for " + << testName << " tensor " << tensorLabel + << ": allclose, atol=" << atol << " rtol=" << rtol); + } + return bundle::ComparisonTolerance::allClose(atol, rtol); +} + inline std::optional checkTomlSkip(const std::string& testName) { if(testName.empty()) diff --git a/dnn-providers/integration-tests/src/harness/bundle/BundleReferenceValidationHarness.cpp b/dnn-providers/integration-tests/src/harness/bundle/BundleReferenceValidationHarness.cpp index cfc829162cfa..2ee52c1e2850 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/BundleReferenceValidationHarness.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/BundleReferenceValidationHarness.cpp @@ -128,9 +128,8 @@ void BundleReferenceValidationHarness::TestBody() // defaultTolerance(), never resolveTolerance(): a TOML override belongs to an // engine and must not loosen the gate on our own data. For the same reason this - // path never selects a non-default validator — allclose always (ALMIOPEN-2561). - const auto toleranceFor = [&wrapper](int64_t /*uid*/, - const std::string& /*label*/, + // path never selects a non-default validator — allclose always. + const auto toleranceFor = [&wrapper](const std::string& /*label*/, hipdnn_flatbuffers_sdk::data_objects::DataType dataType) { const float value = tolerance::defaultTolerance(wrapper, dataType); return ComparisonTolerance::allClose(value, value); diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp index 7742fb6c659d..512032498da2 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp @@ -19,7 +19,6 @@ #include #include "harness/ReferenceCapabilityError.hpp" -#include "harness/TestConfig.hpp" #include "harness/TomlGuards.hpp" #include "harness/bundle/LoadedEngine.hpp" #include "harness/bundle/LoadedEngineTable.hpp" @@ -629,30 +628,14 @@ VerificationOutcome const ExpectedTensorLookup& expectedFor) { auto wrapper = _bundle->graphWrapper(); - - const auto tomlOverride = TestConfig::get().findToleranceOverride(currentTestName()); - if(tomlOverride) - { - HIPDNN_PLUGIN_LOG_INFO("Tolerance override applied for " << currentTestName() - << ": atol=" << tomlOverride->atol - << " rtol=" << tomlOverride->rtol); - } - - const auto toleranceFor = [&](int64_t /*uid*/, - const std::string& label, - hipdnn_flatbuffers_sdk::data_objects::DataType dataType) { - // allclose unless this engine's TOML names the tensor. Nothing else selects a - // validator — see ALMIOPEN-2561. - if(const auto rmsThreshold = findTomlRmsThreshold(currentTestName(), label)) - { - return ComparisonTolerance::rms(*rmsThreshold); - } - - ComparisonTolerance tolerance; - tolerance::resolveTolerance( - wrapper, dataType, currentTestName(), tolerance.atol, tolerance.rtol); - return tolerance; - }; + // defaultTolerance() rather than resolveTolerance(): the TOML tolerance override is + // applied by gradingForTensor(), which reads the validator override first and so is + // the only place that can log the check that actually graded this tensor. + const auto toleranceFor + = [&](const std::string& label, hipdnn_flatbuffers_sdk::data_objects::DataType dataType) { + const float value = tolerance::defaultTolerance(wrapper, dataType); + return gradingForTensor(currentTestName(), label, value, value); + }; const auto mismatches = bundle::compareOutputs(wrapper, _bundle->outputTensorUids, diff --git a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp index 9e8f1b8f29f5..0422f123defb 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.cpp @@ -159,13 +159,12 @@ std::vector for(const int64_t uid : outputUids) { const auto* attrs = tensorAttrMap.at(uid); - auto mismatch - = compareTensor(uid, - *attrs, - expectedFor(uid), - *actual.at(uid), - toleranceFor(uid, tensorLabel(uid, *attrs), attrs->data_type()), - contextLine); + auto mismatch = compareTensor(uid, + *attrs, + expectedFor(uid), + *actual.at(uid), + toleranceFor(tensorLabel(uid, *attrs), attrs->data_type()), + contextLine); if(mismatch.has_value()) { mismatches.push_back(*std::move(mismatch)); diff --git a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp index f57fe1fc7bbe..901acaf08ea6 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/OutputComparison.hpp @@ -45,7 +45,7 @@ using ExpectedTensorLookup = std::function; +/// Keyed on the tensor's label — its name, or "uid=N" when the graph did not give it +/// one. The uid is deliberately not passed: it is not stable between a C++ graph test +/// and the bundle captured from it, so nothing may key a comparison decision on it. +/// The label is both what a TOML `tensors` glob matches and what the failure report +/// prints. +using ToleranceLookup = std::function; /// Compare one tensor. Returns nullopt when it matched. /// diff --git a/dnn-providers/integration-tests/src/integration-tests/layernorm/IntegrationGpuLayernormBackward.cpp b/dnn-providers/integration-tests/src/integration-tests/layernorm/IntegrationGpuLayernormBackward.cpp index 5a6e27aca54a..4e308672e804 100644 --- a/dnn-providers/integration-tests/src/integration-tests/layernorm/IntegrationGpuLayernormBackward.cpp +++ b/dnn-providers/integration-tests/src/integration-tests/layernorm/IntegrationGpuLayernormBackward.cpp @@ -161,9 +161,10 @@ class LayernormBackward this->registerValidator(outputs.dx, this->getTolerance(graphObj, outputs.dx)); // dscale/dbias are reductions over the batch axes: elements land arbitrarily - // near zero through cancellation, so allclose is the wrong check for them. The - // engine's TOML config selects the RMS validator for those two tensors — - // see [[validator_overrides]] in HIP_MLOPS_ENGINE.toml, ALMIOPEN-2561. + // near zero through cancellation, so allclose is the wrong check for them. + // HIP_MLOPS_ENGINE's TOML selects the RMS validator for these two tensors via + // [[validator_overrides]]. The tolerance passed here is the fallback: it is what + // grades them on any engine whose config names no validator for them. this->registerValidator(outputs.dscale, this->getTolerance(graphObj, outputs.dscale)); this->registerValidator(outputs.dbias, this->getTolerance(graphObj, outputs.dbias)); diff --git a/dnn-providers/integration-tests/tests/TestOutputComparison.cpp b/dnn-providers/integration-tests/tests/TestOutputComparison.cpp index d5bfb0556555..6e8d19aba040 100644 --- a/dnn-providers/integration-tests/tests/TestOutputComparison.cpp +++ b/dnn-providers/integration-tests/tests/TestOutputComparison.cpp @@ -220,7 +220,7 @@ TEST(TestOutputComparison, AllOutputsMatchingYieldsNoMismatches) {K_UID_A, K_UID_B}, actual, [&](int64_t uid) -> hipdnn_data_sdk::utilities::ITensor& { return *expected.at(uid); }, - [](int64_t, const std::string&, auto) { return exact(); }, + [](const std::string&, auto) { return exact(); }, "Bundle: b"); EXPECT_TRUE(mismatches.empty()); @@ -246,7 +246,7 @@ TEST(TestOutputComparison, EveryDriftedTensorIsReportedNotJustTheFirst) {K_UID_A, K_UID_B}, actual, [&](int64_t uid) -> hipdnn_data_sdk::utilities::ITensor& { return *expected.at(uid); }, - [](int64_t, const std::string&, auto) { return exact(); }, + [](const std::string&, auto) { return exact(); }, "Bundle: b"); ASSERT_EQ(mismatches.size(), 2u); @@ -276,7 +276,7 @@ TEST(TestOutputComparison, OnlyTheRequestedUidsAreCompared) {K_UID_A}, actual, [&](int64_t uid) -> hipdnn_data_sdk::utilities::ITensor& { return *expected.at(uid); }, - [](int64_t, const std::string&, auto) { return exact(); }, + [](const std::string&, auto) { return exact(); }, "Bundle: b"); EXPECT_TRUE(mismatches.empty()) << "uid 4 drifted but was not in the list"; @@ -285,7 +285,7 @@ TEST(TestOutputComparison, OnlyTheRequestedUidsAreCompared) // --------------------------------------------------------------------------- // Validator kind. ALLCLOSE grades each element against its own magnitude; RMS grades // the tensor against its largest. The two disagree exactly where reduction outputs -// live: an element near zero, on a tensor whose scale is large. ALMIOPEN-2561. +// live: an element near zero, on a tensor whose scale is large. // --------------------------------------------------------------------------- TEST(TestOutputComparison, RmsAcceptsANearZeroElementThatAllcloseRejects) @@ -373,7 +373,7 @@ TEST(TestOutputComparison, ValidatorKindIsChosenPerTensor) {K_UID_A, K_UID_B}, actual, [&](int64_t uid) -> hipdnn_data_sdk::utilities::ITensor& { return *expected.at(uid); }, - [&](int64_t, const std::string& label, auto) { + [&](const std::string& label, auto) { labelsSeen.push_back(label); return label == "uid=4" ? ComparisonTolerance::rms(K_RMS_THRESHOLD) : ComparisonTolerance::allClose(0.0f, 1e-4f); diff --git a/dnn-providers/integration-tests/tests/TestTestSettings.cpp b/dnn-providers/integration-tests/tests/TestTestSettings.cpp index fd5bc403fecf..87dbd1ee0109 100644 --- a/dnn-providers/integration-tests/tests/TestTestSettings.cpp +++ b/dnn-providers/integration-tests/tests/TestTestSettings.cpp @@ -300,7 +300,7 @@ rtol = 5e-2 // // allclose is the default everywhere and this section is the only thing that can // change it, so the parser is strict: an entry that does not say exactly what it -// means is a load error, not a silent fall-back to allclose. ALMIOPEN-2561. +// means is a load error, not a silent fall-back to allclose. // --------------------------------------------------------------------------- namespace @@ -484,6 +484,56 @@ rms_threshold = 1e-4 EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); } +// An entry with 'tensors' but no 'filters' reads as "apply to every test" to a human, +// and is one of the likelier hand-edit mistakes. It is a load error instead, because +// a validator override is scoped to the tests whose numerics justified it. +TEST(TestSettingsValidatorOverrides, ThrowsOnMissingFilters) +{ + const TempTomlFile file(R"( +[meta] +version = 1 + +[[validator_overrides]] +tensors = ["*::DSCALE"] +validator = "rms" +rms_threshold = 1e-4 +)"); + + EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); +} + +TEST(TestSettingsValidatorOverrides, ThrowsOnEmptyFilters) +{ + const TempTomlFile file(R"( +[meta] +version = 1 + +[[validator_overrides]] +filters = [] +tensors = ["*::DSCALE"] +validator = "rms" +rms_threshold = 1e-4 +)"); + + EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); +} + +// A bare value where a table belongs — the shape a mistyped array-of-tables takes. +// Reading it as "no overrides" would drop the entry and quietly grade with allclose. +// The key sits above [meta] so it lands at the document root: written below the +// header it would be a key of [meta] instead, which is a different mistake. +TEST(TestSettingsValidatorOverrides, ThrowsOnNonTableEntry) +{ + const TempTomlFile file(R"( +validator_overrides = ["*LayernormBackward*"] + +[meta] +version = 1 +)"); + + EXPECT_THROW(const TestSettings settings(file.path()), std::runtime_error); +} + // --------------------------------------------------------------------------- // [[test_skips]] parsing // --------------------------------------------------------------------------- diff --git a/dnn-providers/integration-tests/tests/TestTomlGuards.cpp b/dnn-providers/integration-tests/tests/TestTomlGuards.cpp index 0c92e37ff2c7..ea5351a4ea6c 100644 --- a/dnn-providers/integration-tests/tests/TestTomlGuards.cpp +++ b/dnn-providers/integration-tests/tests/TestTomlGuards.cpp @@ -9,6 +9,7 @@ using hipdnn_integration_tests::applyTomlToleranceOverride; using hipdnn_integration_tests::checkTomlSkip; using hipdnn_integration_tests::currentTestName; using hipdnn_integration_tests::findTomlRmsThreshold; +using hipdnn_integration_tests::gradingForTensor; // NOLINTBEGIN(readability-identifier-naming) -- gtest macro-generated names @@ -80,4 +81,17 @@ TEST(TestTomlGuards, FindTomlRmsThresholdReturnsNulloptWhenNoSettings) EXPECT_EQ(findTomlRmsThreshold("SomeTest.Name", "LayernormBackward_0::DSCALE"), std::nullopt); } +// Both harnesses grade every output tensor through gradingForTensor, so the no-config +// answer is the contract the whole suite runs under: the caller's own tolerance, graded +// by allclose. +TEST(TestTomlGuards, GradingForTensorKeepsTheCallersToleranceWhenNoSettings) +{ + const auto grading + = gradingForTensor("SomeTest.Name", "LayernormBackward_0::DSCALE", 1e-3f, 2e-3f); + + EXPECT_EQ(grading.kind, hipdnn_integration_tests::bundle::ValidatorKind::ALLCLOSE); + EXPECT_FLOAT_EQ(grading.atol, 1e-3f); + EXPECT_FLOAT_EQ(grading.rtol, 2e-3f); +} + // NOLINTEND(readability-identifier-naming)