feat(integration-tests): per-tensor validator selection from the engine TOML

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

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,34 @@ 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"`, 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
   `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.
 

dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp

@@ -20,9 +20,7 @@
 #include <hipdnn_frontend/attributes/TensorAttributes.hpp>
 
 #include <hipdnn_plugin_sdk/PluginLogging.hpp>
-#include <hipdnn_test_sdk/utilities/ComparisonReport.hpp>
-#include <hipdnn_test_sdk/utilities/CpuFpReferenceMiopenRmsValidation.hpp>
-#include <hipdnn_test_sdk/utilities/CpuFpReferenceValidation.hpp>
+#include <hipdnn_test_sdk/utilities/ReferenceValidationInterface.hpp>
 #include <hipdnn_test_sdk/utilities/SdkFrontendTypeConversions.hpp>
 #include <hipdnn_test_sdk/utilities/TestTolerances.hpp>
 #include <hipdnn_test_sdk/utilities/TestUtilities.hpp>
@@ -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<Test
     struct TensorValidationEntry
     {
         std::unique_ptr<hipdnn_test_sdk::utilities::IReferenceValidation> 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;
     };
@@ -265,6 +266,10 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam<Test
         registerValidator(attr, tolerance, tolerance);
     }
 
+    // Registers the comparison for one output tensor. allclose at the resolved
+    // tolerance, unless the engine's TOML config names this tensor in a
+    // [[validator_overrides]] entry — the only thing that can select a different
+    // validator. See ALMIOPEN-2561.
     void registerValidator(const std::shared_ptr<hipdnn_frontend::graph::TensorAttributes> attr,
                            float absoluteTolerance,
                            float relativeTolerance)
@@ -273,44 +278,30 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam<Test
         float finalRtol = relativeTolerance;
         applyTomlToleranceOverride(currentTestName(), finalAtol, finalRtol);
 
-        _deferredValidators.emplace_back([this, attr, finalAtol, finalRtol]() {
-            auto sdkDataType
+        const auto testName = currentTestName();
+        _deferredValidators.emplace_back([this, attr, testName, finalAtol, finalRtol]() {
+            const auto sdkDataType
                 = hipdnn_test_sdk::utilities::frontendToSdkDataType(attr->get_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)
+            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))
             {
-                ADD_FAILURE() << "Duplicate validator for tensor " << attr->get_uid() << " ("
-                              << attr->get_name() << "); keeping first registration";
+                tolerance = bundle::ComparisonTolerance::rms(*rmsThreshold);
             }
-            _tensorIdToNameMap.insert({attr->get_uid(), attr->get_name()});
-        });
-    }
 
-    void registerRmsValidator(const std::shared_ptr<hipdnn_frontend::graph::TensorAttributes> 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}});
+            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()});
         });
@@ -423,33 +414,21 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam<Test
             }
 
             auto& entry = entryIt->second;
-            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{
-                    "Test: " + currentTestName(), tensorLabel, dtypeName, entry.atol, 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());
             }
         }
     }
@@ -600,7 +579,7 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam<Test
         auto valIt = _tensorValidationMap.find(tensorId);
         if(valIt != _tensorValidationMap.end())
         {
-            return valIt->second.name;
+            return valIt->second.label;
         }
         auto nameIt = _tensorIdToNameMap.find(tensorId);
         if(nameIt != _tensorIdToNameMap.end())

dnn-providers/integration-tests/src/harness/TestConfig.hpp

@@ -326,6 +326,20 @@ class TestConfig
         return _testSettings->findToleranceOverride(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<ValidatorOverride> 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