From 5d3bdf2c43bfaf8136c8f553e43a2eaef5d03716 Mon Sep 17 00:00:00 2001 From: Samuel Reeder Date: Mon, 19 Oct 2026 04:44:18 +0000 Subject: [PATCH 1/6] feat(hipdnn-integration-tests): add --benchmark timing for verified graphs Add an opt-in --benchmark mode to hipdnn_integration_tests. After a bundle or C++ graph test verifies the engine output, the harness runs the same plan on the same tensors again and times it with Graph::execute_timed_ext(). Each timed test prints min/median/mean device time, the engine and the timing quality, and records them as benchmark_* gtest properties. --benchmark-iterations and --benchmark-warmup set the sample and warmup counts. Combinations that do not execute the graph are refused. Closes #11652 --- .../integration-tests/docs/README.md | 3 +- .../integration-tests/docs/running-tests.md | 59 +++ .../src/harness/Benchmark.hpp | 259 ++++++++++ .../src/harness/GraphBenchmark.hpp | 68 +++ .../IntegrationGraphVerificationHarness.hpp | 9 + .../src/harness/TestConfig.hpp | 11 + .../bundle/FrontendGraphEngineRunner.cpp | 14 + .../bundle/FrontendGraphEngineRunner.hpp | 9 +- .../src/harness/bundle/HarnessPolicy.hpp | 4 + .../src/harness/bundle/IGraphEngineRunner.hpp | 8 + .../IntegrationBundleVerificationHarness.cpp | 32 +- .../IntegrationBundleVerificationHarness.hpp | 14 + .../src/harness/bundle/ProductionPolicy.cpp | 1 + dnn-providers/integration-tests/src/main.cpp | 29 ++ .../integration-tests/tests/CMakeLists.txt | 1 + .../integration-tests/tests/TestBenchmark.cpp | 453 ++++++++++++++++++ .../tests/mocks/MockGraphEngineRunner.hpp | 7 + 17 files changed, 974 insertions(+), 7 deletions(-) create mode 100644 dnn-providers/integration-tests/src/harness/Benchmark.hpp create mode 100644 dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp create mode 100644 dnn-providers/integration-tests/tests/TestBenchmark.cpp diff --git a/dnn-providers/integration-tests/docs/README.md b/dnn-providers/integration-tests/docs/README.md index a583ee7b3373..8f63ef74d07d 100644 --- a/dnn-providers/integration-tests/docs/README.md +++ b/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. | @@ -50,6 +50,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) | diff --git a/dnn-providers/integration-tests/docs/running-tests.md b/dnn-providers/integration-tests/docs/running-tests.md index 96689e32097b..281ac4f4a55e 100644 --- a/dnn-providers/integration-tests/docs/running-tests.md +++ b/dnn-providers/integration-tests/docs/running-tests.md @@ -128,6 +128,58 @@ 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 +runs the same plan on the same tensors again and times it. A test that skips or +fails is not timed. The C++ graph tests and the bundles both support it. + +```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 --benchmark-iterations 50 \ + --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 timing=DEVICE_ONLY samples=20 min=0.0182 ms median=0.0184 ms mean=0.0185 ms +``` + +- The harness runs `--benchmark-warmup` untimed executions (default 3). Then + it collects `--benchmark-iterations` timed executions (default 20) with + `Graph::execute_timed_ext()`. +- `timing=DEVICE_ONLY` means that the stream was stalled for every sample, so + the times exclude host submission overhead. `timing=UNSTALLED` means that at + least one sample was not stalled and can include host time. Do not compare + the two kinds. +- A timed execution that returns no measurement (a stall watchdog timeout) is + retried. The line then shows `discarded=N`. The harness makes at most two + attempts for each requested sample. The test fails with + `benchmark failed: …` if it cannot collect all the samples in those attempts, + or if an execution fails. +- `--gtest_output=json:` or `xml:` writes the same values as the + test properties `benchmark_engine`, `benchmark_timing`, `benchmark_samples`, + `benchmark_min_ms`, `benchmark_median_ms` and `benchmark_mean_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 + (`/MyGraph/MyGraph.json`, [format](file-formats.md)). Then pass + `--golden-data-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 --`. In + `/*/*_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 +207,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 ` | | 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` | | 3 | Untimed executions before timing. Requires `--benchmark`. | | `--gtest_*` | | | Passed through to GTest. | The binary refuses these flag combinations (exit 1): @@ -164,6 +219,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 diff --git a/dnn-providers/integration-tests/src/harness/Benchmark.hpp b/dnn-providers/integration-tests/src/harness/Benchmark.hpp new file mode 100644 index 000000000000..129b9ea3e5ff --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/Benchmark.hpp @@ -0,0 +1,259 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include +#include + +namespace hipdnn_integration_tests +{ + +/// How --benchmark times one verified graph. +struct BenchmarkSettings +{ + static constexpr int DEFAULT_TIMED_ITERATIONS = 20; + static constexpr int DEFAULT_WARMUP_ITERATIONS = 3; + + 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 timedIterations; + std::optional warmupIterations; + bool skipGraphValidation = false; + bool captureBundles = false; + bool writeSupportClaims = false; +}; + +struct BenchmarkResolution +{ + std::optional settings; + std::optional 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}; +} + +/// The timings of one graph on one engine. +struct BenchmarkResult +{ + std::string engine; + std::vector samplesMs; + + /// DEVICE_ONLY only when every sample was. One UNSTALLED sample makes the set + /// UNSTALLED, because that sample can include host submission time. + hipdnn_frontend::TimingQuality quality = hipdnn_frontend::TimingQuality::INVALID; + + /// Timed runs that returned no measurement (stall watchdog timeout or a negative + /// reading). They are retried and are not in samplesMs. + int discardedSamples = 0; + + /// Why the benchmark failed. Empty on success. + std::string error; + + bool ok() const + { + return error.empty() && !samplesMs.empty(); + } + + float minMs() const + { + return samplesMs.empty() ? 0.0f : *std::min_element(samplesMs.begin(), samplesMs.end()); + } + + float medianMs() const + { + if(samplesMs.empty()) + { + return 0.0f; + } + auto sorted = samplesMs; + std::sort(sorted.begin(), sorted.end()); + const std::size_t mid = sorted.size() / 2; + return sorted.size() % 2 == 1 ? sorted[mid] : (sorted[mid - 1] + sorted[mid]) / 2.0f; + } + + float meanMs() const + { + if(samplesMs.empty()) + { + return 0.0f; + } + const double sum = std::accumulate(samplesMs.begin(), samplesMs.end(), 0.0); + return static_cast(sum / static_cast(samplesMs.size())); + } +}; + +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: + default: + return "INVALID"; + } +} + +using BenchmarkRunOnce = std::function; +using BenchmarkTimeOnce = std::function; + +/// Runs `warmupIterations` untimed executions, then collects `timedIterations` valid +/// samples. A timed run with no measurement is retried, up to `timedIterations` extra +/// attempts. Any execution error stops the benchmark and is returned in `error`. +/// +/// Takes callables rather than a graph so the loop is testable without a device. +inline BenchmarkResult runBenchmarkLoop(const BenchmarkSettings& settings, + const BenchmarkRunOnce& runOnce, + const BenchmarkTimeOnce& timeOnce) +{ + BenchmarkResult result; + + for(int i = 0; i < settings.warmupIterations; ++i) + { + if(auto status = runOnce(); !status.is_good()) + { + result.error = "warmup execution failed: " + status.get_message(); + return result; + } + } + + bool anyUnstalled = false; + const int maxAttempts = 2 * settings.timedIterations; + for(int attempt = 0; attempt < maxAttempts + && static_cast(result.samplesMs.size()) < settings.timedIterations; + ++attempt) + { + hipdnn_frontend::ExecutionTiming timing; + if(auto status = timeOnce(timing); !status.is_good()) + { + result.error = "timed execution failed: " + status.get_message(); + result.samplesMs.clear(); + return result; + } + if(timing.quality == hipdnn_frontend::TimingQuality::INVALID || !timing.elapsedMs) + { + ++result.discardedSamples; + continue; + } + anyUnstalled = anyUnstalled || timing.quality == hipdnn_frontend::TimingQuality::UNSTALLED; + result.samplesMs.push_back(*timing.elapsedMs); + } + + if(static_cast(result.samplesMs.size()) < settings.timedIterations) + { + result.error = std::to_string(result.discardedSamples) + " of " + + std::to_string(maxAttempts) + + " timed executions returned no measurement; collected " + + std::to_string(result.samplesMs.size()) + " of " + + std::to_string(settings.timedIterations) + " samples"; + result.samplesMs.clear(); + return result; + } + + result.quality = anyUnstalled ? hipdnn_frontend::TimingQuality::UNSTALLED + : hipdnn_frontend::TimingQuality::DEVICE_ONLY; + return result; +} + +/// Publishes a benchmark for the running test: a "[ BENCHMARK ]" line on stdout and +/// benchmark_* properties in --gtest_output files. A failed benchmark fails the test. +inline void reportBenchmark(const BenchmarkResult& result) +{ + if(!result.ok()) + { + ADD_FAILURE() << "benchmark failed: " + << (result.error.empty() ? "no samples were collected" : result.error); + return; + } + + const auto ms = [](float value) { + std::ostringstream out; + out << std::fixed << std::setprecision(4) << value; + return out.str(); + }; + + const auto minMs = ms(result.minMs()); + const auto medianMs = ms(result.medianMs()); + const auto meanMs = ms(result.meanMs()); + + ::testing::Test::RecordProperty("benchmark_engine", result.engine); + ::testing::Test::RecordProperty("benchmark_timing", toString(result.quality)); + ::testing::Test::RecordProperty("benchmark_samples", static_cast(result.samplesMs.size())); + ::testing::Test::RecordProperty("benchmark_min_ms", minMs); + ::testing::Test::RecordProperty("benchmark_median_ms", medianMs); + ::testing::Test::RecordProperty("benchmark_mean_ms", meanMs); + + std::string testName = ""; + if(const auto* info = ::testing::UnitTest::GetInstance()->current_test_info()) + { + testName = std::string(info->test_suite_name()) + "." + info->name(); + } + + std::cout << "[ BENCHMARK ] " << testName << ": engine=" << result.engine + << " timing=" << toString(result.quality) << " samples=" << result.samplesMs.size() + << " min=" << minMs << " ms median=" << medianMs << " ms mean=" << meanMs << " ms"; + if(result.discardedSamples > 0) + { + std::cout << " discarded=" << result.discardedSamples; + } + std::cout << '\n'; +} + +} // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp new file mode 100644 index 000000000000..1eb9f1fd35e7 --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp @@ -0,0 +1,68 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include + +#include +#include +#include + +#include "harness/Benchmark.hpp" +#include "harness/VariantPack.hpp" +#include "harness/bundle/LoadedEngineTable.hpp" + +namespace hipdnn_integration_tests +{ + +/// Times the plan `graph` already built, on the inputs and outputs in `variantPack`. +inline BenchmarkResult benchmarkGraph(hipdnn_frontend::graph::Graph& graph, + hipdnnHandle_t handle, + VariantPack& variantPack, + const BenchmarkSettings& settings) +{ + BenchmarkResult failed; + + int64_t engineId = -1; + if(auto status = graph.get_execution_plan_engine_id(engineId); !status.is_good()) + { + failed.error = status.get_message(); + return failed; + } + + int64_t workspaceSize = 0; + if(auto status = graph.get_workspace_size(workspaceSize); !status.is_good()) + { + failed.error = status.get_message(); + return failed; + } + if(workspaceSize < 0) + { + failed.error = "engine reported a negative workspace size"; + return failed; + } + const hipdnn_data_sdk::utilities::Workspace workspace(static_cast(workspaceSize)); + + auto result = runBenchmarkLoop( + settings, + [&] { return graph.execute(handle, variantPack, workspace.get()); }, + [&](hipdnn_frontend::ExecutionTiming& timing) { + return graph.execute_timed_ext(handle, variantPack, workspace.get(), timing); + }); + + // Plugin engines are not in the static name registry, so ask the loaded table first. + result.engine = hipdnn_data_sdk::utilities::engineNameOrHex(engineId); + for(const auto& engine : bundle::LoadedEngineTable::get().all()) + { + if(engine.id == engineId) + { + result.engine = engine.name; + } + } + return result; +} + +} // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp index 0d56009fa259..0d88eae690c4 100644 --- a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp +++ b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp @@ -30,6 +30,7 @@ #include #include +#include "harness/GraphBenchmark.hpp" #include "harness/GraphDescription.hpp" #include "harness/IReferenceGraphExecutor.hpp" #include "harness/ReferenceGraphExecutorFactory.hpp" @@ -260,6 +261,14 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam attr, diff --git a/dnn-providers/integration-tests/src/harness/TestConfig.hpp b/dnn-providers/integration-tests/src/harness/TestConfig.hpp index eceee7780da8..0fffd7bd85bb 100644 --- a/dnn-providers/integration-tests/src/harness/TestConfig.hpp +++ b/dnn-providers/integration-tests/src/harness/TestConfig.hpp @@ -16,6 +16,7 @@ #include #include "common/PlatformUtils.hpp" +#include "harness/Benchmark.hpp" #include "harness/TestSettings.hpp" #include "harness/ValidationSite.hpp" @@ -164,6 +165,7 @@ struct TestConfigOptions // failed by one. bool enforceSupportClaims = false; bool writeSupportClaims = false; + std::optional benchmark; }; // Singleton class for storing CLI-based test configuration. @@ -232,6 +234,7 @@ class TestConfig instance._enforceSupportClaims = opts.enforceSupportClaims; instance._writeSupportClaims = opts.writeSupportClaims; + instance._benchmark = opts.benchmark; // Golden bundle configuration — default is ON; env var can override. instance._allowBundles = opts.allowBundles; @@ -493,6 +496,13 @@ class TestConfig return _captureDir.value(); } + /// Set by --benchmark: how to time each verified graph. Empty when off. + const std::optional& benchmarkSettings() const + { + throwIfNotInitialized(); + return _benchmark; + } + private: TestConfig() = default; @@ -512,6 +522,7 @@ class TestConfig std::optional _verificationMode; std::optional _validatorDevice; std::optional _captureDir; + std::optional _benchmark; std::string _currentArch; std::size_t _currentDeviceVramMb = 0; std::string _currentPlatform; diff --git a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp index e764e3c5f480..ab39a421ace9 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp @@ -10,6 +10,7 @@ #include #include "harness/EngineNotApplicableError.hpp" +#include "harness/GraphBenchmark.hpp" #include "harness/SharedHandle.hpp" #include "harness/bundle/IntegrationTestBundle.hpp" @@ -150,4 +151,17 @@ EngineOpResult return result; } +BenchmarkResult FrontendGraphEngineRunner::benchmark(GraphSession& session, + VariantPack& variantPack, + const BenchmarkSettings& settings) +{ + if(session.graph == nullptr) + { + BenchmarkResult failed; + failed.error = "openGraph() produced no graph to benchmark"; + return failed; + } + return benchmarkGraph(*session.graph, getSharedHandle(), variantPack, settings); +} + } // namespace hipdnn_integration_tests::bundle diff --git a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.hpp b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.hpp index d9328e8e7363..3672820257dc 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.hpp @@ -11,13 +11,14 @@ namespace hipdnn_integration_tests::bundle /// The production IGraphEngineRunner: drives a real hipdnn_frontend::graph::Graph /// on the shared handle. /// -/// Implements the three steps the harness needs from an engine, in the order it -/// needs them: +/// Implements the steps the harness needs from an engine, in the order it needs +/// them: /// - openGraph() — one from_binary, then the single get_ranked_engine_ids() /// query whose answer becomes GraphSession::engines. /// - buildPlans() — create_execution_plans / check_support / build_plans, with a /// late decline returned as an answer rather than thrown. /// - execute() — builds plans if needed, then runs the graph on a variant pack. +/// - benchmark() — times the plan execute() built (--benchmark only). /// /// This is the only place in the harness that needs a handle, a device, or a /// loaded engine plugin. Kept in its own translation unit so the unit-test binary @@ -35,6 +36,10 @@ class FrontendGraphEngineRunner : public IGraphEngineRunner EngineOpResult execute(GraphSession& session, const std::optional& engineUnderTest, VariantPack& variantPack) override; + + BenchmarkResult benchmark(GraphSession& session, + VariantPack& variantPack, + const BenchmarkSettings& settings) override; }; } // namespace hipdnn_integration_tests::bundle diff --git a/dnn-providers/integration-tests/src/harness/bundle/HarnessPolicy.hpp b/dnn-providers/integration-tests/src/harness/bundle/HarnessPolicy.hpp index 0a48b8285e10..f3e4c47847be 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/HarnessPolicy.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/HarnessPolicy.hpp @@ -9,6 +9,7 @@ #include #include +#include "harness/Benchmark.hpp" #include "harness/TestConfig.hpp" namespace hipdnn_integration_tests::bundle @@ -125,6 +126,9 @@ struct HarnessPolicy std::string platform; std::size_t deviceVramMb = 0; + /// Set by --benchmark: time each verified graph with these settings. + std::optional benchmark; + bool useDevice() const { return placement == TensorPlacement::DEVICE; diff --git a/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp b/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp index 2053e34f6b28..aa69bda7c260 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp @@ -9,6 +9,7 @@ #include #include +#include "harness/Benchmark.hpp" #include "harness/VariantPack.hpp" #include "harness/bundle/GraphSession.hpp" #include "harness/bundle/LoadedEngine.hpp" @@ -92,6 +93,13 @@ class IGraphEngineRunner const std::optional& engineUnderTest, VariantPack& variantPack) = 0; + + /// Times the plan execute() built, on `variantPack`. Called only after the + /// outputs of execute() were verified. + virtual BenchmarkResult benchmark(GraphSession& session, + VariantPack& variantPack, + const BenchmarkSettings& settings) + = 0; }; } // namespace hipdnn_integration_tests::bundle diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp index 65a744c186ae..8b4ec8dd5d22 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp @@ -480,7 +480,9 @@ VerificationOutcome IntegrationBundleVerificationHarness::runGoldenMode(GraphSes { return engineDidNotRun(engine); } - return compareAgainstGolden(engine.outputs); + auto outcome = compareAgainstGolden(engine.outputs); + benchmarkIfVerified(session, engine.outputs, outcome); + return outcome; } VerificationOutcome IntegrationBundleVerificationHarness::runReferenceMode(GraphSession& session) @@ -495,17 +497,26 @@ VerificationOutcome IntegrationBundleVerificationHarness::runReferenceMode(Graph } auto oracles = resolveOracles(_deps.policy.mode); + std::optional outcome; if(oracles.mode == VerificationMode::AUTO) { // Golden data is auto mode's first oracle; when the bundle has it, no // reference is consulted. if(_bundle->hasGoldenOutputs) { - return compareAgainstGolden(engine.outputs); + outcome = compareAgainstGolden(engine.outputs); + } + else + { + oracles.declined("golden (absent)"); } - oracles.declined("golden (absent)"); } - return runOracleChain(engine.outputs, oracles); + if(!outcome.has_value()) + { + outcome = runOracleChain(engine.outputs, oracles); + } + benchmarkIfVerified(session, engine.outputs, *outcome); + return *outcome; } VerificationOutcome @@ -860,6 +871,19 @@ void IntegrationBundleVerificationHarness::markOutputsModified(OutputTensors& ou detail::markOutputsModified(outputs, _deps.policy.useDevice()); } +void IntegrationBundleVerificationHarness::benchmarkIfVerified(GraphSession& session, + OutputTensors& engineOutputs, + const VerificationOutcome& outcome) +{ + if(!_deps.policy.benchmark.has_value() || outcome.status != OutcomeStatus::PASSED + || outcome.depth != VerificationDepth::VERIFIED) + { + return; + } + auto variantPack = buildVariantPack(engineOutputs, _deps.policy.useDevice()); + _benchmark = _deps.engineRunner->benchmark(session, variantPack, *_deps.policy.benchmark); +} + // ---- comparison ------------------------------------------------------------ VerificationOutcome diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp index 1dd22b075153..85ddd7e2ef74 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp @@ -234,6 +234,12 @@ class IntegrationBundleVerificationHarness : public ::testing::Test // Before the disposition, which returns: a pass must say what it was compared // against, or "auto" landing on a reference is indistinguishable from a skip. _deps.reporter->recordVerifier(_bundlePath.string(), outcome.verifier); + // A benchmark exists only for a verified pass. Its failure fails the test but + // leaves the claim verdict above alone: the claim is about correctness. + if(_benchmark.has_value()) + { + reportBenchmark(*_benchmark); + } reportOutcome(outcome); } @@ -481,6 +487,12 @@ class IntegrationBundleVerificationHarness : public ::testing::Test OutputTensors& refOutputs); void markOutputsModified(OutputTensors& outputs) const; + // Under --benchmark, times the engine on the outputs `outcome` just verified and + // keeps the result for TestBody() to report. Does nothing for any other outcome. + void benchmarkIfVerified(GraphSession& session, + OutputTensors& engineOutputs, + const VerificationOutcome& outcome); + // Golden data is loaded on the host, so under --validator auto it is compared there. VerificationOutcome compareAgainstGolden(OutputTensors& engineOutputs); VerificationOutcome compareOutputs(OutputTensors& engineOutputs, @@ -534,6 +546,8 @@ class IntegrationBundleVerificationHarness : public ::testing::Test // generated by prepareInputs() and owned here rather than by the shared bundle, so // they live exactly as long as the test. TensorMap _inputs; + // Set by benchmarkIfVerified(); reported at the end of TestBody(). + std::optional _benchmark; }; } // namespace hipdnn_integration_tests::bundle diff --git a/dnn-providers/integration-tests/src/harness/bundle/ProductionPolicy.cpp b/dnn-providers/integration-tests/src/harness/bundle/ProductionPolicy.cpp index 6b06cb64b423..73d8704d7c18 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/ProductionPolicy.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/ProductionPolicy.cpp @@ -24,6 +24,7 @@ HarnessPolicy productionPolicy(TensorPlacement placement) policy.arch = TestConfig::get().getCurrentArch(); policy.platform = currentPlatform(); policy.deviceVramMb = TestConfig::get().getCurrentDeviceVramMb(); + policy.benchmark = TestConfig::get().benchmarkSettings(); return policy; } diff --git a/dnn-providers/integration-tests/src/main.cpp b/dnn-providers/integration-tests/src/main.cpp index f93e324447ba..78eaf44b208b 100644 --- a/dnn-providers/integration-tests/src/main.cpp +++ b/dnn-providers/integration-tests/src/main.cpp @@ -23,6 +23,7 @@ #include "common/PlatformUtils.hpp" #include "common/Utilities.hpp" +#include "harness/Benchmark.hpp" #include "harness/SharedHandle.hpp" #include "harness/SupportMatrixCollector.hpp" #include "harness/TestConfig.hpp" @@ -190,6 +191,19 @@ int main(int argc, char** argv) noexcept "change = zero git diff. Run one at a time: concurrent " "--write-support-claims runs against the same bundle tree race " "on the sidecars and the last writer wins."); + parser.add_argument("--benchmark") + .default_value(false) + .implicit_value(true) + .help("After a test verifies the engine's outputs, time the same plan on the " + "same tensors and print min/median/mean device time per test. The " + "values are also written as benchmark_* properties in " + "--gtest_output files."); + parser.add_argument("--benchmark-iterations") + .scan<'i', int>() + .help("Timed executions per test (default 20). Requires --benchmark."); + parser.add_argument("--benchmark-warmup") + .scan<'i', int>() + .help("Untimed executions before timing (default 3). Requires --benchmark."); std::vector remainingArgs; try @@ -386,6 +400,20 @@ int main(int argc, char** argv) noexcept return 1; } + hipdnn_integration_tests::BenchmarkRequest benchmarkRequest; + benchmarkRequest.enabled = parser.get("--benchmark"); + benchmarkRequest.timedIterations = parser.present("--benchmark-iterations"); + benchmarkRequest.warmupIterations = parser.present("--benchmark-warmup"); + benchmarkRequest.skipGraphValidation = skipGraphValidation; + benchmarkRequest.captureBundles = captureDir.has_value(); + benchmarkRequest.writeSupportClaims = writeSupportClaims; + auto benchmark = hipdnn_integration_tests::resolveBenchmark(benchmarkRequest); + if(benchmark.error.has_value()) + { + std::cerr << "Error: " << *benchmark.error << "\n"; + return 1; + } + hipdnn_integration_tests::TestConfigOptions opts; opts.articlePath = std::move(articlePath); opts.engineName = std::move(engineName); @@ -399,6 +427,7 @@ int main(int argc, char** argv) noexcept opts.validatorDevice = validator; opts.captureDir = std::move(captureDir); opts.writeSupportClaims = writeSupportClaims; + opts.benchmark = benchmark.settings; opts.enforceSupportClaims = claimMode.mode == hipdnn_integration_tests::bundle::ClaimMode::ENFORCE; diff --git a/dnn-providers/integration-tests/tests/CMakeLists.txt b/dnn-providers/integration-tests/tests/CMakeLists.txt index 6d2ff941d266..03b31e0720dd 100644 --- a/dnn-providers/integration-tests/tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/tests/CMakeLists.txt @@ -63,6 +63,7 @@ add_executable(hipdnn_integration_tests_unit_tests TestGTestFilter.cpp TestVerificationPaths.cpp TestBundleVerificationHarness.cpp + TestBenchmark.cpp TestInputFillRecipes.cpp TestFillInputs.cpp TestToleranceResolver.cpp diff --git a/dnn-providers/integration-tests/tests/TestBenchmark.cpp b/dnn-providers/integration-tests/tests/TestBenchmark.cpp new file mode 100644 index 000000000000..970ffe29dc14 --- /dev/null +++ b/dnn-providers/integration-tests/tests/TestBenchmark.cpp @@ -0,0 +1,453 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include "BundleFixtureFiles.hpp" +#include "HarnessTestSupport.hpp" +#include "harness/Benchmark.hpp" + +// NOLINTBEGIN(readability-identifier-naming) + +using namespace hipdnn_integration_tests; +using namespace hipdnn_integration_tests::bundle; +using hipdnn_frontend::Error; +using hipdnn_frontend::ErrorCode; +using hipdnn_frontend::ExecutionTiming; +using hipdnn_frontend::TimingQuality; + +namespace +{ + +Error ok() +{ + return {ErrorCode::OK, ""}; +} + +Error broken(const std::string& message) +{ + return {ErrorCode::HIPDNN_BACKEND_ERROR, message}; +} + +ExecutionTiming measured(float ms, TimingQuality quality = TimingQuality::DEVICE_ONLY) +{ + ExecutionTiming timing; + timing.elapsedMs = ms; + timing.quality = quality; + return timing; +} + +ExecutionTiming watchdogTimeout() +{ + ExecutionTiming timing; + timing.timedOut = true; + return timing; +} + +/// A timer that replays `script` in order, then repeats its last entry. +struct ScriptedTimer +{ + std::vector script; + int calls = 0; + + Error operator()(ExecutionTiming& timing) + { + const auto index + = std::min(static_cast(calls), script.size() - 1); + timing = script[index]; + ++calls; + return ok(); + } +}; + +BenchmarkSettings settings(int timed, int warmup) +{ + BenchmarkSettings result; + result.timedIterations = timed; + result.warmupIterations = warmup; + return result; +} + +BenchmarkResult successfulResult(std::vector samples) +{ + BenchmarkResult result; + result.engine = "TEST_ENGINE"; + result.samplesMs = std::move(samples); + result.quality = TimingQuality::DEVICE_ONLY; + return result; +} + +} // namespace + +// ---- runBenchmarkLoop ------------------------------------------------------ + +TEST(TestBenchmarkLoop, RunsWarmupThenCollectsTheRequestedSamples) +{ + int warmups = 0; + ScriptedTimer timer{{measured(3.0f), measured(1.0f), measured(2.0f)}}; + + const auto result = runBenchmarkLoop( + settings(3, 2), + [&] { + ++warmups; + EXPECT_EQ(timer.calls, 0) << "warmup must finish before timing starts"; + return ok(); + }, + std::ref(timer)); + + ASSERT_TRUE(result.ok()) << result.error; + EXPECT_EQ(warmups, 2); + EXPECT_EQ(timer.calls, 3); + EXPECT_EQ(result.samplesMs, (std::vector{3.0f, 1.0f, 2.0f})); + EXPECT_EQ(result.quality, TimingQuality::DEVICE_ONLY); + EXPECT_EQ(result.discardedSamples, 0); +} + +TEST(TestBenchmarkLoop, ZeroWarmupStartsTimingImmediately) +{ + int warmups = 0; + ScriptedTimer timer{{measured(1.0f)}}; + + const auto result = runBenchmarkLoop( + settings(1, 0), + [&] { + ++warmups; + return ok(); + }, + std::ref(timer)); + + ASSERT_TRUE(result.ok()) << result.error; + EXPECT_EQ(warmups, 0); + EXPECT_EQ(timer.calls, 1); +} + +TEST(TestBenchmarkLoop, OneUnstalledSampleMarksTheWholeSetUnstalled) +{ + ScriptedTimer timer{{measured(1.0f), measured(1.5f, TimingQuality::UNSTALLED), measured(1.0f)}}; + + const auto result = runBenchmarkLoop(settings(3, 0), [] { return ok(); }, std::ref(timer)); + + ASSERT_TRUE(result.ok()) << result.error; + EXPECT_EQ(result.quality, TimingQuality::UNSTALLED); +} + +TEST(TestBenchmarkLoop, RunsWithoutAMeasurementAreRetriedAndCounted) +{ + ScriptedTimer timer{{watchdogTimeout(), measured(2.0f), ExecutionTiming{}, measured(4.0f)}}; + + const auto result = runBenchmarkLoop(settings(2, 0), [] { return ok(); }, std::ref(timer)); + + ASSERT_TRUE(result.ok()) << result.error; + EXPECT_EQ(timer.calls, 4); + EXPECT_EQ(result.samplesMs, (std::vector{2.0f, 4.0f})); + EXPECT_EQ(result.discardedSamples, 2); + EXPECT_EQ(result.quality, TimingQuality::DEVICE_ONLY); +} + +TEST(TestBenchmarkLoop, TooFewMeasurementsFailAfterABoundedNumberOfAttempts) +{ + ScriptedTimer timer{{measured(1.0f), watchdogTimeout()}}; + + const auto result = runBenchmarkLoop(settings(5, 0), [] { return ok(); }, std::ref(timer)); + + EXPECT_FALSE(result.ok()); + EXPECT_EQ(timer.calls, 10) << "retries are capped at timedIterations extra attempts"; + EXPECT_TRUE(result.samplesMs.empty()) << "a partial set must not be reported as a result"; + EXPECT_NE(result.error.find("9 of 10 timed executions returned no measurement"), + std::string::npos) + << result.error; + EXPECT_NE(result.error.find("collected 1 of 5 samples"), std::string::npos) << result.error; +} + +TEST(TestBenchmarkLoop, WarmupFailureStopsBeforeAnyTimedRun) +{ + ScriptedTimer timer{{measured(1.0f)}}; + + const auto result + = runBenchmarkLoop(settings(3, 2), [] { return broken("bad kernel"); }, std::ref(timer)); + + EXPECT_FALSE(result.ok()); + EXPECT_EQ(timer.calls, 0); + EXPECT_EQ(result.error, "warmup execution failed: bad kernel"); +} + +TEST(TestBenchmarkLoop, TimedFailureFailsAndDropsEarlierSamples) +{ + int calls = 0; + const auto result = runBenchmarkLoop( + settings(3, 0), + [] { return ok(); }, + [&](ExecutionTiming& timing) { + if(++calls == 2) + { + return broken("device lost"); + } + timing = measured(1.0f); + return ok(); + }); + + EXPECT_FALSE(result.ok()); + EXPECT_EQ(calls, 2); + EXPECT_TRUE(result.samplesMs.empty()); + EXPECT_EQ(result.error, "timed execution failed: device lost"); +} + +// ---- BenchmarkResult statistics ------------------------------------------- + +TEST(TestBenchmarkResult, StatisticsOfAnOddSampleCount) +{ + const auto result = successfulResult({5.0f, 1.0f, 3.0f}); + + EXPECT_FLOAT_EQ(result.minMs(), 1.0f); + EXPECT_FLOAT_EQ(result.medianMs(), 3.0f); + EXPECT_FLOAT_EQ(result.meanMs(), 3.0f); +} + +TEST(TestBenchmarkResult, MedianOfAnEvenSampleCountAveragesTheMiddlePair) +{ + const auto result = successfulResult({4.0f, 1.0f, 3.0f, 2.0f}); + + EXPECT_FLOAT_EQ(result.medianMs(), 2.5f); + EXPECT_FLOAT_EQ(result.meanMs(), 2.5f); +} + +TEST(TestBenchmarkResult, AResultWithoutSamplesIsNotOk) +{ + EXPECT_FALSE(BenchmarkResult{}.ok()); +} + +// ---- resolveBenchmark ----------------------------------------------------- + +TEST(TestResolveBenchmark, OffWithoutFlagsHasNoSettingsAndNoError) +{ + const auto resolution = resolveBenchmark({}); + + EXPECT_FALSE(resolution.settings.has_value()); + EXPECT_FALSE(resolution.error.has_value()); +} + +TEST(TestResolveBenchmark, EnabledAloneUsesTheDefaults) +{ + BenchmarkRequest request; + request.enabled = true; + + const auto resolution = resolveBenchmark(request); + + ASSERT_TRUE(resolution.settings.has_value()) << resolution.error.value_or(""); + EXPECT_EQ(resolution.settings->timedIterations, 20); + EXPECT_EQ(resolution.settings->warmupIterations, 3); +} + +TEST(TestResolveBenchmark, ExplicitCountsOverrideTheDefaults) +{ + BenchmarkRequest request; + request.enabled = true; + request.timedIterations = 7; + request.warmupIterations = 0; + + const auto resolution = resolveBenchmark(request); + + ASSERT_TRUE(resolution.settings.has_value()) << resolution.error.value_or(""); + EXPECT_EQ(resolution.settings->timedIterations, 7); + EXPECT_EQ(resolution.settings->warmupIterations, 0); +} + +TEST(TestResolveBenchmark, CountsWithoutBenchmarkAreRefused) +{ + BenchmarkRequest iterations; + iterations.timedIterations = 5; + BenchmarkRequest warmup; + warmup.warmupIterations = 1; + + for(const auto& request : {iterations, warmup}) + { + const auto resolution = resolveBenchmark(request); + EXPECT_FALSE(resolution.settings.has_value()); + ASSERT_TRUE(resolution.error.has_value()); + EXPECT_NE(resolution.error->find("require --benchmark"), std::string::npos); + } +} + +TEST(TestResolveBenchmark, OutOfRangeCountsAreRefused) +{ + BenchmarkRequest noIterations; + noIterations.enabled = true; + noIterations.timedIterations = 0; + BenchmarkRequest negativeWarmup; + negativeWarmup.enabled = true; + negativeWarmup.warmupIterations = -1; + + EXPECT_EQ(resolveBenchmark(noIterations).error, + std::optional("--benchmark-iterations must be at least 1")); + EXPECT_EQ(resolveBenchmark(negativeWarmup).error, + std::optional("--benchmark-warmup must not be negative")); +} + +TEST(TestResolveBenchmark, ModesThatDoNotExecuteAreRefused) +{ + BenchmarkRequest skip; + skip.skipGraphValidation = true; + BenchmarkRequest capture; + capture.captureBundles = true; + BenchmarkRequest write; + write.writeSupportClaims = true; + + for(auto request : {skip, capture, write}) + { + // Without --benchmark these modes are untouched. + EXPECT_FALSE(resolveBenchmark(request).error.has_value()); + + request.enabled = true; + const auto resolution = resolveBenchmark(request); + EXPECT_FALSE(resolution.settings.has_value()); + ASSERT_TRUE(resolution.error.has_value()); + EXPECT_NE(resolution.error->find("do not execute the graph"), std::string::npos); + } +} + +// ---- reportBenchmark ------------------------------------------------------ + +TEST(TestReportBenchmark, RecordsTheResultAsTestProperties) +{ + reportBenchmark(successfulResult({0.5f, 0.25f, 1.0f})); + + const auto* result = ::testing::UnitTest::GetInstance()->current_test_info()->result(); + std::map properties; + for(int i = 0; i < result->test_property_count(); ++i) + { + const auto& property = result->GetTestProperty(i); + properties[property.key()] = property.value(); + } + + EXPECT_EQ(properties["benchmark_engine"], "TEST_ENGINE"); + EXPECT_EQ(properties["benchmark_timing"], "DEVICE_ONLY"); + EXPECT_EQ(properties["benchmark_samples"], "3"); + EXPECT_EQ(properties["benchmark_min_ms"], "0.2500"); + EXPECT_EQ(properties["benchmark_median_ms"], "0.5000"); + EXPECT_EQ(properties["benchmark_mean_ms"], "0.5833"); +} + +TEST(TestReportBenchmark, AFailedBenchmarkFailsTheTest) +{ + BenchmarkResult failed; + failed.error = "timed execution failed: device lost"; + + EXPECT_NONFATAL_FAILURE(reportBenchmark(failed), + "benchmark failed: timed execution failed: device lost"); + EXPECT_NONFATAL_FAILURE(reportBenchmark(BenchmarkResult{}), + "benchmark failed: no samples were collected"); +} + +// ---- bundle harness ------------------------------------------------------- + +namespace +{ + +/// Expectations are checked after the harness returns: a gmock or EXPECT failure +/// raised inside driveHarness() lands in its intercepted results, not in this test. +class TestBundleHarnessBenchmark : public ::testing::Test +{ +protected: + std::optional _scopedDir; + testing_support::HarnessMocks _mocks; + + /// What the mocked runner returns from benchmark(). + BenchmarkResult _benchmarkResult = successfulResult({1.0f}); + /// The settings of every benchmark() call, and whether its variant pack held + /// the verified output tensor. + std::vector _benchmarkCalls; + bool _sawVerifiedOutput = false; + + void SetUp() override + { + testing_support::ensureTestConfigInitialized(); + _scopedDir.emplace(hipdnn_test_sdk::utilities::claimScratchDirectory("bundle_benchmark")); + + ON_CALL(_mocks.engineRunner, benchmark(::testing::_, ::testing::_, ::testing::_)) + .WillByDefault( + [this](GraphSession&, VariantPack& variantPack, const BenchmarkSettings& s) { + _benchmarkCalls.push_back(s); + _sawVerifiedOutput = variantPack.count(fixtures::K_OUTPUT_UID) == 1; + return _benchmarkResult; + }); + } + + /// Drives a golden bundle through the real harness. The mocked engine writes + /// `engineOutput`; the golden output is fixtures::K_OUTPUT_VALUE. + void run(float engineOutput, + std::optional benchmark, + ::testing::TestPartResultArray* results) + { + testing_support::engineWrites(_mocks.engineRunner, &fixtures::writeOutput, engineOutput); + + auto policy = testing_support::hostPolicy(VerificationMode::AUTO); + policy.benchmark = benchmark; + IntegrationBundleVerificationHarness harness(_mocks.dependencies(policy)); + harness.setBundle(fixtures::loadBundle(_scopedDir->path(), "bench", true), "bench"); + + testing_support::driveHarness(harness, results); + } +}; + +} // namespace + +TEST_F(TestBundleHarnessBenchmark, AVerifiedPassIsTimedWithThePolicySettings) +{ + ::testing::TestPartResultArray results; + run(fixtures::K_OUTPUT_VALUE, settings(4, 1), &results); + + EXPECT_FALSE(testing_support::anyFailed(results)) << testing_support::allMessages(results); + EXPECT_FALSE(testing_support::anySkipped(results)); + ASSERT_EQ(_benchmarkCalls.size(), 1u); + EXPECT_EQ(_benchmarkCalls[0].timedIterations, 4); + EXPECT_EQ(_benchmarkCalls[0].warmupIterations, 1); + EXPECT_TRUE(_sawVerifiedOutput) << "the benchmark must run on the verified tensors"; +} + +TEST_F(TestBundleHarnessBenchmark, AMismatchIsNotTimed) +{ + ::testing::TestPartResultArray results; + run(fixtures::K_OUTPUT_VALUE + 100.0f, settings(4, 1), &results); + + EXPECT_TRUE(testing_support::anyFailed(results)); + EXPECT_TRUE(_benchmarkCalls.empty()); +} + +TEST_F(TestBundleHarnessBenchmark, WithoutBenchmarkNothingIsTimed) +{ + ::testing::TestPartResultArray results; + run(fixtures::K_OUTPUT_VALUE, std::nullopt, &results); + + EXPECT_FALSE(testing_support::anyFailed(results)) << testing_support::allMessages(results); + EXPECT_TRUE(_benchmarkCalls.empty()); +} + +TEST_F(TestBundleHarnessBenchmark, ABenchmarkFailureFailsAVerifiedTest) +{ + _benchmarkResult = BenchmarkResult{}; + _benchmarkResult.error = "timed execution failed: device lost"; + + ::testing::TestPartResultArray results; + run(fixtures::K_OUTPUT_VALUE, settings(4, 1), &results); + + EXPECT_EQ(_benchmarkCalls.size(), 1u); + EXPECT_TRUE(testing_support::anyFailed(results)); + EXPECT_NE(testing_support::allMessages(results).find("benchmark failed: timed execution " + "failed: device lost"), + std::string::npos) + << testing_support::allMessages(results); +} + +// NOLINTEND(readability-identifier-naming) diff --git a/dnn-providers/integration-tests/tests/mocks/MockGraphEngineRunner.hpp b/dnn-providers/integration-tests/tests/mocks/MockGraphEngineRunner.hpp index 65d08a341cd6..73749080876d 100644 --- a/dnn-providers/integration-tests/tests/mocks/MockGraphEngineRunner.hpp +++ b/dnn-providers/integration-tests/tests/mocks/MockGraphEngineRunner.hpp @@ -37,6 +37,13 @@ class MockGraphEngineRunner : public IGraphEngineRunner const std::optional& engineUnderTest, VariantPack& variantPack), (override)); + + MOCK_METHOD(BenchmarkResult, + benchmark, + (GraphSession & session, + VariantPack& variantPack, + const BenchmarkSettings& settings), + (override)); }; } // namespace hipdnn_integration_tests::bundle From 4c34711cbac485d120557cf3e2cd650d001653ca Mon Sep 17 00:00:00 2001 From: Samuel Reeder Date: Mon, 19 Oct 2026 15:46:48 +0000 Subject: [PATCH 2/6] fix(hipdnn-integration-tests): address --benchmark review Rename three test cases that the test name validator rejects. Run the bundle harness benchmark tests over AUTO with golden data, AUTO with the reference oracle chain, and GOLDEN mode, so each benchmark hook has a unit test. Move the plan workspace allocation into allocatePlanWorkspace(), shared by execute() and the benchmark. Document which tests are not timed. --- .../integration-tests/docs/running-tests.md | 9 ++- .../src/harness/GraphBenchmark.hpp | 18 ++--- .../src/harness/PlanWorkspace.hpp | 36 +++++++++ .../bundle/FrontendGraphEngineRunner.cpp | 19 ++--- .../integration-tests/tests/TestBenchmark.cpp | 78 ++++++++++++++++--- 5 files changed, 122 insertions(+), 38 deletions(-) create mode 100644 dnn-providers/integration-tests/src/harness/PlanWorkspace.hpp diff --git a/dnn-providers/integration-tests/docs/running-tests.md b/dnn-providers/integration-tests/docs/running-tests.md index 281ac4f4a55e..f8fcf697a872 100644 --- a/dnn-providers/integration-tests/docs/running-tests.md +++ b/dnn-providers/integration-tests/docs/running-tests.md @@ -132,8 +132,13 @@ build/bin/hipdnn_integration_tests \ `--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 -runs the same plan on the same tensors again and times it. A test that skips or -fails is not timed. The C++ graph tests and the bundles both support it. +runs the same plan on the same tensors again and times it. 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 \ diff --git a/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp index 1eb9f1fd35e7..bda7ab548155 100644 --- a/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp +++ b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp @@ -3,15 +3,14 @@ #pragma once -#include #include #include #include -#include #include #include "harness/Benchmark.hpp" +#include "harness/PlanWorkspace.hpp" #include "harness/VariantPack.hpp" #include "harness/bundle/LoadedEngineTable.hpp" @@ -33,24 +32,17 @@ inline BenchmarkResult benchmarkGraph(hipdnn_frontend::graph::Graph& graph, return failed; } - int64_t workspaceSize = 0; - if(auto status = graph.get_workspace_size(workspaceSize); !status.is_good()) + auto workspace = allocatePlanWorkspace(graph, failed.error); + if(!workspace.has_value()) { - failed.error = status.get_message(); - return failed; - } - if(workspaceSize < 0) - { - failed.error = "engine reported a negative workspace size"; return failed; } - const hipdnn_data_sdk::utilities::Workspace workspace(static_cast(workspaceSize)); auto result = runBenchmarkLoop( settings, - [&] { return graph.execute(handle, variantPack, workspace.get()); }, + [&] { return graph.execute(handle, variantPack, workspace->get()); }, [&](hipdnn_frontend::ExecutionTiming& timing) { - return graph.execute_timed_ext(handle, variantPack, workspace.get(), timing); + return graph.execute_timed_ext(handle, variantPack, workspace->get(), timing); }); // Plugin engines are not in the static name registry, so ask the loaded table first. diff --git a/dnn-providers/integration-tests/src/harness/PlanWorkspace.hpp b/dnn-providers/integration-tests/src/harness/PlanWorkspace.hpp new file mode 100644 index 000000000000..884e2db12aff --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/PlanWorkspace.hpp @@ -0,0 +1,36 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include +#include + +#include +#include + +namespace hipdnn_integration_tests +{ + +/// Allocates the workspace that the built plan of `graph` needs. On failure, +/// returns std::nullopt and sets `error`. +inline std::optional> + allocatePlanWorkspace(hipdnn_frontend::graph::Graph& graph, std::string& error) +{ + int64_t workspaceSize = 0; + if(auto status = graph.get_workspace_size(workspaceSize); !status.is_good()) + { + error = status.get_message(); + return std::nullopt; + } + if(workspaceSize < 0) + { + error = "engine reported a negative workspace size"; + return std::nullopt; + } + return hipdnn_data_sdk::utilities::Workspace<>(static_cast(workspaceSize)); +} + +} // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp index ab39a421ace9..be74f2b29dae 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp @@ -4,13 +4,14 @@ #include "harness/bundle/FrontendGraphEngineRunner.hpp" #include +#include #include -#include #include #include "harness/EngineNotApplicableError.hpp" #include "harness/GraphBenchmark.hpp" +#include "harness/PlanWorkspace.hpp" #include "harness/SharedHandle.hpp" #include "harness/bundle/IntegrationTestBundle.hpp" @@ -122,19 +123,15 @@ EngineOpResult auto result = [&]() -> EngineOpResult { try { - int64_t workspaceSize = 0; - if(auto status = graph.get_workspace_size(workspaceSize); !status.is_good()) + std::string error; + auto workspace = allocatePlanWorkspace(graph, error); + if(!workspace.has_value()) { - return EngineOpResult::failed(status.get_message()); - } - if(workspaceSize < 0) - { - return EngineOpResult::failed("engine reported a negative workspace size"); + return EngineOpResult::failed(error); } - const hipdnn_data_sdk::utilities::Workspace workspace( - static_cast(workspaceSize)); - if(auto status = graph.execute(handle, variantPack, workspace.get()); !status.is_good()) + if(auto status = graph.execute(handle, variantPack, workspace->get()); + !status.is_good()) { return EngineOpResult::failed(status.get_message()); } diff --git a/dnn-providers/integration-tests/tests/TestBenchmark.cpp b/dnn-providers/integration-tests/tests/TestBenchmark.cpp index 970ffe29dc14..00d8b2a62d79 100644 --- a/dnn-providers/integration-tests/tests/TestBenchmark.cpp +++ b/dnn-providers/integration-tests/tests/TestBenchmark.cpp @@ -5,11 +5,13 @@ #include #include +#include #include #include #include #include #include +#include #include #include @@ -319,7 +321,7 @@ TEST(TestResolveBenchmark, ModesThatDoNotExecuteAreRefused) // ---- reportBenchmark ------------------------------------------------------ -TEST(TestReportBenchmark, RecordsTheResultAsTestProperties) +TEST(TestReportBenchmark, RecordsTheResultAsProperties) { reportBenchmark(successfulResult({0.5f, 0.25f, 1.0f})); @@ -339,7 +341,7 @@ TEST(TestReportBenchmark, RecordsTheResultAsTestProperties) EXPECT_EQ(properties["benchmark_mean_ms"], "0.5833"); } -TEST(TestReportBenchmark, AFailedBenchmarkFailsTheTest) +TEST(TestReportBenchmark, AFailedBenchmarkAddsAFailure) { BenchmarkResult failed; failed.error = "timed execution failed: device lost"; @@ -355,9 +357,30 @@ TEST(TestReportBenchmark, AFailedBenchmarkFailsTheTest) namespace { +/// One way the harness verifies a bundle. Each way has its own benchmark hook. +struct OracleCase +{ + const char* name; + VerificationMode mode; + bool hasGoldenOutput; + + friend void PrintTo(const OracleCase& oracleCase, std::ostream* os) + { + *os << oracleCase.name; + } +}; + +/// AUTO with golden data and explicit GOLDEN compare against the golden file. +/// AUTO without golden data falls through to the reference oracle chain. +const std::array ORACLE_CASES{{ + {"AutoGolden", VerificationMode::AUTO, true}, + {"AutoReference", VerificationMode::AUTO, false}, + {"GoldenMode", VerificationMode::GOLDEN, true}, +}}; + /// Expectations are checked after the harness returns: a gmock or EXPECT failure /// raised inside driveHarness() lands in its intercepted results, not in this test. -class TestBundleHarnessBenchmark : public ::testing::Test +class TestBundleHarnessBenchmark : public ::testing::TestWithParam { protected: std::optional _scopedDir; @@ -372,30 +395,44 @@ class TestBundleHarnessBenchmark : public ::testing::Test void SetUp() override { + using ::testing::_; testing_support::ensureTestConfigInitialized(); _scopedDir.emplace(hipdnn_test_sdk::utilities::claimScratchDirectory("bundle_benchmark")); - ON_CALL(_mocks.engineRunner, benchmark(::testing::_, ::testing::_, ::testing::_)) + ON_CALL(_mocks.engineRunner, benchmark(_, _, _)) .WillByDefault( [this](GraphSession&, VariantPack& variantPack, const BenchmarkSettings& s) { _benchmarkCalls.push_back(s); _sawVerifiedOutput = variantPack.count(fixtures::K_OUTPUT_UID) == 1; return _benchmarkResult; }); + + // Both references produce the golden value, so the reference path and the + // golden path agree on what a match is. + for(auto* reference : {&_mocks.cpuReference, &_mocks.gpuReference}) + { + ON_CALL(*reference, execute(_, _, _)) + .WillByDefault([](void*, size_t, const VariantPack& variantPack) { + auto* ptr = static_cast(variantPack.at(fixtures::K_OUTPUT_UID)); + std::fill(ptr, ptr + fixtures::K_OUTPUT_ELEMS, fixtures::K_OUTPUT_VALUE); + }); + } } - /// Drives a golden bundle through the real harness. The mocked engine writes - /// `engineOutput`; the golden output is fixtures::K_OUTPUT_VALUE. + /// Drives a bundle through the real harness with the parameter's oracle. The + /// mocked engine writes `engineOutput`; the expected output is + /// fixtures::K_OUTPUT_VALUE. void run(float engineOutput, std::optional benchmark, ::testing::TestPartResultArray* results) { testing_support::engineWrites(_mocks.engineRunner, &fixtures::writeOutput, engineOutput); - auto policy = testing_support::hostPolicy(VerificationMode::AUTO); + auto policy = testing_support::hostPolicy(GetParam().mode); policy.benchmark = benchmark; IntegrationBundleVerificationHarness harness(_mocks.dependencies(policy)); - harness.setBundle(fixtures::loadBundle(_scopedDir->path(), "bench", true), "bench"); + harness.setBundle( + fixtures::loadBundle(_scopedDir->path(), "bench", GetParam().hasGoldenOutput), "bench"); testing_support::driveHarness(harness, results); } @@ -403,7 +440,7 @@ class TestBundleHarnessBenchmark : public ::testing::Test } // namespace -TEST_F(TestBundleHarnessBenchmark, AVerifiedPassIsTimedWithThePolicySettings) +TEST_P(TestBundleHarnessBenchmark, AVerifiedPassIsTimedWithThePolicySettings) { ::testing::TestPartResultArray results; run(fixtures::K_OUTPUT_VALUE, settings(4, 1), &results); @@ -416,7 +453,7 @@ TEST_F(TestBundleHarnessBenchmark, AVerifiedPassIsTimedWithThePolicySettings) EXPECT_TRUE(_sawVerifiedOutput) << "the benchmark must run on the verified tensors"; } -TEST_F(TestBundleHarnessBenchmark, AMismatchIsNotTimed) +TEST_P(TestBundleHarnessBenchmark, AMismatchIsNotTimed) { ::testing::TestPartResultArray results; run(fixtures::K_OUTPUT_VALUE + 100.0f, settings(4, 1), &results); @@ -425,16 +462,26 @@ TEST_F(TestBundleHarnessBenchmark, AMismatchIsNotTimed) EXPECT_TRUE(_benchmarkCalls.empty()); } -TEST_F(TestBundleHarnessBenchmark, WithoutBenchmarkNothingIsTimed) +TEST_P(TestBundleHarnessBenchmark, WithoutBenchmarkAMatchPassesUntimed) { ::testing::TestPartResultArray results; run(fixtures::K_OUTPUT_VALUE, std::nullopt, &results); EXPECT_FALSE(testing_support::anyFailed(results)) << testing_support::allMessages(results); + EXPECT_FALSE(testing_support::anySkipped(results)); EXPECT_TRUE(_benchmarkCalls.empty()); } -TEST_F(TestBundleHarnessBenchmark, ABenchmarkFailureFailsAVerifiedTest) +TEST_P(TestBundleHarnessBenchmark, WithoutBenchmarkAMismatchStillFails) +{ + ::testing::TestPartResultArray results; + run(fixtures::K_OUTPUT_VALUE + 100.0f, std::nullopt, &results); + + EXPECT_TRUE(testing_support::anyFailed(results)); + EXPECT_TRUE(_benchmarkCalls.empty()); +} + +TEST_P(TestBundleHarnessBenchmark, ABenchmarkFailureFailsAVerifiedBundle) { _benchmarkResult = BenchmarkResult{}; _benchmarkResult.error = "timed execution failed: device lost"; @@ -450,4 +497,11 @@ TEST_F(TestBundleHarnessBenchmark, ABenchmarkFailureFailsAVerifiedTest) << testing_support::allMessages(results); } +INSTANTIATE_TEST_SUITE_P(Oracles, + TestBundleHarnessBenchmark, + ::testing::ValuesIn(ORACLE_CASES), + [](const ::testing::TestParamInfo& info) { + return std::string(info.param.name); + }); + // NOLINTEND(readability-identifier-naming) From cdfbcd991c53c3bc51d9500dde315fb913a0e188 Mon Sep 17 00:00:00 2001 From: Samuel Reeder Date: Tue, 20 Oct 2026 17:25:56 +0000 Subject: [PATCH 3/6] refactor(hipdnn-integration-tests): time --benchmark with Graph::autotune Replace the hand-written timing loop, statistics and engine-name lookup with Graph::autotune() limited to the engine under test (FIXED_AVERAGE, the --benchmark warmup and iteration counts). The benchmark now uses the same stall-gate restart, retry budget and statistics as autotune, and reports min, robust, mean and stddev from AutotuneResult. --- .../integration-tests/docs/running-tests.md | 34 +-- .../src/harness/Benchmark.hpp | 158 ++---------- .../src/harness/GraphBenchmark.hpp | 50 ++-- .../bundle/FrontendGraphEngineRunner.cpp | 9 +- .../bundle/FrontendGraphEngineRunner.hpp | 6 +- .../src/harness/bundle/IGraphEngineRunner.hpp | 10 +- .../IntegrationBundleVerificationHarness.hpp | 2 +- dnn-providers/integration-tests/src/main.cpp | 8 +- .../integration-tests/tests/TestBenchmark.cpp | 226 +++--------------- .../tests/mocks/MockGraphEngineRunner.hpp | 2 +- 10 files changed, 112 insertions(+), 393 deletions(-) diff --git a/dnn-providers/integration-tests/docs/running-tests.md b/dnn-providers/integration-tests/docs/running-tests.md index f8fcf697a872..8a32eb74bfeb 100644 --- a/dnn-providers/integration-tests/docs/running-tests.md +++ b/dnn-providers/integration-tests/docs/running-tests.md @@ -132,8 +132,8 @@ build/bin/hipdnn_integration_tests \ `--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 -runs the same plan on the same tensors again and times it. The C++ graph tests -and the bundles both support it. +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 @@ -152,24 +152,26 @@ build/bin/hipdnn_integration_tests \ Each timed test prints one line: ```text -[ BENCHMARK ] quick_ConvolutionFwd_Default.8_16_16_16_bfp16_nhwc_dil1x1_postpad0x0_prepad0x0: engine=MIOPEN_ENGINE timing=DEVICE_ONLY samples=20 min=0.0182 ms median=0.0184 ms mean=0.0185 ms +[ BENCHMARK ] quick_ConvolutionFwd_Default.8_16_16_16_bfp16_nhwc_dil1x1_postpad0x0_prepad0x0: engine=MIOPEN_ENGINE timing=DEVICE_ONLY samples=20 min=0.0182 ms robust=0.0184 ms mean=0.0185 ms stddev=0.0003 ms ``` -- The harness runs `--benchmark-warmup` untimed executions (default 3). Then - it collects `--benchmark-iterations` timed executions (default 20) with - `Graph::execute_timed_ext()`. -- `timing=DEVICE_ONLY` means that the stream was stalled for every sample, so - the times exclude host submission overhead. `timing=UNSTALLED` means that at - least one sample was not stalled and can include host time. Do not compare - the two kinds. -- A timed execution that returns no measurement (a stall watchdog timeout) is - retried. The line then shows `discarded=N`. The harness makes at most two - attempts for each requested sample. The test fails with - `benchmark failed: …` if it cannot collect all the samples in those attempts, - or if an execution fails. +- Autotune runs `--benchmark-warmup` untimed executions (default 3). 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. +- `robust` is the mean after autotune drops the slow tail: what the engine + usually costs. `min`, `mean` and `stddev` are over all 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:` or `xml:` writes the same values as the test properties `benchmark_engine`, `benchmark_timing`, `benchmark_samples`, - `benchmark_min_ms`, `benchmark_median_ms` and `benchmark_mean_ms`. + `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 diff --git a/dnn-providers/integration-tests/src/harness/Benchmark.hpp b/dnn-providers/integration-tests/src/harness/Benchmark.hpp index 129b9ea3e5ff..a44ff0bc3749 100644 --- a/dnn-providers/integration-tests/src/harness/Benchmark.hpp +++ b/dnn-providers/integration-tests/src/harness/Benchmark.hpp @@ -3,21 +3,16 @@ #pragma once -#include -#include -#include #include #include -#include #include #include #include -#include #include -#include #include +#include namespace hipdnn_integration_tests { @@ -87,56 +82,6 @@ inline BenchmarkResolution resolveBenchmark(const BenchmarkRequest& request) return {settings, std::nullopt}; } -/// The timings of one graph on one engine. -struct BenchmarkResult -{ - std::string engine; - std::vector samplesMs; - - /// DEVICE_ONLY only when every sample was. One UNSTALLED sample makes the set - /// UNSTALLED, because that sample can include host submission time. - hipdnn_frontend::TimingQuality quality = hipdnn_frontend::TimingQuality::INVALID; - - /// Timed runs that returned no measurement (stall watchdog timeout or a negative - /// reading). They are retried and are not in samplesMs. - int discardedSamples = 0; - - /// Why the benchmark failed. Empty on success. - std::string error; - - bool ok() const - { - return error.empty() && !samplesMs.empty(); - } - - float minMs() const - { - return samplesMs.empty() ? 0.0f : *std::min_element(samplesMs.begin(), samplesMs.end()); - } - - float medianMs() const - { - if(samplesMs.empty()) - { - return 0.0f; - } - auto sorted = samplesMs; - std::sort(sorted.begin(), sorted.end()); - const std::size_t mid = sorted.size() / 2; - return sorted.size() % 2 == 1 ? sorted[mid] : (sorted[mid - 1] + sorted[mid]) / 2.0f; - } - - float meanMs() const - { - if(samplesMs.empty()) - { - return 0.0f; - } - const double sum = std::accumulate(samplesMs.begin(), samplesMs.end(), 0.0); - return static_cast(sum / static_cast(samplesMs.size())); - } -}; - inline const char* toString(hipdnn_frontend::TimingQuality quality) { switch(quality) @@ -151,75 +96,16 @@ inline const char* toString(hipdnn_frontend::TimingQuality quality) } } -using BenchmarkRunOnce = std::function; -using BenchmarkTimeOnce = std::function; - -/// Runs `warmupIterations` untimed executions, then collects `timedIterations` valid -/// samples. A timed run with no measurement is retried, up to `timedIterations` extra -/// attempts. Any execution error stops the benchmark and is returned in `error`. -/// -/// Takes callables rather than a graph so the loop is testable without a device. -inline BenchmarkResult runBenchmarkLoop(const BenchmarkSettings& settings, - const BenchmarkRunOnce& runOnce, - const BenchmarkTimeOnce& timeOnce) -{ - BenchmarkResult result; - - for(int i = 0; i < settings.warmupIterations; ++i) - { - if(auto status = runOnce(); !status.is_good()) - { - result.error = "warmup execution failed: " + status.get_message(); - return result; - } - } - - bool anyUnstalled = false; - const int maxAttempts = 2 * settings.timedIterations; - for(int attempt = 0; attempt < maxAttempts - && static_cast(result.samplesMs.size()) < settings.timedIterations; - ++attempt) - { - hipdnn_frontend::ExecutionTiming timing; - if(auto status = timeOnce(timing); !status.is_good()) - { - result.error = "timed execution failed: " + status.get_message(); - result.samplesMs.clear(); - return result; - } - if(timing.quality == hipdnn_frontend::TimingQuality::INVALID || !timing.elapsedMs) - { - ++result.discardedSamples; - continue; - } - anyUnstalled = anyUnstalled || timing.quality == hipdnn_frontend::TimingQuality::UNSTALLED; - result.samplesMs.push_back(*timing.elapsedMs); - } - - if(static_cast(result.samplesMs.size()) < settings.timedIterations) - { - result.error = std::to_string(result.discardedSamples) + " of " - + std::to_string(maxAttempts) - + " timed executions returned no measurement; collected " - + std::to_string(result.samplesMs.size()) + " of " - + std::to_string(settings.timedIterations) + " samples"; - result.samplesMs.clear(); - return result; - } - - result.quality = anyUnstalled ? hipdnn_frontend::TimingQuality::UNSTALLED - : hipdnn_frontend::TimingQuality::DEVICE_ONLY; - return result; -} - -/// Publishes a benchmark for the running test: a "[ BENCHMARK ]" line on stdout and -/// benchmark_* properties in --gtest_output files. A failed benchmark fails the test. -inline void reportBenchmark(const BenchmarkResult& result) +/// 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.ok()) + if(!result.succeeded) { ADD_FAILURE() << "benchmark failed: " - << (result.error.empty() ? "no samples were collected" : result.error); + << (result.errorMessage.empty() ? "the engine was not benchmarked" + : result.errorMessage); return; } @@ -229,16 +115,18 @@ inline void reportBenchmark(const BenchmarkResult& result) return out.str(); }; - const auto minMs = ms(result.minMs()); - const auto medianMs = ms(result.medianMs()); - const auto meanMs = ms(result.meanMs()); + const auto minMs = ms(result.minTimeMs); + const auto robustMs = ms(result.robustTimeMs); + const auto meanMs = ms(result.avgTimeMs); + const auto stddevMs = ms(result.stddevMs); - ::testing::Test::RecordProperty("benchmark_engine", result.engine); - ::testing::Test::RecordProperty("benchmark_timing", toString(result.quality)); - ::testing::Test::RecordProperty("benchmark_samples", static_cast(result.samplesMs.size())); + ::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_median_ms", medianMs); + ::testing::Test::RecordProperty("benchmark_robust_ms", robustMs); ::testing::Test::RecordProperty("benchmark_mean_ms", meanMs); + ::testing::Test::RecordProperty("benchmark_stddev_ms", stddevMs); std::string testName = ""; if(const auto* info = ::testing::UnitTest::GetInstance()->current_test_info()) @@ -246,14 +134,10 @@ inline void reportBenchmark(const BenchmarkResult& result) testName = std::string(info->test_suite_name()) + "." + info->name(); } - std::cout << "[ BENCHMARK ] " << testName << ": engine=" << result.engine - << " timing=" << toString(result.quality) << " samples=" << result.samplesMs.size() - << " min=" << minMs << " ms median=" << medianMs << " ms mean=" << meanMs << " ms"; - if(result.discardedSamples > 0) - { - std::cout << " discarded=" << result.discardedSamples; - } - std::cout << '\n'; + std::cout << "[ BENCHMARK ] " << testName << ": engine=" << result.engineName + << " timing=" << toString(result.timingQuality) << " samples=" << result.iterationsRun + << " min=" << minMs << " ms robust=" << robustMs << " ms mean=" << meanMs + << " ms stddev=" << stddevMs << " ms\n"; } } // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp index bda7ab548155..5f9f8248460e 100644 --- a/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp +++ b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp @@ -5,56 +5,62 @@ #include #include +#include -#include #include +#include #include "harness/Benchmark.hpp" #include "harness/PlanWorkspace.hpp" #include "harness/VariantPack.hpp" -#include "harness/bundle/LoadedEngineTable.hpp" namespace hipdnn_integration_tests { -/// Times the plan `graph` already built, on the inputs and outputs in `variantPack`. -inline BenchmarkResult benchmarkGraph(hipdnn_frontend::graph::Graph& graph, - hipdnnHandle_t handle, - VariantPack& variantPack, - const BenchmarkSettings& settings) +/// Times the engine of the plan `graph` already built, on the tensors in +/// `variantPack`. Graph::autotune() does the timing, limited to that one engine, +/// so the warmup, the stall gate, the retries and the statistics are the same as +/// in autotune. +inline hipdnn_frontend::AutotuneResult benchmarkGraph(hipdnn_frontend::graph::Graph& graph, + hipdnnHandle_t handle, + VariantPack& variantPack, + const BenchmarkSettings& settings) { - BenchmarkResult failed; + hipdnn_frontend::AutotuneResult failed; int64_t engineId = -1; if(auto status = graph.get_execution_plan_engine_id(engineId); !status.is_good()) { - failed.error = status.get_message(); + failed.errorMessage = status.get_message(); return failed; } - auto workspace = allocatePlanWorkspace(graph, failed.error); + auto workspace = allocatePlanWorkspace(graph, failed.errorMessage); if(!workspace.has_value()) { return failed; } - auto result = runBenchmarkLoop( - settings, - [&] { return graph.execute(handle, variantPack, workspace->get()); }, - [&](hipdnn_frontend::ExecutionTiming& timing) { - return graph.execute_timed_ext(handle, variantPack, workspace->get(), timing); - }); + hipdnn_frontend::AutotuneConfig config; + config.strategy = hipdnn_frontend::AutotuneStrategy::FIXED_AVERAGE; + config.timedIterations = settings.timedIterations; + config.warmupIterations = settings.warmupIterations; + config.engineIdFilter = {engineId}; - // Plugin engines are not in the static name registry, so ask the loaded table first. - result.engine = hipdnn_data_sdk::utilities::engineNameOrHex(engineId); - for(const auto& engine : bundle::LoadedEngineTable::get().all()) + // autotune() returns an error when no engine succeeds. The engine's own result + // explains why, so look for it first. + std::vector results; + const auto status = graph.autotune(handle, variantPack, workspace->get(), config, {}, &results); + for(const auto& result : results) { - if(engine.id == engineId) + if(result.engineId == engineId && !result.excludedByCaller) { - result.engine = engine.name; + return result; } } - return result; + failed.errorMessage = status.is_good() ? "autotune returned no result for the engine under test" + : status.get_message(); + return failed; } } // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp index be74f2b29dae..ab66d2ced76a 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp @@ -148,14 +148,13 @@ EngineOpResult return result; } -BenchmarkResult FrontendGraphEngineRunner::benchmark(GraphSession& session, - VariantPack& variantPack, - const BenchmarkSettings& settings) +hipdnn_frontend::AutotuneResult FrontendGraphEngineRunner::benchmark( + GraphSession& session, VariantPack& variantPack, const BenchmarkSettings& settings) { if(session.graph == nullptr) { - BenchmarkResult failed; - failed.error = "openGraph() produced no graph to benchmark"; + hipdnn_frontend::AutotuneResult failed; + failed.errorMessage = "openGraph() produced no graph to benchmark"; return failed; } return benchmarkGraph(*session.graph, getSharedHandle(), variantPack, settings); diff --git a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.hpp b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.hpp index 3672820257dc..813ff2d2eee3 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.hpp @@ -37,9 +37,9 @@ class FrontendGraphEngineRunner : public IGraphEngineRunner const std::optional& engineUnderTest, VariantPack& variantPack) override; - BenchmarkResult benchmark(GraphSession& session, - VariantPack& variantPack, - const BenchmarkSettings& settings) override; + hipdnn_frontend::AutotuneResult benchmark(GraphSession& session, + VariantPack& variantPack, + const BenchmarkSettings& settings) override; }; } // namespace hipdnn_integration_tests::bundle diff --git a/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp b/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp index aa69bda7c260..9b2274845138 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp @@ -94,11 +94,11 @@ class IGraphEngineRunner VariantPack& variantPack) = 0; - /// Times the plan execute() built, on `variantPack`. Called only after the - /// outputs of execute() were verified. - virtual BenchmarkResult benchmark(GraphSession& session, - VariantPack& variantPack, - const BenchmarkSettings& settings) + /// Times the engine of the plan execute() built, on `variantPack`, with + /// Graph::autotune(). Called only after the outputs of execute() were verified. + virtual hipdnn_frontend::AutotuneResult benchmark(GraphSession& session, + VariantPack& variantPack, + const BenchmarkSettings& settings) = 0; }; diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp index 85ddd7e2ef74..994c6050f972 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp @@ -547,7 +547,7 @@ class IntegrationBundleVerificationHarness : public ::testing::Test // they live exactly as long as the test. TensorMap _inputs; // Set by benchmarkIfVerified(); reported at the end of TestBody(). - std::optional _benchmark; + std::optional _benchmark; }; } // namespace hipdnn_integration_tests::bundle diff --git a/dnn-providers/integration-tests/src/main.cpp b/dnn-providers/integration-tests/src/main.cpp index 78eaf44b208b..fe582cfeeeaf 100644 --- a/dnn-providers/integration-tests/src/main.cpp +++ b/dnn-providers/integration-tests/src/main.cpp @@ -194,10 +194,10 @@ int main(int argc, char** argv) noexcept parser.add_argument("--benchmark") .default_value(false) .implicit_value(true) - .help("After a test verifies the engine's outputs, time the same plan on the " - "same tensors and print min/median/mean device time per test. The " - "values are also written as benchmark_* properties in " - "--gtest_output files."); + .help("After a test verifies the engine's outputs, time the same engine on " + "the same tensors with Graph::autotune() and print min/robust/mean/" + "stddev device time per test. The values are also written as " + "benchmark_* properties in --gtest_output files."); parser.add_argument("--benchmark-iterations") .scan<'i', int>() .help("Timed executions per test (default 20). Requires --benchmark."); diff --git a/dnn-providers/integration-tests/tests/TestBenchmark.cpp b/dnn-providers/integration-tests/tests/TestBenchmark.cpp index 00d8b2a62d79..b61909402a03 100644 --- a/dnn-providers/integration-tests/tests/TestBenchmark.cpp +++ b/dnn-providers/integration-tests/tests/TestBenchmark.cpp @@ -7,7 +7,6 @@ #include #include #include -#include #include #include #include @@ -25,55 +24,12 @@ using namespace hipdnn_integration_tests; using namespace hipdnn_integration_tests::bundle; -using hipdnn_frontend::Error; -using hipdnn_frontend::ErrorCode; -using hipdnn_frontend::ExecutionTiming; +using hipdnn_frontend::AutotuneResult; using hipdnn_frontend::TimingQuality; namespace { -Error ok() -{ - return {ErrorCode::OK, ""}; -} - -Error broken(const std::string& message) -{ - return {ErrorCode::HIPDNN_BACKEND_ERROR, message}; -} - -ExecutionTiming measured(float ms, TimingQuality quality = TimingQuality::DEVICE_ONLY) -{ - ExecutionTiming timing; - timing.elapsedMs = ms; - timing.quality = quality; - return timing; -} - -ExecutionTiming watchdogTimeout() -{ - ExecutionTiming timing; - timing.timedOut = true; - return timing; -} - -/// A timer that replays `script` in order, then repeats its last entry. -struct ScriptedTimer -{ - std::vector script; - int calls = 0; - - Error operator()(ExecutionTiming& timing) - { - const auto index - = std::min(static_cast(calls), script.size() - 1); - timing = script[index]; - ++calls; - return ok(); - } -}; - BenchmarkSettings settings(int timed, int warmup) { BenchmarkSettings result; @@ -82,153 +38,28 @@ BenchmarkSettings settings(int timed, int warmup) return result; } -BenchmarkResult successfulResult(std::vector samples) -{ - BenchmarkResult result; - result.engine = "TEST_ENGINE"; - result.samplesMs = std::move(samples); - result.quality = TimingQuality::DEVICE_ONLY; +AutotuneResult successfulResult() +{ + AutotuneResult result; + result.engineName = "TEST_ENGINE"; + result.succeeded = true; + result.iterationsRun = 3; + result.minTimeMs = 0.25f; + result.robustTimeMs = 0.375f; + result.avgTimeMs = 0.5f; + result.stddevMs = 0.125f; + result.timingQuality = TimingQuality::DEVICE_ONLY; return result; } -} // namespace - -// ---- runBenchmarkLoop ------------------------------------------------------ - -TEST(TestBenchmarkLoop, RunsWarmupThenCollectsTheRequestedSamples) -{ - int warmups = 0; - ScriptedTimer timer{{measured(3.0f), measured(1.0f), measured(2.0f)}}; - - const auto result = runBenchmarkLoop( - settings(3, 2), - [&] { - ++warmups; - EXPECT_EQ(timer.calls, 0) << "warmup must finish before timing starts"; - return ok(); - }, - std::ref(timer)); - - ASSERT_TRUE(result.ok()) << result.error; - EXPECT_EQ(warmups, 2); - EXPECT_EQ(timer.calls, 3); - EXPECT_EQ(result.samplesMs, (std::vector{3.0f, 1.0f, 2.0f})); - EXPECT_EQ(result.quality, TimingQuality::DEVICE_ONLY); - EXPECT_EQ(result.discardedSamples, 0); -} - -TEST(TestBenchmarkLoop, ZeroWarmupStartsTimingImmediately) -{ - int warmups = 0; - ScriptedTimer timer{{measured(1.0f)}}; - - const auto result = runBenchmarkLoop( - settings(1, 0), - [&] { - ++warmups; - return ok(); - }, - std::ref(timer)); - - ASSERT_TRUE(result.ok()) << result.error; - EXPECT_EQ(warmups, 0); - EXPECT_EQ(timer.calls, 1); -} - -TEST(TestBenchmarkLoop, OneUnstalledSampleMarksTheWholeSetUnstalled) +AutotuneResult failedResult(std::string message) { - ScriptedTimer timer{{measured(1.0f), measured(1.5f, TimingQuality::UNSTALLED), measured(1.0f)}}; - - const auto result = runBenchmarkLoop(settings(3, 0), [] { return ok(); }, std::ref(timer)); - - ASSERT_TRUE(result.ok()) << result.error; - EXPECT_EQ(result.quality, TimingQuality::UNSTALLED); -} - -TEST(TestBenchmarkLoop, RunsWithoutAMeasurementAreRetriedAndCounted) -{ - ScriptedTimer timer{{watchdogTimeout(), measured(2.0f), ExecutionTiming{}, measured(4.0f)}}; - - const auto result = runBenchmarkLoop(settings(2, 0), [] { return ok(); }, std::ref(timer)); - - ASSERT_TRUE(result.ok()) << result.error; - EXPECT_EQ(timer.calls, 4); - EXPECT_EQ(result.samplesMs, (std::vector{2.0f, 4.0f})); - EXPECT_EQ(result.discardedSamples, 2); - EXPECT_EQ(result.quality, TimingQuality::DEVICE_ONLY); -} - -TEST(TestBenchmarkLoop, TooFewMeasurementsFailAfterABoundedNumberOfAttempts) -{ - ScriptedTimer timer{{measured(1.0f), watchdogTimeout()}}; - - const auto result = runBenchmarkLoop(settings(5, 0), [] { return ok(); }, std::ref(timer)); - - EXPECT_FALSE(result.ok()); - EXPECT_EQ(timer.calls, 10) << "retries are capped at timedIterations extra attempts"; - EXPECT_TRUE(result.samplesMs.empty()) << "a partial set must not be reported as a result"; - EXPECT_NE(result.error.find("9 of 10 timed executions returned no measurement"), - std::string::npos) - << result.error; - EXPECT_NE(result.error.find("collected 1 of 5 samples"), std::string::npos) << result.error; -} - -TEST(TestBenchmarkLoop, WarmupFailureStopsBeforeAnyTimedRun) -{ - ScriptedTimer timer{{measured(1.0f)}}; - - const auto result - = runBenchmarkLoop(settings(3, 2), [] { return broken("bad kernel"); }, std::ref(timer)); - - EXPECT_FALSE(result.ok()); - EXPECT_EQ(timer.calls, 0); - EXPECT_EQ(result.error, "warmup execution failed: bad kernel"); -} - -TEST(TestBenchmarkLoop, TimedFailureFailsAndDropsEarlierSamples) -{ - int calls = 0; - const auto result = runBenchmarkLoop( - settings(3, 0), - [] { return ok(); }, - [&](ExecutionTiming& timing) { - if(++calls == 2) - { - return broken("device lost"); - } - timing = measured(1.0f); - return ok(); - }); - - EXPECT_FALSE(result.ok()); - EXPECT_EQ(calls, 2); - EXPECT_TRUE(result.samplesMs.empty()); - EXPECT_EQ(result.error, "timed execution failed: device lost"); -} - -// ---- BenchmarkResult statistics ------------------------------------------- - -TEST(TestBenchmarkResult, StatisticsOfAnOddSampleCount) -{ - const auto result = successfulResult({5.0f, 1.0f, 3.0f}); - - EXPECT_FLOAT_EQ(result.minMs(), 1.0f); - EXPECT_FLOAT_EQ(result.medianMs(), 3.0f); - EXPECT_FLOAT_EQ(result.meanMs(), 3.0f); -} - -TEST(TestBenchmarkResult, MedianOfAnEvenSampleCountAveragesTheMiddlePair) -{ - const auto result = successfulResult({4.0f, 1.0f, 3.0f, 2.0f}); - - EXPECT_FLOAT_EQ(result.medianMs(), 2.5f); - EXPECT_FLOAT_EQ(result.meanMs(), 2.5f); + AutotuneResult result; + result.errorMessage = std::move(message); + return result; } -TEST(TestBenchmarkResult, AResultWithoutSamplesIsNotOk) -{ - EXPECT_FALSE(BenchmarkResult{}.ok()); -} +} // namespace // ---- resolveBenchmark ----------------------------------------------------- @@ -323,7 +154,7 @@ TEST(TestResolveBenchmark, ModesThatDoNotExecuteAreRefused) TEST(TestReportBenchmark, RecordsTheResultAsProperties) { - reportBenchmark(successfulResult({0.5f, 0.25f, 1.0f})); + reportBenchmark(successfulResult()); const auto* result = ::testing::UnitTest::GetInstance()->current_test_info()->result(); std::map properties; @@ -337,19 +168,17 @@ TEST(TestReportBenchmark, RecordsTheResultAsProperties) EXPECT_EQ(properties["benchmark_timing"], "DEVICE_ONLY"); EXPECT_EQ(properties["benchmark_samples"], "3"); EXPECT_EQ(properties["benchmark_min_ms"], "0.2500"); - EXPECT_EQ(properties["benchmark_median_ms"], "0.5000"); - EXPECT_EQ(properties["benchmark_mean_ms"], "0.5833"); + EXPECT_EQ(properties["benchmark_robust_ms"], "0.3750"); + EXPECT_EQ(properties["benchmark_mean_ms"], "0.5000"); + EXPECT_EQ(properties["benchmark_stddev_ms"], "0.1250"); } TEST(TestReportBenchmark, AFailedBenchmarkAddsAFailure) { - BenchmarkResult failed; - failed.error = "timed execution failed: device lost"; - - EXPECT_NONFATAL_FAILURE(reportBenchmark(failed), - "benchmark failed: timed execution failed: device lost"); - EXPECT_NONFATAL_FAILURE(reportBenchmark(BenchmarkResult{}), - "benchmark failed: no samples were collected"); + EXPECT_NONFATAL_FAILURE(reportBenchmark(failedResult("Benchmark failed on iteration 2")), + "benchmark failed: Benchmark failed on iteration 2"); + EXPECT_NONFATAL_FAILURE(reportBenchmark(AutotuneResult{}), + "benchmark failed: the engine was not benchmarked"); } // ---- bundle harness ------------------------------------------------------- @@ -387,7 +216,7 @@ class TestBundleHarnessBenchmark : public ::testing::TestWithParam testing_support::HarnessMocks _mocks; /// What the mocked runner returns from benchmark(). - BenchmarkResult _benchmarkResult = successfulResult({1.0f}); + AutotuneResult _benchmarkResult = successfulResult(); /// The settings of every benchmark() call, and whether its variant pack held /// the verified output tensor. std::vector _benchmarkCalls; @@ -483,8 +312,7 @@ TEST_P(TestBundleHarnessBenchmark, WithoutBenchmarkAMismatchStillFails) TEST_P(TestBundleHarnessBenchmark, ABenchmarkFailureFailsAVerifiedBundle) { - _benchmarkResult = BenchmarkResult{}; - _benchmarkResult.error = "timed execution failed: device lost"; + _benchmarkResult = failedResult("timed execution failed: device lost"); ::testing::TestPartResultArray results; run(fixtures::K_OUTPUT_VALUE, settings(4, 1), &results); diff --git a/dnn-providers/integration-tests/tests/mocks/MockGraphEngineRunner.hpp b/dnn-providers/integration-tests/tests/mocks/MockGraphEngineRunner.hpp index 73749080876d..450ecad5d9df 100644 --- a/dnn-providers/integration-tests/tests/mocks/MockGraphEngineRunner.hpp +++ b/dnn-providers/integration-tests/tests/mocks/MockGraphEngineRunner.hpp @@ -38,7 +38,7 @@ class MockGraphEngineRunner : public IGraphEngineRunner VariantPack& variantPack), (override)); - MOCK_METHOD(BenchmarkResult, + MOCK_METHOD(hipdnn_frontend::AutotuneResult, benchmark, (GraphSession & session, VariantPack& variantPack, From 5e480eee7d013a3e3e26316dca8e0c5b9a62e4a2 Mon Sep 17 00:00:00 2001 From: Samuel Reeder Date: Tue, 20 Oct 2026 17:57:44 +0000 Subject: [PATCH 4/6] fix(hipdnn-integration-tests): default --benchmark-warmup to 30 Each test starts with an idle GPU. With 1-10 warmup executions the MI210 quick-tier bundles read 4-60% slow against a 200-iteration reference; with 30 they settle within about 1-2%. More timed iterations do not remove the bias. --- dnn-providers/integration-tests/docs/running-tests.md | 9 +++++++-- .../integration-tests/src/harness/Benchmark.hpp | 5 ++++- dnn-providers/integration-tests/src/main.cpp | 2 +- dnn-providers/integration-tests/tests/TestBenchmark.cpp | 2 +- 4 files changed, 13 insertions(+), 5 deletions(-) diff --git a/dnn-providers/integration-tests/docs/running-tests.md b/dnn-providers/integration-tests/docs/running-tests.md index 8a32eb74bfeb..df19a43caa74 100644 --- a/dnn-providers/integration-tests/docs/running-tests.md +++ b/dnn-providers/integration-tests/docs/running-tests.md @@ -155,10 +155,15 @@ Each timed test prints one line: [ BENCHMARK ] quick_ConvolutionFwd_Default.8_16_16_16_bfp16_nhwc_dil1x1_postpad0x0_prepad0x0: engine=MIOPEN_ENGINE timing=DEVICE_ONLY samples=20 min=0.0182 ms robust=0.0184 ms mean=0.0185 ms stddev=0.0003 ms ``` -- Autotune runs `--benchmark-warmup` untimed executions (default 3). Then it +- 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. - `timing=DEVICE_ONLY` means that the stream was stalled, so the times exclude @@ -216,7 +221,7 @@ Otherwise, you test the old tree. | `--capture-bundles ` | | 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` | | 3 | Untimed executions before timing. Requires `--benchmark`. | +| `--benchmark-warmup` | | 30 | Untimed executions before timing. Requires `--benchmark`. | | `--gtest_*` | | | Passed through to GTest. | The binary refuses these flag combinations (exit 1): diff --git a/dnn-providers/integration-tests/src/harness/Benchmark.hpp b/dnn-providers/integration-tests/src/harness/Benchmark.hpp index a44ff0bc3749..4acb12a44f14 100644 --- a/dnn-providers/integration-tests/src/harness/Benchmark.hpp +++ b/dnn-providers/integration-tests/src/harness/Benchmark.hpp @@ -21,7 +21,10 @@ namespace hipdnn_integration_tests struct BenchmarkSettings { static constexpr int DEFAULT_TIMED_ITERATIONS = 20; - static constexpr int DEFAULT_WARMUP_ITERATIONS = 3; + // 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; diff --git a/dnn-providers/integration-tests/src/main.cpp b/dnn-providers/integration-tests/src/main.cpp index fe582cfeeeaf..1065696ef74a 100644 --- a/dnn-providers/integration-tests/src/main.cpp +++ b/dnn-providers/integration-tests/src/main.cpp @@ -203,7 +203,7 @@ int main(int argc, char** argv) noexcept .help("Timed executions per test (default 20). Requires --benchmark."); parser.add_argument("--benchmark-warmup") .scan<'i', int>() - .help("Untimed executions before timing (default 3). Requires --benchmark."); + .help("Untimed executions before timing (default 30). Requires --benchmark."); std::vector remainingArgs; try diff --git a/dnn-providers/integration-tests/tests/TestBenchmark.cpp b/dnn-providers/integration-tests/tests/TestBenchmark.cpp index b61909402a03..657eae458867 100644 --- a/dnn-providers/integration-tests/tests/TestBenchmark.cpp +++ b/dnn-providers/integration-tests/tests/TestBenchmark.cpp @@ -80,7 +80,7 @@ TEST(TestResolveBenchmark, EnabledAloneUsesTheDefaults) ASSERT_TRUE(resolution.settings.has_value()) << resolution.error.value_or(""); EXPECT_EQ(resolution.settings->timedIterations, 20); - EXPECT_EQ(resolution.settings->warmupIterations, 3); + EXPECT_EQ(resolution.settings->warmupIterations, 30); } TEST(TestResolveBenchmark, ExplicitCountsOverrideTheDefaults) From 6d118adc7fad3d8dcdcd1f73bb034df411c85792 Mon Sep 17 00:00:00 2001 From: Samuel Reeder Date: Tue, 20 Oct 2026 18:41:17 +0000 Subject: [PATCH 5/6] feat(hipdnn-integration-tests): report whether --benchmark timed a tuned solution Add tuned=0|1 to the [ BENCHMARK ] line and a benchmark_tuned gtest property. tuned=1 only when HIPDNN_FORCE_BENCHMARKING is on and the engine exposes the global.benchmarking knob; the harness sets no knobs, so the default is the engine's heuristic choice. --- .../integration-tests/docs/running-tests.md | 16 ++++++-- .../src/harness/Benchmark.hpp | 20 +++++++-- .../integration-tests/tests/TestBenchmark.cpp | 41 ++++++++++++++++++- 3 files changed, 69 insertions(+), 8 deletions(-) diff --git a/dnn-providers/integration-tests/docs/running-tests.md b/dnn-providers/integration-tests/docs/running-tests.md index df19a43caa74..e066428ddb43 100644 --- a/dnn-providers/integration-tests/docs/running-tests.md +++ b/dnn-providers/integration-tests/docs/running-tests.md @@ -152,7 +152,7 @@ build/bin/hipdnn_integration_tests \ Each timed test prints one line: ```text -[ BENCHMARK ] quick_ConvolutionFwd_Default.8_16_16_16_bfp16_nhwc_dil1x1_postpad0x0_prepad0x0: engine=MIOPEN_ENGINE timing=DEVICE_ONLY samples=20 min=0.0182 ms robust=0.0184 ms mean=0.0185 ms stddev=0.0003 ms +[ 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.0182 ms robust=0.0184 ms mean=0.0185 ms stddev=0.0003 ms ``` - Autotune runs `--benchmark-warmup` untimed executions (default 30). Then it @@ -166,6 +166,14 @@ Each timed test prints one line: 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 @@ -174,9 +182,9 @@ Each timed test prints one line: engine, for example when an execution fails or too many timed executions return no measurement. - `--gtest_output=json:` or `xml:` writes the same values as the - test properties `benchmark_engine`, `benchmark_timing`, `benchmark_samples`, - `benchmark_min_ms`, `benchmark_robust_ms`, `benchmark_mean_ms` and - `benchmark_stddev_ms`. + 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 diff --git a/dnn-providers/integration-tests/src/harness/Benchmark.hpp b/dnn-providers/integration-tests/src/harness/Benchmark.hpp index 4acb12a44f14..f0156d7eca8a 100644 --- a/dnn-providers/integration-tests/src/harness/Benchmark.hpp +++ b/dnn-providers/integration-tests/src/harness/Benchmark.hpp @@ -13,6 +13,7 @@ #include #include +#include namespace hipdnn_integration_tests { @@ -99,6 +100,16 @@ inline const char* toString(hipdnn_frontend::TimingQuality quality) } } +/// 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. @@ -122,6 +133,7 @@ inline void reportBenchmark(const hipdnn_frontend::AutotuneResult& result) 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)); @@ -130,6 +142,7 @@ inline void reportBenchmark(const hipdnn_frontend::AutotuneResult& result) ::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); std::string testName = ""; if(const auto* info = ::testing::UnitTest::GetInstance()->current_test_info()) @@ -138,9 +151,10 @@ inline void reportBenchmark(const hipdnn_frontend::AutotuneResult& result) } std::cout << "[ BENCHMARK ] " << testName << ": engine=" << result.engineName - << " timing=" << toString(result.timingQuality) << " samples=" << result.iterationsRun - << " min=" << minMs << " ms robust=" << robustMs << " ms mean=" << meanMs - << " ms stddev=" << stddevMs << " ms\n"; + << " 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 diff --git a/dnn-providers/integration-tests/tests/TestBenchmark.cpp b/dnn-providers/integration-tests/tests/TestBenchmark.cpp index 657eae458867..ca7733234c90 100644 --- a/dnn-providers/integration-tests/tests/TestBenchmark.cpp +++ b/dnn-providers/integration-tests/tests/TestBenchmark.cpp @@ -14,6 +14,8 @@ #include #include +#include +#include #include #include "BundleFixtureFiles.hpp" @@ -154,7 +156,12 @@ TEST(TestResolveBenchmark, ModesThatDoNotExecuteAreRefused) TEST(TestReportBenchmark, RecordsTheResultAsProperties) { - reportBenchmark(successfulResult()); + const hipdnn_test_sdk::utilities::ScopedEnvironmentVariableSetter forceBenchmarking( + hipdnn_plugin_sdk::FORCE_BENCHMARKING_ENV_NAME, "1"); + auto benchmarked = successfulResult(); + benchmarked.supportsExhaustive = true; + + reportBenchmark(benchmarked); const auto* result = ::testing::UnitTest::GetInstance()->current_test_info()->result(); std::map properties; @@ -171,6 +178,38 @@ TEST(TestReportBenchmark, RecordsTheResultAsProperties) EXPECT_EQ(properties["benchmark_robust_ms"], "0.3750"); EXPECT_EQ(properties["benchmark_mean_ms"], "0.5000"); EXPECT_EQ(properties["benchmark_stddev_ms"], "0.1250"); + EXPECT_EQ(properties["benchmark_tuned"], "1"); +} + +// The harness sets no knobs, so only HIPDNN_FORCE_BENCHMARKING can turn benchmarking +// on, and only an engine with the global.benchmarking knob reads it. +TEST(TestReportBenchmark, TunedOnlyWhenForcedOnAnEngineWithTheKnob) +{ + struct Case + { + bool engineHasKnob; + const char* forceBenchmarking; + bool tuned; + }; + const std::array cases{{ + {true, "", false}, + {true, "1", true}, + {true, "0", false}, + {false, "1", false}, + {false, "", false}, + }}; + + for(const auto& c : cases) + { + const hipdnn_test_sdk::utilities::ScopedEnvironmentVariableSetter forceBenchmarking( + hipdnn_plugin_sdk::FORCE_BENCHMARKING_ENV_NAME, c.forceBenchmarking); + auto result = successfulResult(); + result.supportsExhaustive = c.engineHasKnob; + + EXPECT_EQ(selectedByBenchmarking(result), c.tuned) + << "knob=" << c.engineHasKnob << " HIPDNN_FORCE_BENCHMARKING='" << c.forceBenchmarking + << "'"; + } } TEST(TestReportBenchmark, AFailedBenchmarkAddsAFailure) From 2929e22ada64e19e84f3e9d2b03c30f4cbe19cbd Mon Sep 17 00:00:00 2001 From: Samuel Reeder Date: Thu, 22 Oct 2026 17:11:34 +0000 Subject: [PATCH 6/6] fix(hipdnn-integration-tests): address --benchmark review comments - Catch exceptions from the benchmark in benchmarkIfVerified() so a throw fails the test without replacing the verified outcome. - Move the result pick into selectEngineResult() and unit-test it. - Move reportBenchmark() and toString(TimingQuality) to BenchmarkReport.hpp and currentTestName() to CurrentTestName.hpp; reuse currentTestName() in the report. - toString(TimingQuality) throws on an unknown value. - executeGpuGraph() uses allocatePlanWorkspace(). - Make the docs example match its sample output. --- .../integration-tests/docs/running-tests.md | 4 +- .../src/harness/Benchmark.hpp | 80 ----------- .../src/harness/BenchmarkReport.hpp | 124 ++++++++++++++++++ .../src/harness/CurrentTestName.hpp | 24 ++++ .../src/harness/GraphBenchmark.hpp | 14 +- .../IntegrationGraphVerificationHarness.hpp | 12 +- .../src/harness/TomlGuards.hpp | 11 +- .../src/harness/bundle/IGraphEngineRunner.hpp | 2 + .../IntegrationBundleVerificationHarness.cpp | 15 ++- .../IntegrationBundleVerificationHarness.hpp | 1 + .../integration-tests/tests/TestBenchmark.cpp | 116 +++++++++++++++- 11 files changed, 287 insertions(+), 116 deletions(-) create mode 100644 dnn-providers/integration-tests/src/harness/BenchmarkReport.hpp create mode 100644 dnn-providers/integration-tests/src/harness/CurrentTestName.hpp diff --git a/dnn-providers/integration-tests/docs/running-tests.md b/dnn-providers/integration-tests/docs/running-tests.md index e066428ddb43..1c9907b87ffd 100644 --- a/dnn-providers/integration-tests/docs/running-tests.md +++ b/dnn-providers/integration-tests/docs/running-tests.md @@ -145,14 +145,14 @@ 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 --benchmark-iterations 50 \ + --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.0182 ms robust=0.0184 ms mean=0.0185 ms stddev=0.0003 ms +[ 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 diff --git a/dnn-providers/integration-tests/src/harness/Benchmark.hpp b/dnn-providers/integration-tests/src/harness/Benchmark.hpp index f0156d7eca8a..e752d8aa1d61 100644 --- a/dnn-providers/integration-tests/src/harness/Benchmark.hpp +++ b/dnn-providers/integration-tests/src/harness/Benchmark.hpp @@ -3,18 +3,9 @@ #pragma once -#include -#include #include -#include #include -#include - -#include -#include -#include - namespace hipdnn_integration_tests { @@ -86,75 +77,4 @@ inline BenchmarkResolution resolveBenchmark(const BenchmarkRequest& request) return {settings, std::nullopt}; } -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: - default: - return "INVALID"; - } -} - -/// 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); - - std::string testName = ""; - if(const auto* info = ::testing::UnitTest::GetInstance()->current_test_info()) - { - testName = std::string(info->test_suite_name()) + "." + info->name(); - } - - 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 diff --git a/dnn-providers/integration-tests/src/harness/BenchmarkReport.hpp b/dnn-providers/integration-tests/src/harness/BenchmarkReport.hpp new file mode 100644 index 000000000000..b94cdd2021b1 --- /dev/null +++ b/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 +#include +#include +#include +#include +#include +#include + +#include + +#include +#include +#include +#include + +#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(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& 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 = ""; + } + + 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 diff --git a/dnn-providers/integration-tests/src/harness/CurrentTestName.hpp b/dnn-providers/integration-tests/src/harness/CurrentTestName.hpp new file mode 100644 index 000000000000..5796ae3c3221 --- /dev/null +++ b/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 + +#include + +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 diff --git a/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp index 5f9f8248460e..10faa34ff9f1 100644 --- a/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp +++ b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp @@ -11,6 +11,7 @@ #include #include "harness/Benchmark.hpp" +#include "harness/BenchmarkReport.hpp" #include "harness/PlanWorkspace.hpp" #include "harness/VariantPack.hpp" @@ -47,20 +48,9 @@ inline hipdnn_frontend::AutotuneResult benchmarkGraph(hipdnn_frontend::graph::Gr config.warmupIterations = settings.warmupIterations; config.engineIdFilter = {engineId}; - // autotune() returns an error when no engine succeeds. The engine's own result - // explains why, so look for it first. std::vector results; const auto status = graph.autotune(handle, variantPack, workspace->get(), config, {}, &results); - for(const auto& result : results) - { - if(result.engineId == engineId && !result.excludedByCaller) - { - return result; - } - } - failed.errorMessage = status.is_good() ? "autotune returned no result for the engine under test" - : status.get_message(); - return failed; + return selectEngineResult(engineId, results, status); } } // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp index 0d88eae690c4..5518781b7f60 100644 --- a/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp +++ b/dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp @@ -12,7 +12,6 @@ #include #include -#include #include #include #include @@ -33,6 +32,7 @@ #include "harness/GraphBenchmark.hpp" #include "harness/GraphDescription.hpp" #include "harness/IReferenceGraphExecutor.hpp" +#include "harness/PlanWorkspace.hpp" #include "harness/ReferenceGraphExecutorFactory.hpp" #include "harness/SharedHandle.hpp" #include "harness/SupportMatrixCollector.hpp" @@ -557,14 +557,12 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam(workspaceSize)); + std::string error; + auto workspace = allocatePlanWorkspace(graph, error); + ASSERT_TRUE(workspace.has_value()) << error; auto variantPack = bundle.toDeviceVariantPack(); - result = graph.execute(handle, variantPack, workspace.get()); + auto result = graph.execute(handle, variantPack, workspace->get()); ASSERT_EQ(result.code, hipdnn_frontend::ErrorCode::OK) << result.err_msg; } diff --git a/dnn-providers/integration-tests/src/harness/TomlGuards.hpp b/dnn-providers/integration-tests/src/harness/TomlGuards.hpp index be4313f0bb6c..4810a8296ed9 100644 --- a/dnn-providers/integration-tests/src/harness/TomlGuards.hpp +++ b/dnn-providers/integration-tests/src/harness/TomlGuards.hpp @@ -11,22 +11,13 @@ #include +#include "harness/CurrentTestName.hpp" #include "harness/TestConfig.hpp" #include "harness/bundle/OutputComparison.hpp" namespace hipdnn_integration_tests { -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(); -} - inline std::optional findTomlToleranceOverride(const std::string& testName) { if(testName.empty()) diff --git a/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp b/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp index 9b2274845138..772d73a605ac 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IGraphEngineRunner.hpp @@ -9,6 +9,8 @@ #include #include +#include + #include "harness/Benchmark.hpp" #include "harness/VariantPack.hpp" #include "harness/bundle/GraphSession.hpp" diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp index 8b4ec8dd5d22..b38b41e43671 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp @@ -880,8 +880,19 @@ void IntegrationBundleVerificationHarness::benchmarkIfVerified(GraphSession& ses { return; } - auto variantPack = buildVariantPack(engineOutputs, _deps.policy.useDevice()); - _benchmark = _deps.engineRunner->benchmark(session, variantPack, *_deps.policy.benchmark); + // A throw here must not reach TestBody()'s catch: it would replace the verified + // PASS with a harness failure and change the support-claim verdict. + try + { + auto variantPack = buildVariantPack(engineOutputs, _deps.policy.useDevice()); + _benchmark = _deps.engineRunner->benchmark(session, variantPack, *_deps.policy.benchmark); + } + catch(const std::exception& e) + { + hipdnn_frontend::AutotuneResult failed; + failed.errorMessage = e.what(); + _benchmark = failed; + } } // ---- comparison ------------------------------------------------------------ diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp index 994c6050f972..1366c6368aaf 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp @@ -24,6 +24,7 @@ #include #include +#include "harness/BenchmarkReport.hpp" #include "harness/BundleMetadata.hpp" #include "harness/TestConfig.hpp" #include "harness/TomlGuards.hpp" diff --git a/dnn-providers/integration-tests/tests/TestBenchmark.cpp b/dnn-providers/integration-tests/tests/TestBenchmark.cpp index ca7733234c90..78b4e3505c3c 100644 --- a/dnn-providers/integration-tests/tests/TestBenchmark.cpp +++ b/dnn-providers/integration-tests/tests/TestBenchmark.cpp @@ -11,6 +11,7 @@ #include #include #include +#include #include #include @@ -21,6 +22,7 @@ #include "BundleFixtureFiles.hpp" #include "HarnessTestSupport.hpp" #include "harness/Benchmark.hpp" +#include "harness/BenchmarkReport.hpp" // NOLINTBEGIN(readability-identifier-naming) @@ -152,6 +154,85 @@ TEST(TestResolveBenchmark, ModesThatDoNotExecuteAreRefused) } } +// ---- selectEngineResult --------------------------------------------------- + +namespace +{ + +AutotuneResult resultFor(int64_t engineId, bool excludedByCaller = false) +{ + auto result = successfulResult(); + result.engineId = engineId; + result.excludedByCaller = excludedByCaller; + return result; +} + +const hipdnn_frontend::Error K_GOOD{hipdnn_frontend::ErrorCode::OK, ""}; +const hipdnn_frontend::Error K_NO_ENGINE_SUCCEEDED{hipdnn_frontend::ErrorCode::HIPDNN_BACKEND_ERROR, + "no engine succeeded"}; + +} // namespace + +TEST(TestSelectEngineResult, ReturnsTheResultOfTheRequestedEngine) +{ + auto other = resultFor(7); + other.minTimeMs = 99.0f; + auto wanted = resultFor(42); + wanted.minTimeMs = 1.0f; + + const auto selected = selectEngineResult(42, {other, wanted}, K_GOOD); + + EXPECT_EQ(selected.engineId, 42); + EXPECT_TRUE(selected.succeeded); + EXPECT_FLOAT_EQ(selected.minTimeMs, 1.0f); +} + +TEST(TestSelectEngineResult, TheEngineOwnFailureWinsOverTheAutotuneStatus) +{ + auto failed = resultFor(42); + failed.succeeded = false; + failed.errorMessage = "Benchmark failed on iteration 2"; + + const auto selected = selectEngineResult(42, {failed}, K_NO_ENGINE_SUCCEEDED); + + EXPECT_FALSE(selected.succeeded); + EXPECT_EQ(selected.errorMessage, "Benchmark failed on iteration 2"); +} + +TEST(TestSelectEngineResult, NeverReportsAnExcludedOrDifferentEngine) +{ + for(const auto& results : {std::vector{resultFor(42, true)}, + std::vector{resultFor(7)}}) + { + const auto selected = selectEngineResult(42, results, K_GOOD); + + EXPECT_FALSE(selected.succeeded); + EXPECT_EQ(selected.engineId, 42); + EXPECT_EQ(selected.errorMessage, "autotune returned no result for the engine under test"); + } +} + +TEST(TestSelectEngineResult, NoResultsReportsTheAutotuneStatus) +{ + const auto good = selectEngineResult(42, {}, K_GOOD); + const auto bad = selectEngineResult(42, {}, K_NO_ENGINE_SUCCEEDED); + + EXPECT_FALSE(good.succeeded); + EXPECT_EQ(good.errorMessage, "autotune returned no result for the engine under test"); + EXPECT_FALSE(bad.succeeded); + EXPECT_EQ(bad.errorMessage, "no engine succeeded"); +} + +// ---- toString(TimingQuality) ------------------------------------------------ + +TEST(TestTimingQualityName, NamesEveryValueAndRefusesUnknownValues) +{ + EXPECT_STREQ(toString(TimingQuality::DEVICE_ONLY), "DEVICE_ONLY"); + EXPECT_STREQ(toString(TimingQuality::UNSTALLED), "UNSTALLED"); + EXPECT_STREQ(toString(TimingQuality::INVALID), "INVALID"); + EXPECT_THROW(toString(static_cast(99)), std::invalid_argument); +} + // ---- reportBenchmark ------------------------------------------------------ TEST(TestReportBenchmark, RecordsTheResultAsProperties) @@ -231,6 +312,7 @@ struct OracleCase const char* name; VerificationMode mode; bool hasGoldenOutput; + Verifier verifier; friend void PrintTo(const OracleCase& oracleCase, std::ostream* os) { @@ -241,9 +323,9 @@ struct OracleCase /// AUTO with golden data and explicit GOLDEN compare against the golden file. /// AUTO without golden data falls through to the reference oracle chain. const std::array ORACLE_CASES{{ - {"AutoGolden", VerificationMode::AUTO, true}, - {"AutoReference", VerificationMode::AUTO, false}, - {"GoldenMode", VerificationMode::GOLDEN, true}, + {"AutoGolden", VerificationMode::AUTO, true, Verifier::GOLDEN}, + {"AutoReference", VerificationMode::AUTO, false, Verifier::GPU_REFERENCE}, + {"GoldenMode", VerificationMode::GOLDEN, true, Verifier::GOLDEN}, }}; /// Expectations are checked after the harness returns: a gmock or EXPECT failure @@ -260,18 +342,27 @@ class TestBundleHarnessBenchmark : public ::testing::TestWithParam /// the verified output tensor. std::vector _benchmarkCalls; bool _sawVerifiedOutput = false; + /// When set, the mocked runner throws it from benchmark(). + std::optional _benchmarkThrows; + /// The verifier the harness reported for each bundle it ran. + std::vector _verifiers; void SetUp() override { using ::testing::_; testing_support::ensureTestConfigInitialized(); _scopedDir.emplace(hipdnn_test_sdk::utilities::claimScratchDirectory("bundle_benchmark")); + testing_support::captureVerifiers(_mocks.reporter, _verifiers); ON_CALL(_mocks.engineRunner, benchmark(_, _, _)) .WillByDefault( [this](GraphSession&, VariantPack& variantPack, const BenchmarkSettings& s) { _benchmarkCalls.push_back(s); _sawVerifiedOutput = variantPack.count(fixtures::K_OUTPUT_UID) == 1; + if(_benchmarkThrows.has_value()) + { + throw std::runtime_error(*_benchmarkThrows); + } return _benchmarkResult; }); @@ -319,6 +410,7 @@ TEST_P(TestBundleHarnessBenchmark, AVerifiedPassIsTimedWithThePolicySettings) EXPECT_EQ(_benchmarkCalls[0].timedIterations, 4); EXPECT_EQ(_benchmarkCalls[0].warmupIterations, 1); EXPECT_TRUE(_sawVerifiedOutput) << "the benchmark must run on the verified tensors"; + EXPECT_EQ(_verifiers, std::vector{GetParam().verifier}); } TEST_P(TestBundleHarnessBenchmark, AMismatchIsNotTimed) @@ -364,6 +456,24 @@ TEST_P(TestBundleHarnessBenchmark, ABenchmarkFailureFailsAVerifiedBundle) << testing_support::allMessages(results); } +// A throw from the benchmark must not turn the verified PASS into a harness failure, +// which would also change the support-claim verdict. +TEST_P(TestBundleHarnessBenchmark, AThrowingBenchmarkFailsButKeepsTheVerifiedOutcome) +{ + _benchmarkThrows = "hipMalloc failed"; + + ::testing::TestPartResultArray results; + run(fixtures::K_OUTPUT_VALUE, settings(4, 1), &results); + + EXPECT_EQ(_benchmarkCalls.size(), 1u); + EXPECT_TRUE(testing_support::anyFailed(results)); + EXPECT_NE(testing_support::allMessages(results).find("benchmark failed: hipMalloc failed"), + std::string::npos) + << testing_support::allMessages(results); + EXPECT_EQ(_verifiers, std::vector{GetParam().verifier}) + << "the outcome the harness recorded must still be the verified pass"; +} + INSTANTIATE_TEST_SUITE_P(Oracles, TestBundleHarnessBenchmark, ::testing::ValuesIn(ORACLE_CASES),