Skip to content

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

Merged
SamuelReeder merged 7 commits into
developfrom
users/sareeder/hipdnn-integration-perf-driver
Oct 27, 2026
Merged

SamuelReeder merged 7 commits into
developfrom
users/sareeder/hipdnn-integration-perf-driver

Conversation

@SamuelReeder

@SamuelReeder SamuelReeder commented Oct 19, 2026 •

Copy link
Copy Markdown
Contributor

Summary

Add an opt-in --benchmark mode to hipdnn_integration_tests. Each bundle or C++ graph test that verifies the engine output is then timed on the same engine and tensors with Graph::autotune() (STANDARD mode). This gives a C++ accuracy-plus-performance driver for bundled shapes and custom graphs, as requested in #11652, without a new binary or a second timing implementation.

JIRA ID : #11652

Closes #11652

Risk Assessment

Risk level 2. The base rubric puts new feature paths and test infrastructure at level 3. This PR is at level 2 because --benchmark is off by default, so no CI lane or default run executes the new timing path. The change is limited to the integration-test harness under dnn-providers/integration-tests/. Without --benchmark, the only behavior changes are a restructure of runReferenceMode() that keeps the same oracle order, and the move of the workspace allocation in FrontendGraphEngineRunner::execute() into a shared helper. Unit tests cover both for AUTO with golden data, AUTO with the reference chain, and GOLDEN mode. The hipDNN library, the providers and the plugin ABI do not change. Timing reuses the public Graph::autotune(), limited to the engine under test, so the stall gate from #11813, the retries and the statistics are autotune's own.

ASIC Coverage

Host-side harness code only. No kernel, dispatch, default or support-surface change. Standard PR CI is sufficient. The timed path was exercised locally on gfx90a (MI210) with the MIOpen and hipBLASLt providers.

Testing Summary

  • Unit (no device): flag resolution, engine-result selection, gtest property output and failure reporting, and the bundle harness. The harness tests run for each oracle path (AUTO with golden data, AUTO with the reference chain, GOLDEN mode): only a verified pass is timed, the policy settings reach the runner, a benchmark failure or throw fails the test without changing the verified outcome, and a run without --benchmark passes or fails as before. Mutations caught: removing the "verified pass only" guard; removing the runGoldenMode() hook; removing the reference-chain hook; narrowing the benchmark catch. The timing loop and statistics are autotune's and are covered by the frontend autotune tests.
  • GPU: the bundle path (pinned and unpinned engine; AUTO, GOLDEN and reference verification), the C++ graph path, a custom graph-only bundle via --golden-data-dir, --gtest_output=json, and rocprofv3 kernel stats (with warmup 3, the engine kernel ran 1 + 3 warmup + 20 timed times, with 20 stall waits).
  • Default counts: on MI210, each warmup/iteration pair ran 3 times over the 188 timed MIOpen quick bundles and the 16 hipBLASLt Matmul bundles. The reference was warmup 10 with 200 timed iterations. With 1-10 warmups, the median robust time read 4-60% above the reference, and more timed iterations did not remove that. With 30 warmups, the median bias was 0.0-0.4% (p90 at most 2.3%), and the spread between runs was about 1%. The defaults are therefore 30 warmup and 20 timed executions. The wall time of the run did not change measurably (about 9 s for the MIOpen quick tier).
  • Tuned or heuristic solution: the harness sets no knobs, so engines run their heuristic choice and the line shows tuned=0. With HIPDNN_FORCE_BENCHMARKING=1, MIOpen logged tuning policy 3 (search) instead of 1 (none), and all 52 quick convolution bundles showed tuned=1. hipBLASLt has no global.benchmarking knob and showed tuned=0 for all 16 Matmul bundles with the same setting.
  • Regression: the MIOpen quick lane without --benchmark, and the full harness unit suite.

