diff --git a/dnn-providers/integration-tests/docs/README.md b/dnn-providers/integration-tests/docs/README.md index 4de20a7f8e31..463b30a0f2d6 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. | | [Support Matrix and Feature Coverage](support-matrix.md) | Render the human-readable matrix of which engine supports which graph on which GPU and OS, validate the sidecars it is built from, see the pre-commit and CI checks around it, and measure per-op feature coverage. | @@ -51,6 +51,7 @@ other rows do not, so they are easy to miss. | `UNVERIFIABLE BUNDLES` / `Unverifiable: … tried: …` | **yes** (those tests failed) | The engine ran the bundle, but no oracle could check its output: no golden data and no reference executor for the op. Nothing was tested. | [Unverifiable bundles](running-tests.md#unverifiable-bundles); fix: [When no reference executor supports the op](adding-tests.md#when-no-reference-executor-supports-the-op) | | `Error: zero tests ran.` | **yes** | A discovery or filter misconfiguration, or, with GTest sharding, an empty shard. | [Hard stops](running-tests.md#hard-stops) | | `FATAL: --enforce-support-claims is active and …` | **yes** | Enforcement verified nothing. | [Hard stops](running-tests.md#hard-stops) | +| `benchmark failed: …` | **yes** (that test failed) | Under `--benchmark`, the verified graph failed when it ran again, or the timer returned too few measurements. | [Time the engine](running-tests.md#4-time-the-engine---benchmark) | | `unclaimed_support` in the `SUPPORT CLAIM SUMMARY` | no | The engine accepts graphs that no sidecar claims. Record them. | [Updating support claims](adding-tests.md#updating-support-claims) | | `Skipped:` equal to (or close to) the total | no | Green, but little or nothing was tested. Read the skip reasons. | [Test coverage summary](running-tests.md#test-coverage-summary) | | `WARNING: Bundle tests are enabled but …` / `WARNING: No bundles could be loaded from …` | no | No bundles registered. Only compiled-in C++ tests ran. | [Troubleshooting](running-tests.md#troubleshooting) | diff --git a/dnn-providers/integration-tests/docs/running-tests.md b/dnn-providers/integration-tests/docs/running-tests.md index a90d7e7aa663..b7a8f32cf481 100644 --- a/dnn-providers/integration-tests/docs/running-tests.md +++ b/dnn-providers/integration-tests/docs/running-tests.md @@ -128,6 +128,78 @@ build/bin/hipdnn_integration_tests \ runs anything. - Standard GTest sharding (`GTEST_TOTAL_SHARDS` / `GTEST_SHARD_INDEX`) works. +### 4. Time the engine (`--benchmark`) + +`--benchmark` adds a performance measurement to each test. The test first +verifies the engine output in the usual way. If the output passes, the harness +times the same engine on the same tensors with `Graph::autotune()`, limited to +that one engine. The C++ graph tests and the bundles both support it. + +A test is not timed when it skips, fails, or passes without output +verification (a bundle at enforcement level APPLICABLE or BUILDABLE). Such a +test prints no `[ BENCHMARK ]` line. A filter that selects only such tests +gives a green run with no timings, so count the `[ BENCHMARK ]` lines. + +```bash +build/bin/hipdnn_integration_tests \ + --test-article build/lib/hipdnn_plugins/engines/libmiopen_plugin.so \ + --test-engine MIOPEN_ENGINE \ + --gtest_filter='quick_ConvolutionFwd_Default.*' \ + --benchmark \ + --gtest_output=json:benchmark.json +``` + +Each timed test prints one line: + +```text +[ BENCHMARK ] quick_ConvolutionFwd_Default.8_16_16_16_bfp16_nhwc_dil1x1_postpad0x0_prepad0x0: engine=MIOPEN_ENGINE tuned=0 timing=DEVICE_ONLY samples=20 min=0.0181 ms robust=0.0184 ms mean=0.0184 ms stddev=0.0002 ms +``` + +- Autotune runs `--benchmark-warmup` untimed executions (default 30). Then it + runs `--benchmark-iterations` timed executions (default 20), with the + `FIXED_AVERAGE` strategy and the same stall gate, retry and statistics + rules as any other autotune call. +- The GPU idles between tests, so each test starts cold. Too few warmup + executions make the times read high: on MI210 the quick-tier bundles read + 4-60% slow with 1-10 warmups and settle after about 30. More timed + iterations do not fix this; more warmup does. Raise `--benchmark-warmup` + if `stddev` stays large compared with `mean`. +- `robust` is the mean after autotune drops the slow tail: what the engine + usually costs. `min`, `mean` and `stddev` are over all timed samples. +- `tuned=0` means that the engine chose its solution with its heuristics. The + harness sets no knobs, so this is the default. To time the solution that the + engine picks by benchmarking its candidates, set + `HIPDNN_FORCE_BENCHMARKING=1`. The line then shows `tuned=1` for engines + that support the `global.benchmarking` knob (MIOpen, the hip-kernel + ingestor). Engines without the knob, such as hipBLASLt, still show + `tuned=0`. The selection runs during the verification execution, so its + cost is not in the timed samples. +- `timing=DEVICE_ONLY` means that the stream was stalled, so the times exclude + host submission overhead. `timing=UNSTALLED` means that autotune could not + stall the stream and re-measured without the stall gate, so the times can + include host time. Do not compare the two kinds. +- The test fails with `benchmark failed: …` when autotune cannot benchmark the + engine, for example when an execution fails or too many timed executions + return no measurement. +- `--gtest_output=json:` or `xml:` writes the same values as the + test properties `benchmark_engine`, `benchmark_tuned`, `benchmark_timing`, + `benchmark_samples`, `benchmark_min_ms`, `benchmark_robust_ms`, + `benchmark_mean_ms` and `benchmark_stddev_ms`. +- `engine=` names the engine that ran. To compare engines, run once for each + `--test-engine`. Without `--test-engine`, hipDNN selects the engine. +- To time your own graph, write it as a single-graph bundle + (`/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 +227,9 @@ Otherwise, you test the old tree. | `--fail-on-unsupported` | | off | Fail instead of skip when no engine supports a graph. **C++ graph tests only.** | | `--generate-support-matrix[=file]` | | off | Write a markdown matrix. **C++ graph tests only**; on a default build it writes an empty table. | | `--capture-bundles ` | | off | Dump compiled-in C++ graph tests as bundle JSON (see [Adding Tests](adding-tests.md#from-an-existing-c-graph-test)). | +| `--benchmark` | | off | Time each verified test. See [Time the engine](#4-time-the-engine---benchmark). | +| `--benchmark-iterations` | | 20 | Timed executions per test. Requires `--benchmark`. | +| `--benchmark-warmup` | | 30 | Untimed executions before timing. Requires `--benchmark`. | | `--gtest_*` | | | Passed through to GTest. | The binary refuses these flag combinations (exit 1): @@ -164,6 +239,10 @@ The binary refuses these flag combinations (exit 1): - `--write-support-claims` together with `--enforce-support-claims` or `=true`. - `--write-support-claims` without `--test-article`, or without a bundle root (`--golden-data-dir` or `HIPDNN_TEST_GOLDEN_DATA_DIR`). +- `--benchmark-iterations` or `--benchmark-warmup` without `--benchmark`, or + `--benchmark-iterations` less than 1, or `--benchmark-warmup` less than 0. +- `--benchmark` together with `--skip-graph-validation`, `--capture-bundles` + or `--write-support-claims`. These modes do not execute the graph. `--enforce-support-claims=false` is accepted in every combination. If enforcement is only inherited and there is no `--test-engine`, the run 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..e752d8aa1d61 --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/Benchmark.hpp @@ -0,0 +1,80 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include + +namespace hipdnn_integration_tests +{ + +/// How --benchmark times one verified graph. +struct BenchmarkSettings +{ + static constexpr int DEFAULT_TIMED_ITERATIONS = 20; + // The GPU idles between tests (input fill, reference, comparison), so each test + // starts cold. On MI210, quick-tier bundles read 4-60% slow with 1-10 warmup + // executions and settle after about 30. More timed iterations do not fix this. + static constexpr int DEFAULT_WARMUP_ITERATIONS = 30; + + int timedIterations = DEFAULT_TIMED_ITERATIONS; + int warmupIterations = DEFAULT_WARMUP_ITERATIONS; +}; + +/// The benchmark flags as typed, plus the flags that stop a test before it executes. +struct BenchmarkRequest +{ + bool enabled = false; + std::optional 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}; +} + +} // 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 new file mode 100644 index 000000000000..10faa34ff9f1 --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/GraphBenchmark.hpp @@ -0,0 +1,56 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#pragma once + +#include +#include +#include + +#include +#include + +#include "harness/Benchmark.hpp" +#include "harness/BenchmarkReport.hpp" +#include "harness/PlanWorkspace.hpp" +#include "harness/VariantPack.hpp" + +namespace hipdnn_integration_tests +{ + +/// 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) +{ + hipdnn_frontend::AutotuneResult failed; + + int64_t engineId = -1; + if(auto status = graph.get_execution_plan_engine_id(engineId); !status.is_good()) + { + failed.errorMessage = status.get_message(); + return failed; + } + + auto workspace = allocatePlanWorkspace(graph, failed.errorMessage); + if(!workspace.has_value()) + { + return failed; + } + + hipdnn_frontend::AutotuneConfig config; + config.strategy = hipdnn_frontend::AutotuneStrategy::FIXED_AVERAGE; + config.timedIterations = settings.timedIterations; + config.warmupIterations = settings.warmupIterations; + config.engineIdFilter = {engineId}; + + std::vector results; + const auto status = graph.autotune(handle, variantPack, workspace->get(), config, {}, &results); + 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 0d56009fa259..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 @@ -30,8 +29,10 @@ #include #include +#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" @@ -260,6 +261,14 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam attr, @@ -548,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/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/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/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/FrontendGraphEngineRunner.cpp b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp index e764e3c5f480..ab66d2ced76a 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/FrontendGraphEngineRunner.cpp @@ -4,12 +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" @@ -121,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()); } @@ -150,4 +148,16 @@ EngineOpResult return result; } +hipdnn_frontend::AutotuneResult FrontendGraphEngineRunner::benchmark( + GraphSession& session, VariantPack& variantPack, const BenchmarkSettings& settings) +{ + if(session.graph == nullptr) + { + hipdnn_frontend::AutotuneResult failed; + failed.errorMessage = "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..813ff2d2eee3 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; + + 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/HarnessPolicy.hpp b/dnn-providers/integration-tests/src/harness/bundle/HarnessPolicy.hpp index 4bff112d5ee8..aee67f728003 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 @@ -132,6 +133,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..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,9 @@ #include #include +#include + +#include "harness/Benchmark.hpp" #include "harness/VariantPack.hpp" #include "harness/bundle/GraphSession.hpp" #include "harness/bundle/LoadedEngine.hpp" @@ -92,6 +95,13 @@ class IGraphEngineRunner const std::optional& engineUnderTest, VariantPack& variantPack) = 0; + + /// 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; }; } // 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 9b07c6d96bd5..e1626db6e469 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.cpp @@ -492,7 +492,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) @@ -507,17 +509,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); } - oracles.declined("golden (absent)"); + else + { + oracles.declined("golden (absent)"); + } + } + if(!outcome.has_value()) + { + outcome = runOracleChain(engine.outputs, oracles); } - return runOracleChain(engine.outputs, oracles); + benchmarkIfVerified(session, engine.outputs, *outcome); + return *outcome; } VerificationOutcome @@ -872,6 +883,30 @@ 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; + } + // 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 ------------------------------------------------------------ VerificationOutcome diff --git a/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp b/dnn-providers/integration-tests/src/harness/bundle/IntegrationBundleVerificationHarness.hpp index 55c6c4efa33b..d1b02d496ad7 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" @@ -233,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); } @@ -492,6 +499,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, @@ -545,6 +558,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 f9fb0e378951..6c28a75503be 100644 --- a/dnn-providers/integration-tests/src/harness/bundle/ProductionPolicy.cpp +++ b/dnn-providers/integration-tests/src/harness/bundle/ProductionPolicy.cpp @@ -25,6 +25,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 a1c6e254295f..56dfa29f23cb 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" @@ -191,6 +192,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 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."); + parser.add_argument("--benchmark-warmup") + .scan<'i', int>() + .help("Untimed executions before timing (default 30). Requires --benchmark."); std::vector remainingArgs; try @@ -387,6 +401,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); @@ -400,6 +428,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..78b4e3505c3c --- /dev/null +++ b/dnn-providers/integration-tests/tests/TestBenchmark.cpp @@ -0,0 +1,484 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include + +#include "BundleFixtureFiles.hpp" +#include "HarnessTestSupport.hpp" +#include "harness/Benchmark.hpp" +#include "harness/BenchmarkReport.hpp" + +// NOLINTBEGIN(readability-identifier-naming) + +using namespace hipdnn_integration_tests; +using namespace hipdnn_integration_tests::bundle; +using hipdnn_frontend::AutotuneResult; +using hipdnn_frontend::TimingQuality; + +namespace +{ + +BenchmarkSettings settings(int timed, int warmup) +{ + BenchmarkSettings result; + result.timedIterations = timed; + result.warmupIterations = warmup; + return result; +} + +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; +} + +AutotuneResult failedResult(std::string message) +{ + AutotuneResult result; + result.errorMessage = std::move(message); + return result; +} + +} // namespace + +// ---- 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, 30); +} + +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); + } +} + +// ---- 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) +{ + 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; + 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_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) +{ + 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 ------------------------------------------------------- + +namespace +{ + +/// One way the harness verifies a bundle. Each way has its own benchmark hook. +struct OracleCase +{ + const char* name; + VerificationMode mode; + bool hasGoldenOutput; + Verifier verifier; + + 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, 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 +/// raised inside driveHarness() lands in its intercepted results, not in this test. +class TestBundleHarnessBenchmark : public ::testing::TestWithParam +{ +protected: + std::optional _scopedDir; + testing_support::HarnessMocks _mocks; + + /// What the mocked runner returns from benchmark(). + AutotuneResult _benchmarkResult = successfulResult(); + /// The settings of every benchmark() call, and whether its variant pack held + /// 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; + }); + + // 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 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(GetParam().mode); + policy.benchmark = benchmark; + IntegrationBundleVerificationHarness harness(_mocks.dependencies(policy)); + harness.setBundle( + fixtures::loadBundle(_scopedDir->path(), "bench", GetParam().hasGoldenOutput), "bench"); + + testing_support::driveHarness(harness, results); + } +}; + +} // namespace + +TEST_P(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"; + EXPECT_EQ(_verifiers, std::vector{GetParam().verifier}); +} + +TEST_P(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_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_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 = failedResult("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); +} + +// 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), + [](const ::testing::TestParamInfo& info) { + return std::string(info.param.name); + }); + +// 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..450ecad5d9df 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(hipdnn_frontend::AutotuneResult, + benchmark, + (GraphSession & session, + VariantPack& variantPack, + const BenchmarkSettings& settings), + (override)); }; } // namespace hipdnn_integration_tests::bundle