feat(hipdnn-integration-tests): add --benchmark timing for verified graphs

dnn-providers/integration-tests/docs/README.md

@@ -8,7 +8,7 @@ developer who must run the suite, read its output, or add to it.
 | Document | Read it to… |
 |---|---|
 | [File Formats](file-formats.md) | Understand every file in the bundle tree and the provider configs: graphs, template sweeps, metadata, golden data and DVC pointers, support-claim sidecars, per-engine TOML, tier YAML. Also learn how test names are derived. |
-| [Running the Tests](running-tests.md) | Run the suite (CTest, check targets, the binary directly), learn what the tiers and filters select, and read the output: coverage summary, support-claim summary, hard stops, and runs that look green but are not. |
+| [Running the Tests](running-tests.md) | Run the suite (CTest, check targets, the binary directly), time the engine with `--benchmark`, learn what the tiers and filters select, and read the output: coverage summary, support-claim summary, hard stops, and runs that look green but are not. |
 | [Adding Tests and Updating Claims](adding-tests.md) | Add a bundle (end to end, including a new op), add golden data, update support claims when a run reports `unclaimed_support`, record an engine limitation, add an engine lane, or add a C++ test when a bundle cannot express it. |
 | [Support Claim Enforcement](support-claim-enforcement.md) | For harness maintainers: the claim verdict model, the lifecycle inside `TestBody()`, the coverage ladder, and who owns what in the harness. |
 | [Support Matrix and Feature Coverage](support-matrix.md) | Render the human-readable matrix of which engine supports which graph on which GPU and OS, validate the sidecars it is built from, see the pre-commit and CI checks around it, and measure per-op feature coverage. |
