Repository navigation
perf(hipdnn): speed up superbuild CI with parallel ctest and compiled reference executors - #14680
Merged
Merged
Conversation
PlanRegistrySignatureKeyEqual and GpuPlanRegistrySignatureKeyEqual used std::visit(*this, a, b), which instantiates the visitor for every pair of alternatives (22x22 = 484 and 15x15 = 225). Compare the active alternative indices first, then visit only `a` and fetch the same alternative from `b` (22 and 15 instantiations). Semantics are unchanged: keys of different alternatives are unequal, keys of the same alternative compare with ==. Add TestGpuPlanRegistrySignatureKey, pinning same-type equal and hash-equal, same-type different-value unequal, cross-type unequal, and unordered_map lookup with the key hash/equal functors.
…integration harness once
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
…eport the test selection" This reverts commit eeb145e.
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #14680 +/- ##
========================================
Coverage 77.43% 77.43%
========================================
Files 2902 2902
Lines 489842 489841 -1
Branches 74142 74143 +1
========================================
+ Hits 379301 379303 +2
+ Misses 85573 85570 -3
Partials 24968 24968
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
…g hipdnn_golden_data_tests The golden-data binary was registered before DNN_PROVIDER_TEST_CATEGORY_YAMLS was set, so _add_test_target_internal() still added the raw, unlabeled hipdnn_golden_data_tests test next to its tiered *_suite entries and the binary ran twice. Move the categorization block above the registration.
…ddickin/ref_executor_compile_once
…us shard hkp_reserve_census_shard() was called for each FlyDSL engine and architecture without the generated TOML, so the shard check added in #14573 failed configure for any architecture with FlyDSL content (for example gfx1151: "the gfx1151 census shard names hipkernel_flydsl_rmsnorm_gfx1151.toml, but nothing places it beside the shard"). Pass the file, as the gfx950 registration already does.
…_executor_compile_once # Conflicts: # dnn-providers/hip-kernel-provider/src/CMakeLists.txt # dnn-providers/integration-tests/CMakeLists.txt # dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp
…reference The scalar CPU SdpaBackward reference needs seconds to tens of seconds per 256x128 bundle, and twelve of them in the quick tier exceeded the 600 s ctest timeout of hipdnn_golden_data_tests_quick_suite. referenceShapeIsAffordable now caps the CPU backward at 4M multiply-accumulates per Q*K^T pass (Q elements times Skv), in addition to the existing quick-tier limit. The 64-row padded and GQA bundles stay on CPU; the excluded bundles stay covered by the GPU lane.
…_executor_compile_once
Contributor
Author
|
Ci failure due to #13414, will override merge |
Contributor
TheRock Submodule Bump ActivityNewest first
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Shortens the hipDNN superbuild CI. Two changes: both lanes run
ctest --parallel 16(the larger saving), and the integration-tests / hip-kernel-provider test code stops instantiating the CPU/GPU reference executors in every test file (the smaller one). No product code changes; nothing underprojects/hipdnnbecomes compiled (the one edit there is a header-only equality fix).JIRA ID : ALMIOPEN-3113
Risk Assessment
Risk 2. CI workflow plus test-support code only. The residual risk is a test that is not safe to run in parallel; I audited the registrations (temp dirs use
tmp_path/unique names, shared-state tests areRUN_SERIAL/fixture-ordered) but only one parallel run exists so far (see Testing).ASIC Coverage
Architecture-independent: host-side test code and a CI workflow. Passing PR CI is sufficient; no multi-arch sweep needed. Built and tested locally on gfx90a (MI250). GPU cases for SDPA/flash2 skip on gfx90a as on base.
Testing Summary
-L '^full$'166/166 passed, second call 31/31. Windows: 249/249 and 31/31 (Windows runs 78 more tests than before because develop added the FlyDSL entries).hipdnn_golden_data_testsnow runs only as its tier suites in the first call. Earlier runs of plain--parallel 16(experiment PR ci(hipdnn): [DO NOT MERGE] superbuild timing experiment #14563 and earlier heads of this PR) also passed, so no parallel flakes in any run.shared/ctest, whoseTestCategories.cmakeregisters the YAML tier suites (*_quick_suite,*_full_suite, ...), so CI ran 183 / 189 plain unlabeled tests and-L '^full$'alone selected nothing ("No tests were found"). This PR now checks outshared/ctest: CI registers 238 (Linux) / 244 (Windows) entries, where the tiers repeat one gtest filter (for examplemiopen_plugin_testsandhip_kernel_provider_testsfour times each), andhipdnn_golden_data_testsran twice, as its tier suites and as an unlabeled raw test (fixed below).-L '^full$'runs each binary once (fullis a superset);-LE '^(quick|standard|comprehensive|ffm-quick|full)$'runs the entries with no tier label (rocke_*,hkp-probe-*,shared_ctest_*and the gpu_ref entries). Both calls always run and the step fails if either does. All of the 171 entries that were registered before and are still registered run; the 12 plain provider/plugin entries thatshared/ctestreplaces are covered by their*_full_suite. Enabling it also registers the twoshared_ctest_*tests, which configure a nested CMake project; they neededCMAKE_GENERATOR=Ninja(Linux; the runner has no make) andPython3_ROOT_DIRset to the superbuild venv (actions/setup-pythonexports its ownPython3_ROOT_DIR, whose Python has no PyYAML).hipdnn-providers-all+HIPDNN_ENABLE_SDPA,BUILD_CPP_GRAPH_TESTS,HIPDNN_ENABLE_KERNEL_INGESTOR,ENABLE_HIP_FLASH2_ENGINE. Per-suite pass/skip/fail base vs head:hipdnn_test_sdk_tests: 2420/4/0 -> 2420/4/0hipdnn_integration_tests_unit_tests: 1115/2/0 -> 1120/2/0 (only additions)hip_kernel_provider_integration_tests: 36/72/0 -> 36/72/0TestGpuPlanRegistrySignatureKeyfails when equality is mutated to compare only the variant index..cppfiles: 0 findings. The build container has no clang-tidy, so the build ran with tidy off and tidy was a separate pass.hipdnn,hipdnn-integration-tests,miopen-providerclean. The hip-kernel-provider component reports 5...FP16...names inTestHipFlash2*that exist in develop (only visible withENABLE_HIP_FLASH2_ENGINE=ON, default OFF); not touched here, not checked at base.Testing Checklist
hipdnn-providers-all- Status: Passedfulltier selection - this PR's superbuild CI (Linux and Windows) - Status: PassedTechnical Changes
.github/workflows/hipdnn-superbuild-ci.yml: both lanes check outshared/ctestand runctest --parallel 16 -L '^full$'followed byctest --parallel 16 -LE '<tier labels>'(see Testing for why); the test step setsPython3_ROOT_DIRto the venv (both lanes) andCMAKE_GENERATOR=Ninja(Linux) for the nestedshared_ctest_*configure.dnn-providers/integration-tests/CMakeLists.txt: the YAML categorization block moves above thehipdnn_golden_data_testsregistration._add_test_target_internal()skips the rawadd_test()only whileDNN_PROVIDER_TEST_CATEGORY_YAMLSis already set, so the binary was also registered as an unlabeled raw test (about 2 s on Linux, 8 s on Windows). Withoutshared/ctestthe raw test is still registered.dnn-providers/hip-kernel-provider/src/CMakeLists.txt:hkp_reserve_census_shard(... FILES "${_toml}")for each FlyDSL engine and architecture. After fix(hip-kernel-provider): place shard test configs beside the build-tree census shard #14573 (shard TOML check) and feat(hip-kernel-provider): FlyDSL kernels: RMSNorm fwd/bwd + SDPA forward (RDNA, CDNA) #14111 (FlyDSL) merged, configure failed for any architecture with FlyDSL content ("the gfx1151 census shard names hipkernel_flydsl_rmsnorm_gfx1151.toml, but nothing places it beside the shard"); every superbuild Windows configure on develop failed from 21:15 UTC. Same call the gfx950 registration already makes. This PR carries the one-line fix because its CI runs against current develop.PlanRegistrySignatureKey.hpp(test_sdk, stays header-only) andGpuPlanRegistrySignatureKey.hpp: equality comparesindex()then runs onestd::visit(was a pairwise visit: 22x22 and 15x15 instantiations). Semantics unchanged.ReferenceGraphExecutorFactory,CpuReferenceGraphExecutorAdapterandGpuReferenceGraphExecutormove from inline headers to pimpl +.cppindnn-providers/integration-tests; 12 plan tests and the chained SDPA test use the adapter instead of the raw test_sdk executor.ReferenceExecutorPool.cppstays out of the unit-test target.executeCpuReferenceGraph.TestGpuPlanRegistrySignatureKey; a real adapter case inTestErrorPaths(the old test never called the adapter and is renamed to say what it checks).Measured effect
Be honest about the executor refactor: it cut tidy CPU on the files it touched by 28% (e.g.
ReferenceExecutorPool.cpp465 s -> 34 s) but only about 3% of the whole Linux tidy work, and wall time did not move because each of the three test executables compiles the new.cppfiles. Expect roughly a minute or less on the Linux job from it. Plain compile time was not measured.What else could save time (not in this PR)
Measured on the experiment PR and a Linux GPU machine unless marked estimate:
.clang-tidyto every compile-database entry gives 338k false findings on rocKE; per-component configs are needed for parity.performance-unnecessary-value-param: 25% of a heavy file's check time, under 2% of a light one. Disabling it cut the 8 heaviest files by 33% (wall, -j8). Whole-run effect not measured; policy decision..cppfiles across the three test targets via one OBJECT library (estimate: ~23 CPU-min, ~0.4 min wall).hipdnn_ingestor_generator_pytest153-230 s,descriptor_generator_pytest,hkp-pack) are the floor once ctest is parallel (~4 min).Graph.hppis reached by 221 test files (~51 s of tidy each standalone) and is header-only by design; splitting it needs a design discussion.Full write-up and raw data:
reports/tidy_header_diet/REPORT.md(workspace, not in the repo).