Testing Checklist

  • Harness unit tests - build/bin/hipdnn_integration_tests_unit_tests - Status: Passed (1056 passed, 4 skipped as before)
  • Benchmark unit tests - hipdnn_integration_tests_unit_tests --gtest_filter='*Benchmark*' - Status: Passed
  • tuned field - MIOpen quick_Convolution* with and without HIPDNN_FORCE_BENCHMARKING=1, hipBLASLt quick_Matmul* with it - Status: Passed (52 tuned=0, 52 tuned=1, 16 tuned=0) - gfx90a
  • Test name validation - ctest --test-dir build/dnn-providers/integration-tests -R hipdnn-integration-tests_test_name_validation - Status: Passed
  • GOLDEN and AUTO modes on a golden bundle (dvc pull of quick/BatchnormFwdInference/Inference/golden/small_fp32_nchw) - --verification-mode golden|auto --benchmark - Status: Passed (verified by golden, timed) - gfx90a
  • MIOpen quick bundles with timing - hipdnn_integration_tests --test-article libmiopen_plugin.so --test-engine MIOPEN_ENGINE --test-config MIOPEN_ENGINE.toml --gtest_filter='quick_*' --benchmark --benchmark-iterations 5 - Status: Passed (188 passed, 188 timed, all DEVICE_ONLY, 0 failed) - gfx90a
  • Unpinned hipBLASLt bundles with timing - hipdnn_integration_tests --gtest_filter='quick_Matmul*' --benchmark - Status: Passed (16/16 timed, engine=HIPBLASLT_ENGINE) - gfx90a
  • C++ graph tests with timing - -DBUILD_CPP_GRAPH_TESTS=ON, --no-bundles --gtest_filter='Smoke/*ConvFwd*' --benchmark --benchmark-iterations 10 - Status: Passed (108 timed, 0 failed) - gfx90a
  • Custom graph-only bundle - --golden-data-dir /tmp/custom --gtest_filter='MyBn*' --benchmark - Status: Passed (verified by gpu_ref, timed) - gfx90a
  • Flag validation - invalid combinations exit 1 with a message - Status: Passed
  • MIOpen quick lane, no --benchmark - ctest --test-dir build/dnn-providers/miopen-provider -R '^miopen-provider-external-integration_quick_suite$' - Status: Passed
  • Build with clang-tidy and clang-format on (hipdnn-providers preset) - Status: Passed
  • pre-commit on changed files - Status: Passed
  • PR CI - GitHub PR checks - Status: Pending

Technical Changes

  • src/harness/Benchmark.hpp (new): BenchmarkSettings and resolveBenchmark() flag validation. TestConfig.hpp includes only this header.
  • src/harness/BenchmarkReport.hpp (new): selectEngineResult() (picks the engine under test from the autotune results), toString(TimingQuality), the tuned check, and reportBenchmark() (one [ BENCHMARK ] line plus benchmark_* gtest properties; a result that did not succeed fails the test).
  • src/harness/CurrentTestName.hpp (new): currentTestName(), moved out of TomlGuards.hpp so the report can reuse it.
  • src/harness/GraphBenchmark.hpp (new): benchmarkGraph() calls Graph::autotune() with engineIdFilter set to the active plan's engine, FIXED_AVERAGE, and the --benchmark warmup and iteration counts, and returns selectEngineResult().
  • src/harness/PlanWorkspace.hpp (new): allocatePlanWorkspace(), the workspace size query and allocation shared by FrontendGraphEngineRunner::execute(), IntegrationGraphVerificationHarness::executeGpuGraph() and the benchmark.
  • Bundle harness: benchmarkIfVerified() catches exceptions from the benchmark and reports them as a failed benchmark, so a throw cannot replace the verified outcome.
  • Bundle harness: IGraphEngineRunner::benchmark() and HarnessPolicy::benchmark. A test is timed only after a PASSED, VERIFIED comparison. A benchmark failure fails the test but does not change the support-claim verdict.
  • C++ graph harness: times the graph after validateOutputs() when the test has no failure.
  • main.cpp: --benchmark, --benchmark-iterations (default 20), --benchmark-warmup (default 30). The binary refuses --benchmark with --skip-graph-validation, --capture-bundles or --write-support-claims, because those modes do not execute the graph.
  • Docs: docs/running-tests.md (usage, output, which tests are not timed, custom graphs, rocprofv3 kernel stats, CLI table) and docs/README.md (signal row).

Out of scope: per-op FLOP and bandwidth figures, and an all-engines sweep in one run. Use one run for each --test-engine, or dnn-benchmarking.

…raphs

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
@therock-pr-bot

therock-pr-bot Bot commented Oct 19, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🔎 pre-commit ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —
🤖 therock-pr-bot ✅ Pass —

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

therock-pr-bot Bot commented Oct 19, 2026 •

Copy link
Copy Markdown

🎉 All checks passed! This PR is ready for 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.
…tune

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.
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.
@SamuelReeder
SamuelReeder marked this pull request as ready for review October 20, 2026 18:04
@SamuelReeder
SamuelReeder requested a review from a team as a code owner October 20, 2026 18:04
…ned 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.

@adickin-amd adickin-amd left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

few small nitpicks, lgtm otherwise

@CMiservaAMD CMiservaAMD left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM overall, just the one thought on exposing RUN_UNTIL_STABLE as an alternative optional strategy.

SamuelReeder and others added 2 commits October 22, 2026 17:11
- 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.
@SamuelReeder
SamuelReeder merged commit 69b1bc7 into develop Oct 27, 2026
98 checks passed
@SamuelReeder
SamuelReeder deleted the users/sareeder/hipdnn-integration-perf-driver branch October 27, 2026 05:16
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

adickin-amd added a commit that referenced this pull request Oct 30, 2026
Brings in the SDPA backward O/LSE derivation (#14312), the cuDNN left-bound semantics (#14483), the SDPA kernarg range checks (#14075) and --benchmark (#13915). Conflicts: the outcome publishing moved into runBody() on this branch, so the benchmark report now sits there; TestFillInputs and the docs keep both sides.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[Feature]: [hipdnn integration-tests] performance tools

3 participants