@@ -51,6 +51,7 @@ other rows do not, so they are easy to miss.
 | `UNVERIFIABLE BUNDLES` / `Unverifiable: … tried: …` | **yes** (those tests failed) | The engine ran the bundle, but no oracle could check its output: no golden data and no reference executor for the op. Nothing was tested. | [Unverifiable bundles](running-tests.md#unverifiable-bundles); fix: [When no reference executor supports the op](adding-tests.md#when-no-reference-executor-supports-the-op) |
 | `Error: zero tests ran.` | **yes** | A discovery or filter misconfiguration, or, with GTest sharding, an empty shard. | [Hard stops](running-tests.md#hard-stops) |
 | `FATAL: --enforce-support-claims is active and …` | **yes** | Enforcement verified nothing. | [Hard stops](running-tests.md#hard-stops) |
+| `benchmark failed: …` | **yes** (that test failed) | Under `--benchmark`, the verified graph failed when it ran again, or the timer returned too few measurements. | [Time the engine](running-tests.md#4-time-the-engine---benchmark) |
 | `unclaimed_support` in the `SUPPORT CLAIM SUMMARY` | no | The engine accepts graphs that no sidecar claims. Record them. | [Updating support claims](adding-tests.md#updating-support-claims) |
 | `Skipped:` equal to (or close to) the total | no | Green, but little or nothing was tested. Read the skip reasons. | [Test coverage summary](running-tests.md#test-coverage-summary) |
 | `WARNING: Bundle tests are enabled but …` / `WARNING: No bundles could be loaded from …` | no | No bundles registered. Only compiled-in C++ tests ran. | [Troubleshooting](running-tests.md#troubleshooting) |

dnn-providers/integration-tests/docs/running-tests.md

@@ -128,6 +128,78 @@ build/bin/hipdnn_integration_tests \
   runs anything.
 - Standard GTest sharding (`GTEST_TOTAL_SHARDS` / `GTEST_SHARD_INDEX`) works.
 
+### 4. Time the engine (`--benchmark`)
+
+`--benchmark` adds a performance measurement to each test. The test first
+verifies the engine output in the usual way. If the output passes, the harness
+times the same engine on the same tensors with `Graph::autotune()`, limited to
+that one engine. The C++ graph tests and the bundles both support it.
+
+A test is not timed when it skips, fails, or passes without output
+verification (a bundle at enforcement level APPLICABLE or BUILDABLE). Such a
+test prints no `[ BENCHMARK ]` line. A filter that selects only such tests
+gives a green run with no timings, so count the `[ BENCHMARK ]` lines.
+
+```bash
+build/bin/hipdnn_integration_tests \
+    --test-article build/lib/hipdnn_plugins/engines/libmiopen_plugin.so \
+    --test-engine  MIOPEN_ENGINE \
+    --gtest_filter='quick_ConvolutionFwd_Default.*' \
+    --benchmark \
+    --gtest_output=json:benchmark.json
+```
+
+Each timed test prints one line:
+
+```text
+[ BENCHMARK ] quick_ConvolutionFwd_Default.8_16_16_16_bfp16_nhwc_dil1x1_postpad0x0_prepad0x0: engine=MIOPEN_ENGINE tuned=0 timing=DEVICE_ONLY samples=20 min=0.0181 ms robust=0.0184 ms mean=0.0184 ms stddev=0.0002 ms
+```
+
+- Autotune runs `--benchmark-warmup` untimed executions (default 30). Then it
+  runs `--benchmark-iterations` timed executions (default 20), with the
+  `FIXED_AVERAGE` strategy and the same stall gate, retry and statistics
+  rules as any other autotune call.
+- The GPU idles between tests, so each test starts cold. Too few warmup
+  executions make the times read high: on MI210 the quick-tier bundles read
+  4-60% slow with 1-10 warmups and settle after about 30. More timed
+  iterations do not fix this; more warmup does. Raise `--benchmark-warmup`
+  if `stddev` stays large compared with `mean`.
+- `robust` is the mean after autotune drops the slow tail: what the engine
+  usually costs. `min`, `mean` and `stddev` are over all timed samples.
+- `tuned=0` means that the engine chose its solution with its heuristics. The
+  harness sets no knobs, so this is the default. To time the solution that the
+  engine picks by benchmarking its candidates, set
+  `HIPDNN_FORCE_BENCHMARKING=1`. The line then shows `tuned=1` for engines
+  that support the `global.benchmarking` knob (MIOpen, the hip-kernel
+  ingestor). Engines without the knob, such as hipBLASLt, still show
+  `tuned=0`. The selection runs during the verification execution, so its
+  cost is not in the timed samples.
+- `timing=DEVICE_ONLY` means that the stream was stalled, so the times exclude
+  host submission overhead. `timing=UNSTALLED` means that autotune could not
+  stall the stream and re-measured without the stall gate, so the times can
+  include host time. Do not compare the two kinds.
+- The test fails with `benchmark failed: …` when autotune cannot benchmark the
+  engine, for example when an execution fails or too many timed executions
+  return no measurement.
+- `--gtest_output=json:<file>` or `xml:<file>` writes the same values as the
+  test properties `benchmark_engine`, `benchmark_tuned`, `benchmark_timing`,
+  `benchmark_samples`, `benchmark_min_ms`, `benchmark_robust_ms`,
+  `benchmark_mean_ms` and `benchmark_stddev_ms`.
+- `engine=` names the engine that ran. To compare engines, run once for each
+  `--test-engine`. Without `--test-engine`, hipDNN selects the engine.
+- To time your own graph, write it as a single-graph bundle
+  (`<dir>/MyGraph/MyGraph.json`, [format](file-formats.md)). Then pass
+  `--golden-data-dir <dir> --gtest_filter='MyGraph*'`. Without golden data, a
+  reference executor verifies the output. To convert a C++ graph test into a
+  bundle, use `--capture-bundles`.
+- Verification runs before timing. For large shapes, the CPU reference can be
+  slow. Use `--verification-mode gpu` when a GPU reference exists for the op.
+- To see the kernels that the engine launched, run the same command after
+  `rocprofv3 --kernel-trace --stats --output-format csv -d <dir> --`. In
+  `<dir>/*/*_kernel_stats.csv`, an engine kernel of a single-kernel graph has
+  1 + warmup + iterations calls. The other kernels fill the inputs or run the
+  reference.
+
 ### The build reads a copy of the bundles
 
 CMake copies `integration-test-bundles/` into `build/lib/integration-test-bundles/`
@@ -155,6 +227,9 @@ Otherwise, you test the old tree.
 | `--fail-on-unsupported` | | off | Fail instead of skip when no engine supports a graph. **C++ graph tests only.** |
 | `--generate-support-matrix[=file]` | | off | Write a markdown matrix. **C++ graph tests only**; on a default build it writes an empty table. |
 | `--capture-bundles <dir>` | | off | Dump compiled-in C++ graph tests as bundle JSON (see [Adding Tests](adding-tests.md#from-an-existing-c-graph-test)). |
+| `--benchmark` | | off | Time each verified test. See [Time the engine](#4-time-the-engine---benchmark). |
+| `--benchmark-iterations` | | 20 | Timed executions per test. Requires `--benchmark`. |
+| `--benchmark-warmup` | | 30 | Untimed executions before timing. Requires `--benchmark`. |
 | `--gtest_*` | | | Passed through to GTest. |
 
 The binary refuses these flag combinations (exit 1):
@@ -164,6 +239,10 @@ The binary refuses these flag combinations (exit 1):
 - `--write-support-claims` together with `--enforce-support-claims` or `=true`.
 - `--write-support-claims` without `--test-article`, or without a bundle root
   (`--golden-data-dir` or `HIPDNN_TEST_GOLDEN_DATA_DIR`).
+- `--benchmark-iterations` or `--benchmark-warmup` without `--benchmark`, or
+  `--benchmark-iterations` less than 1, or `--benchmark-warmup` less than 0.
+- `--benchmark` together with `--skip-graph-validation`, `--capture-bundles`
+  or `--write-support-claims`. These modes do not execute the graph.
 
 `--enforce-support-claims=false` is accepted in every combination. If
 enforcement is only inherited and there is no `--test-engine`, the run

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

@@ -0,0 +1,80 @@
+// Copyright © Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <optional>
+#include <string>
+
+namespace hipdnn_integration_tests
+{
+
+/// How --benchmark times one verified graph.
+struct BenchmarkSettings
+{
+    static constexpr int DEFAULT_TIMED_ITERATIONS = 20;
+    // The GPU idles between tests (input fill, reference, comparison), so each test
+    // starts cold. On MI210, quick-tier bundles read 4-60% slow with 1-10 warmup
+    // executions and settle after about 30. More timed iterations do not fix this.
+    static constexpr int DEFAULT_WARMUP_ITERATIONS = 30;
+
+    int timedIterations = DEFAULT_TIMED_ITERATIONS;
+    int warmupIterations = DEFAULT_WARMUP_ITERATIONS;
+};
+
+/// The benchmark flags as typed, plus the flags that stop a test before it executes.
+struct BenchmarkRequest
+{
+    bool enabled = false;
+    std::optional<int> timedIterations;
+    std::optional<int> warmupIterations;
+    bool skipGraphValidation = false;
+    bool captureBundles = false;
+    bool writeSupportClaims = false;
+};
+
+struct BenchmarkResolution
+{
+    std::optional<BenchmarkSettings> settings;
+    std::optional<std::string> error;
+};
+
+/// Settings for --benchmark, or the reason the flags cannot be used together.
+///
+/// A mode that stops a test before it executes would leave --benchmark timing
+/// nothing and the run green. Those combinations are refused instead.
+inline BenchmarkResolution resolveBenchmark(const BenchmarkRequest& request)
+{
+    if(!request.enabled)
+    {
+        if(request.timedIterations.has_value() || request.warmupIterations.has_value())
+        {
+            return {std::nullopt,
+                    "--benchmark-iterations and --benchmark-warmup require --benchmark"};
+        }
+        return {};
+    }
+
+    if(request.skipGraphValidation || request.captureBundles || request.writeSupportClaims)
+    {
+        return {std::nullopt,
+                "--benchmark cannot be combined with --skip-graph-validation, "
+                "--capture-bundles or --write-support-claims: those modes do not "
+                "execute the graph"};
+    }
+
+    BenchmarkSettings settings;
+    settings.timedIterations = request.timedIterations.value_or(settings.timedIterations);
+    settings.warmupIterations = request.warmupIterations.value_or(settings.warmupIterations);
+    if(settings.timedIterations < 1)
+    {
+        return {std::nullopt, "--benchmark-iterations must be at least 1"};
+    }
+    if(settings.warmupIterations < 0)
+    {
+        return {std::nullopt, "--benchmark-warmup must not be negative"};
+    }
+    return {settings, std::nullopt};
+}
+
+} // namespace hipdnn_integration_tests

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

@@ -0,0 +1,124 @@
+// Copyright © Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <cstdint>
+#include <iomanip>
+#include <iostream>
+#include <sstream>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+#include <gtest/gtest.h>
+
+#include <hipdnn_frontend/Error.hpp>
+#include <hipdnn_frontend/Types.hpp>
+#include <hipdnn_frontend/autotune/AutotuneTypes.hpp>
+#include <hipdnn_plugin_sdk/GlobalKnobDefines.hpp>
+
+#include "harness/CurrentTestName.hpp"
+
+namespace hipdnn_integration_tests
+{
+
+inline const char* toString(hipdnn_frontend::TimingQuality quality)
+{
+    switch(quality)
+    {
+    case hipdnn_frontend::TimingQuality::DEVICE_ONLY:
+        return "DEVICE_ONLY";
+    case hipdnn_frontend::TimingQuality::UNSTALLED:
+        return "UNSTALLED";
+    case hipdnn_frontend::TimingQuality::INVALID:
+        return "INVALID";
+    default:
+        // An unknown value must not print as a valid quality.
+        throw std::invalid_argument("unknown TimingQuality value "
+                                    + std::to_string(static_cast<int>(quality)));
+    }
+}
+
+/// Picks the result of the engine under test from Graph::autotune() output.
+///
+/// autotune() returns an error when no engine succeeds, but the engine's own result
+/// says why, so that result wins over `status`. A result for another engine, or one
+/// the caller excluded, is never used: it would report the wrong engine's timing.
+inline hipdnn_frontend::AutotuneResult
+    selectEngineResult(int64_t engineId,
+                       const std::vector<hipdnn_frontend::AutotuneResult>& results,
+                       const hipdnn_frontend::Error& status)
+{
+    for(const auto& result : results)
+    {
+        if(result.engineId == engineId && !result.excludedByCaller)
+        {
+            return result;
+        }
+    }
+    hipdnn_frontend::AutotuneResult failed;
+    failed.engineId = engineId;
+    failed.errorMessage = status.is_good() ? "autotune returned no result for the engine under test"
+                                           : status.get_message();
+    return failed;
+}
+
+/// True when the engine selected its solution by benchmarking candidates rather than
+/// by its heuristics. The harness sets no knobs, so global.benchmarking keeps its
+/// default (0 in every hipDNN provider). Only HIPDNN_FORCE_BENCHMARKING turns it on,
+/// and only for engines that expose the knob (AutotuneResult::supportsExhaustive).
+inline bool selectedByBenchmarking(const hipdnn_frontend::AutotuneResult& result)
+{
+    return result.supportsExhaustive
+           && hipdnn_plugin_sdk::benchmarkingOverrideFromEnv().value_or(false);
+}
+
+/// Publishes the autotune result of one engine for the running test: a
+/// "[ BENCHMARK ]" line on stdout and benchmark_* properties in --gtest_output
+/// files. A result that did not succeed fails the test.
+inline void reportBenchmark(const hipdnn_frontend::AutotuneResult& result)
+{
+    if(!result.succeeded)
+    {
+        ADD_FAILURE() << "benchmark failed: "
+                      << (result.errorMessage.empty() ? "the engine was not benchmarked"
+                                                      : result.errorMessage);
+        return;
+    }
+
+    const auto ms = [](float value) {
+        std::ostringstream out;
+        out << std::fixed << std::setprecision(4) << value;
+        return out.str();
+    };
+
+    const auto minMs = ms(result.minTimeMs);
+    const auto robustMs = ms(result.robustTimeMs);
+    const auto meanMs = ms(result.avgTimeMs);
+    const auto stddevMs = ms(result.stddevMs);
+    const int tuned = selectedByBenchmarking(result) ? 1 : 0;
+
+    ::testing::Test::RecordProperty("benchmark_engine", result.engineName);
+    ::testing::Test::RecordProperty("benchmark_timing", toString(result.timingQuality));
+    ::testing::Test::RecordProperty("benchmark_samples", result.iterationsRun);
+    ::testing::Test::RecordProperty("benchmark_min_ms", minMs);
+    ::testing::Test::RecordProperty("benchmark_robust_ms", robustMs);
+    ::testing::Test::RecordProperty("benchmark_mean_ms", meanMs);
+    ::testing::Test::RecordProperty("benchmark_stddev_ms", stddevMs);
+    ::testing::Test::RecordProperty("benchmark_tuned", tuned);
+
+    auto testName = currentTestName();
+    if(testName.empty())
+    {
+        testName = "<unknown test>";
+    }
+
+    std::cout << "[ BENCHMARK ] " << testName << ": engine=" << result.engineName
+              << " tuned=" << tuned << " timing=" << toString(result.timingQuality)
+              << " samples=" << result.iterationsRun << " min=" << minMs
+              << " ms robust=" << robustMs << " ms mean=" << meanMs << " ms stddev=" << stddevMs
+              << " ms\n";
+}
+
+} // namespace hipdnn_integration_tests

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

@@ -0,0 +1,24 @@
+// Copyright © Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <string>
+
+#include <gtest/gtest.h>
+
+namespace hipdnn_integration_tests
+{
+
+/// "Suite.Name" of the running gtest test, or an empty string outside a test.
+inline std::string currentTestName()
+{
+    auto* info = ::testing::UnitTest::GetInstance()->current_test_info();
+    if(info == nullptr)
+    {
+        return {};
+    }
+    return std::string(info->test_suite_name()) + "." + info->name();
+}
+
+} // namespace hipdnn_integration_tests