Repository navigation
fix(hip-kernel-provider): place shard test configs beside the build-tree census shard - #14573
Merged
GarbyGilbertAMD merged 4 commits intoOct 30, 2026
Conversation
…ree census shard The per-architecture census shard is written into the test descriptor build tree, which is installed wholesale, so the build tree carries the same CTestTestfile.cmake the install does and CTest can run it in place. Engine-pinned external suites routed to a shard name their --test-config TOML by bare name, and only the install tree received that file, so running the shard from the build tree failed with "Config path does not exist" before any case ran. hkp_reserve_census_shard() takes FILES, and the finalizer copies each reserved file beside the shard it writes. Configuration fails when a shard names a TOML that is not beside it, so an engine routed to a shard later cannot reintroduce the gap. The wholesale install skips the Testing/ directory CTest leaves in a shard run in place. The ingestor runbook and the superbuild test skill name the build-tree registration of the engine-pinned suites and the CTest files that only run once installed. ALMIOPEN-2879
github-actions
Bot
added
documentation
project: hipdnn
project: hip-kernel-provider
labels
✅ 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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #14573 +/- ##
===========================================
+ Coverage 76.11% 76.12% +0.01%
===========================================
Files 2886 2886
Lines 483963 484141 +178
Branches 73003 73034 +31
===========================================
+ Hits 368342 368541 +199
+ Misses 90392 90363 -29
- Partials 25229 25237 +8
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
Only a gfx950 configure with the kernel ingestor generates a census shard that names a TOML, so the copy beside the shard and the configure guard were reached by no configuration the superbuild lanes run. A quick cmake -P test drives hkp_reserve_census_shard() and hkp_finalize_census_install() directly: a reserved TOML lands beside the shard byte for byte, a named TOML that nothing places fails configuration in both quoting forms, and an unrecognised argument fails. hkp_reserve_census_shard() rejects unrecognised arguments, as the other public parsers in HkpPackaging.cmake do, so a misspelled FILES keyword fails at the reservation instead of later at the guard or not at all. Its header holds the rationale for the copy and the guard, and the other sites point to it. The guard's comment states that it compares bare names. The ingestor runbook and the superbuild-test skill give the build-tree path of the census shard and the full engine-label rule, which maps every character outside [A-Za-z0-9_.-] to an underscore. ALMIOPEN-2879
The test for an unrecognised hkp_reserve_census_shard() argument staged an entry that names a TOML nothing places, so the shard guard failed the run on its own. A check downgraded to a warning still printed the expected text, and the test still passed. The entry now names no TOML, so only the argument check can fail the run. ALMIOPEN-2879
The census shard comments, the shard test's module docstring and the two skill docs repeated facts that the hkp_reserve_census_shard() header already states. The repeats are removed, the guard's comment on bare names is shortened, and the superbuild-test path definition is rewrapped. ALMIOPEN-2879
BrianHarrisonAMD
approved these changes
Oct 30, 2026
BrianHarrisonAMD
left a comment
BrianHarrisonAMD
left a comment
Contributor
There was a problem hiding this comment.
LGTM thanks for the updates!
GarbyGilbertAMD
deleted the
users/ggilbert/ALMIOPEN-2879/census-shard-test-config
branch
October 30, 2026 17:30
3 of 4 tasks
BrianHarrisonAMD
added a commit
that referenced
this pull request
Oct 30, 2026
## Summary `hipdnn-deep-review` can now continue a review whose earlier rounds exist only as a published record someone else ran. Before this, the next round silently computed an empty set of affected files and could not close while answered questions were still listed as open. JIRA ID : ALMIOPEN-2945 ## Risk Assessment Risk 1. Skill documentation only; no script, product or build change. ## Testing Summary - Ran round 3 of #14573 from its published rounds 1-2 with the new procedure. Every ledger and comments script accepted the seeded copies, and the round closed. - Skill tooling tests and skill validation match `develop`. ## Testing Checklist - [x] Resume a published review - deep review round 3 of #14573 - Status: Passed (`Closure: CLOSED`) - [x] Skill tooling tests - `python3 -m pytest projects/hipdnn/tools/ai/tests` - Status: Passed (333 passed, 7 skipped; 1 failure also present on `develop`, in `hipdnn-ingestor-engine`) - [x] Skill validation - `python3 projects/hipdnn/tools/ai/validate-skills.py` - Status: Passed for this skill (same `hipdnn-ingestor-engine` link error as `develop`) - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - New optional input `prior:<dir>` for a published record of earlier rounds. - RUNBOOK *Iterate* describes how to seed `$EVIDENCE` from that record and settle answered open questions first. - `affected.txt` takes the previous round's files from git, so a round folder without `files.txt` works.
19 of 20 tasks
adickin-amd
added a commit
that referenced
this pull request
Oct 30, 2026
…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.
Merged
3 of 4 tasks
shumway
pushed a commit
that referenced
this pull request
Oct 31, 2026
#14801) ## Summary The HipDNN Windows Superbuild fails at CMake configure on `develop` since #14111 merged: ``` CMake Error at dnn-providers/hip-kernel-provider/descriptor-packaging/cmake/HkpPackaging.cmake:2400 (message): hkp: the gfx1151 census shard names hipkernel_flydsl_rmsnorm_gfx1151.toml, but nothing places it beside the shard, so its entries fail from the build tree. Pass the file to hkp_reserve_census_shard(... FILES <path>) where the entry is routed to the shard. ``` #14573 added that check and moved the gfx950 entry's TOML into `hkp_reserve_census_shard(... FILES ...)`. #14111's last superbuild run predates #14573, and it merged afterwards without a rebase (the required checks are not strict), so its per-arch FlyDSL entries still pass the generated TOML only as `TEST_CONFIG`. Only builds with legacy CTest labels (no `shared/ctest`, as in the Windows superbuild) write the TOML name into the shard text, so Linux is not affected. This PR passes the generated FlyDSL TOML to `hkp_reserve_census_shard(... FILES ...)`, exactly as the gfx950 entry does. JIRA ID: ALMIOPEN-2953 ## Testing - The Windows superbuild on this PR is the check: it configures past the census step. - `descriptor-packaging/tests/test_hkp_census_shard_files.py` (the `FILES` mechanism): 4 passed. - The same configure error currently fails the Windows superbuild on other PRs based on today's `develop` (runs 41135285646, 41134306713, 41136279489). 🤖 Generated with [Claude Code](https://claude.com/claude-code) Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Contributor
TheRock Submodule Bump ActivityNewest first
|
adickin-amd
added a commit
that referenced
this pull request
Nov 3, 2026
… reference executors (#14680) ## 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 under `projects/hipdnn` becomes 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 are `RUN_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 - Parallel ctest and tier selection: this PR's own superbuild CI, run 41136095325 (head 0528e30, with current develop merged in), both lanes green. Linux: `-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_tests` now runs only as its tier suites in the first call. Earlier runs of plain `--parallel 16` (experiment PR #14563 and earlier heads of this PR) also passed, so no parallel flakes in any run. - ctest filtering (review comment): the superbuild checkout did not include `shared/ctest`, whose `TestCategories.cmake` registers 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 out `shared/ctest`: CI registers 238 (Linux) / 244 (Windows) entries, where the tiers repeat one gtest filter (for example `miopen_plugin_tests` and `hip_kernel_provider_tests` four times each), and `hipdnn_golden_data_tests` ran twice, as its tier suites and as an unlabeled raw test (fixed below). `-L '^full$'` runs each binary once (`full` is 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 that `shared/ctest` replaces are covered by their `*_full_suite`. Enabling it also registers the two `shared_ctest_*` tests, which configure a nested CMake project; they needed `CMAKE_GENERATOR=Ninja` (Linux; the runner has no make) and `Python3_ROOT_DIR` set to the superbuild venv (`actions/setup-python` exports its own `Python3_ROOT_DIR`, whose Python has no PyYAML). - Executor refactor: MI250 (gfx90a) Linux machine, preset `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/0 - `hipdnn_integration_tests_unit_tests`: 1115/2/0 -> 1120/2/0 (only additions) - `hip_kernel_provider_integration_tests`: 36/72/0 -> 36/72/0 - The new compiled hip-kernel-provider function ran 35 times during those 36 passing tests (gdb breakpoint count). - New `TestGpuPlanRegistrySignatureKey` fails when equality is mutated to compare only the variant index. - clang-tidy 20.1.0 on all 23 changed `.cpp` files: 0 findings. The build container has no clang-tidy, so the build ran with tidy off and tidy was a separate pass. - Test-name validation: `hipdnn`, `hipdnn-integration-tests`, `miopen-provider` clean. The hip-kernel-provider component reports 5 `...FP16...` names in `TestHipFlash2*` that exist in develop (only visible with `ENABLE_HIP_FLASH2_ENGINE=ON`, default OFF); not touched here, not checked at base. ## Testing Checklist - [x] Build + test lanes on gfx90a - superbuild preset `hipdnn-providers-all` - Status: Passed - [x] clang-tidy on changed files - wheel 20.1.0 - Status: Passed (0 findings) - [x] Parallel ctest + `full` tier selection - this PR's superbuild CI (Linux and Windows) - Status: Passed - [ ] PR CI - remaining GitHub PR checks (2 multi-arch stage builds were still pending when this was written) - Status: Pending ## Technical Changes - `.github/workflows/hipdnn-superbuild-ci.yml`: both lanes check out `shared/ctest` and run `ctest --parallel 16 -L '^full$'` followed by `ctest --parallel 16 -LE '<tier labels>'` (see Testing for why); the test step sets `Python3_ROOT_DIR` to the venv (both lanes) and `CMAKE_GENERATOR=Ninja` (Linux) for the nested `shared_ctest_*` configure. - `dnn-providers/integration-tests/CMakeLists.txt`: the YAML categorization block moves above the `hipdnn_golden_data_tests` registration. `_add_test_target_internal()` skips the raw `add_test()` only while `DNN_PROVIDER_TEST_CATEGORY_YAMLS` is already set, so the binary was also registered as an unlabeled raw test (about 2 s on Linux, 8 s on Windows). Without `shared/ctest` the 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 #14573 (shard TOML check) and #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) and `GpuPlanRegistrySignatureKey.hpp`: equality compares `index()` then runs one `std::visit` (was a pairwise visit: 22x22 and 15x15 instantiations). Semantics unchanged. - `ReferenceGraphExecutorFactory`, `CpuReferenceGraphExecutorAdapter` and `GpuReferenceGraphExecutor` move from inline headers to pimpl + `.cpp` in `dnn-providers/integration-tests`; 12 plan tests and the chained SDPA test use the adapter instead of the raw test_sdk executor. `ReferenceExecutorPool.cpp` stays out of the unit-test target. - hip-kernel-provider harness: the single CPU-reference call goes through a compiled `executeCpuReferenceGraph`. - Tests: new `TestGpuPlanRegistrySignatureKey`; a real adapter case in `TestErrorPaths` (the old test never called the adapter and is renamed to say what it checks). ## Measured effect | | Before | After | |---|---|---| | Linux ctest step | 7.7-11.3 min | 4.7 min (245 s + 37 s, two calls; this PR's CI, with more tests from develop) | | Windows ctest step | 14.1-19.7 min | 7.4 min (349 s + 92 s, two calls; 249 tests incl. 56 FlyDSL entries added by develop) | | Tidy of the 266 affected files, 64 jobs (MI250 Linux machine) | 334.8 CPU-min, 539.7 s wall | 285.6 CPU-min, 536.3 s wall | Be honest about the executor refactor: it cut tidy CPU on the files it touched by 28% (e.g. `ReferenceExecutorPool.cpp` 465 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 `.cpp` files. 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: 1. **Take clang-tidy out of the Linux build step** (it is inline on every compile). Linux build 33 min -> 6.8 min, and the whole Linux superbuild went from 47 to 14 min together with parallel ctest (experiment PR, same EPYC 7763 runner class); tidy then runs as its own job (~25 min full, so scope it to changed files or per-component targets to avoid a new long pole). Caveat found: applying hipDNN's single `.clang-tidy` to every compile-database entry gives 338k false findings on rocKE; per-component configs are needed for parity. 2. **`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. 3. **Share the new `.cpp` files across the three test targets** via one OBJECT library (estimate: ~23 CPU-min, ~0.4 min wall). 4. **Runner variance**: the same Linux build took 26.5 / 32.9 / 37.4 min depending on the host CPU (EPYC 9V74 / 7763 / 7452). 5. **Long Python suites** (`hipdnn_ingestor_generator_pytest` 153-230 s, `descriptor_generator_pytest`, `hkp-pack`) are the floor once ctest is parallel (~4 min). 6. **Frontend `Graph.hpp`** is reached by 221 test files (~51 s of tidy each standalone) and is header-only by design; splitting it needs a design discussion. 7. Not worth doing: include-only cleanup (87 files + 13 headers edited, 0.0 to -0.5% within noise). Full write-up and raw data: `reports/tidy_header_diet/REPORT.md` (workspace, not in the repo).
cgmillette
added a commit
that referenced
this pull request
Nov 4, 2026
…r engines (#14242) ## Summary Packages rocKE's gfx950 unified attention as two hipDNN kernel-ingestor engines: `hipkernel:Gfx950AttentionUnified2D` (prefill and the decode shapes 3D does not cover) and `hipkernel:Gfx950AttentionUnified3D` (split-KV decode, a segment and a reduce launch per kernel through #14158's multi-launch `sources`). Both ship AOT, so an SDPA-forward graph on gfx950 can be served by unified attention with no JIT. Both catalogs are reduced to what a correctness-checked MI355X sweep measured as needed: Unified2D ships 222 of its 792-kernel full space and Unified3D 123 of 348, with no graph losing coverage or more than the measured noise (the final reduced-vs-full sweep is below). JIRA ID : AICK-2542 JIRA ID : AICK-2682 ## Risk Assessment Risk 2 of 5. The PR adds two engines and their descriptors. No existing engine's code, catalog or ranking changes: an unpinned SdpaFwd graph still ranks `hipkernel:Gfx950AttentionDense` first (device test `DenseStillRanksAheadWhenUnpinned`). The shared pieces it touches: - generator tooling: `registry_enumerator` full-space and multi-launch sections, the regeneration gates, provenance stamping, and two fixes in `variant_reduce`. Only `provenance.generated_from` is opt-in. The expanded registry output now always writes per-kernel `provenance.rocke_spec_id` / `rocke_tuning_id`, including Dense's 840 catalog kernels (its manifest pins a `spec_id`). The compact `variants:` selections drop it, so no shipped descriptor changes, but any other consumer of the expanded output sees the new field; - one rocKE builder, `build_unified_attention_3d_tiled_module`; - the ingestor save census's list of dispatch symbols. These are covered by their own tests and by the existing dense regeneration gates, which still pass. 3D is the first production user of `launchNames()` (#14158). ## ASIC Coverage gfx950 only: both engines ship only for gfx950, and other architectures' shards carry neither engine. PR CI has no gfx950 device lane (`hkp-probe-gfx950` packs and asserts but launches nothing), so the device evidence comes from MI355X (and, before the rebase, MI350X) runs. ## Testing Summary - Host: unit and census tests for both packs, the arch-neutral geometry, and the generator suite under `hipdnn_strict_skips`, including regeneration gates g1–g7 for all three gfx950 manifests; static descriptor gates on the installed shards (variant sets, coverage rungs 1–2, validator, launch surface). - Device, MI355X (gfx950), on the shipped (reduced) catalogs at this head's code: - engine-pinned SdpaFwd quick and standard suites against the CPU and GPU references, both engines plus a Dense regression run; - the knobs suite: every shipped kernel class forced and checked against the CPU reference, plus declines, precedence and a NaN-prefilled workspace; - a timing comparison of every engine that serves each corpus graph, and PyTorch; - RUNBOOK stage 7's final sweep: the full catalog against the reduced one in one session. - How the catalogs were reduced: a 3-round dnn-benchmark sweep with correctness on, over 188 prefill graphs (2D) and 88 decode graphs (3D), then `variant_reduce`. ## Packaging and performance statistics All numbers below come from tool output (packaging logs, device suite runs, dnn-benchmark sweeps). **Packaging** (gfx950 shard, shipped catalogs; pack time is the build's own `hkp_pack` command per engine, median of 3, relative to Dense's on the same job) | Engine | Kernels (full space) | AOT budget | Code objects, zstd | KDP | Pack time vs Dense, build default workers | Serial time per kernel vs Dense | |---|---:|---:|---:|---:|---:|---:| | Unified2D | 222 (792) | 1000 | 1.55 MiB | 0.29 MB | 0.11× | 1.58× | | Unified3D | 123 (348), 246 launches | 400 | 0.64 MiB | 0.26 MB | 0.41× | 1.06× | | Dense (catalog unchanged, for scale) | 840 (52 distinct binaries since #14059) | 1000 | 6.16 MiB | 1.26 MB | 1 | 1 | - The gfx950 kpack is 8.49 MiB at this head; the two engines are 2.19 MiB, 26% of it. Since #14059 (runtime head counts for Dense), Dense's 840 catalog kernels compile to 52 distinct binaries, but the kpack still stores one frame per kernel, so Dense grew from 5.79 to 6.16 MiB. Collapsing those entries is #14059's stated follow-up. The pack-time ratios and the full-space figure (14.45 MiB) were measured against Dense before #14059. - Provider load (`hipdnn_list_engines`, median of 5) is 1.24× its time without these descriptors. **Correctness and coverage** (MI355X, shipped catalogs) | Suite | Unified2D | Unified3D | Dense | |---|---|---|---| | SdpaFwd quick + standard, CPU and GPU reference (served / declined / failed) | 43 / 6 / 0 | 4 / 45 / 0 | 45 / 4 / 0 | | Knobs suite (forced kernel classes, declines, precedence) | 30 / 30 (both engines) | | 62 / 62 | | Corpus graphs served and validated against PyTorch | 188 / 188 | 82 / 88 | | Every decline is intentional: Unified2D leaves decode to Unified3D and declines Hkv > Hq; Unified3D serves decode only, and its 6 corpus declines are batch-256 decode graphs whose K/V is over the 32-bit addressing cap (ASM SDPA serves those). **Kernel time** (MI355X, autotuned, median of 3 rounds; ratio = Unified / other, below 1 = Unified faster; every engine's round-to-round drift ≤ 0.4%). The Unified2D rows were re-measured at this head, on the kernels develop's tiled-2D changes recompiled (see the checklist). | | Graphs both serve | Geomean ratio | Time-weighted (other / Unified) | Unified faster on | |---|---:|---:|---:|---:| | Unified3D vs ASM SDPA (decode) | 60 | 0.200 | 8.17 | 60 / 60 | | Unified3D vs PyTorch | 82 | 0.062 | 34.2 | 82 / 82 | | Unified2D vs Dense (prefill) | 179 | 2.37 | 0.38 | 0 / 179 | | Unified2D vs PyTorch | 188 | 1.38 | 0.55 | 39 / 188 | - Unified3D serves 22 decode graphs no existing hipDNN engine serves. - Unified2D is slower than Dense wherever Dense serves, and Dense keeps ranking first there; it adds coverage Dense lacks (the 9 bottom-right causal graphs in these corpora). **Variant reduction** | Engine | Kernels | Chosen eps | Measured drift | Final sweep, reduced / full kernel time (geomean, rounds 1–3) | |---|---|---:|---:|---| | Unified2D | 792 → 222 | 0.02 | 0.62% | 1.000, 0.999, 1.000 (u2d_attn); 1.000, 0.999, 1.000 (u2d_rocke) | | Unified3D | 348 → 123 | 0.03 | 1.08% | 1.004, 1.002, 0.999 (u3d_attn); 1.002, 1.007, 1.004 (u3d_rocke) | Both arms of the final sweep served the same graphs and passed correctness on every corpus. ## Testing Checklist - [x] Unit (host, no GPU) - `HIP_VISIBLE_DEVICES=-1 hip_kernel_provider_tests` - 1454 passed, including `TestKernelIngestorEngine.EveryHandlerReportsSupportsSavingExactlyWhenItSavesAndRestores`. Our build host also fails 11 Dense signature/dispatch tests: it has a `/nonexistent` directory, which those tests assume is absent. They pass in CI - Status: Passed - [x] Census - `ctest -R 'hkp-census|rocke_unified'` - 36/36 (the census inventories are the reduced catalogs') - Status: Passed - [x] Generator suite (strict skips) - `pytest -p hipdnn_strict_skips IngestorGenerator/tests` - 1646 passed, 0 skipped - Status: Passed - [x] Regeneration gates on the reduced selections - `tests/test_registry_regeneration.py` (g1–g7) 39 passed; AOT budget tests 12 passed - Status: Passed - [x] rocKE attention CPU tests and packaging guards - 2D/3D mask_type, 3D module, decode split, KV cache policy, spec-predicate parity: 145 passed, plus 6 develop KV-cache tests that need the compiled `rocke_engine`; `test_hkp_pack_producer_guards.py` and `test_gfx950_attention_unified_3d_bundle_coverage.py`: 26 passed - Status: Passed - [x] Golden IR, before the reduction - `ctest -R rocke_golden_static` - the reduction changes no rocKE code - Status: Passed - [x] Launch surface - `launch_surface.py configs/gfx950_attention_unified_{2d,3d}.engine.yaml --check` - Status: Passed - [x] Variant sets - `verify_variant_sets.py --mode structural` and `--mode full --arch gfx950 --manifest …` - 2D 222/222 distinct binaries, 3D 123/123 - Status: Passed - [x] Coverage gate rungs 1–2 - `coverage_gate.py --mode full --expect-engine …` - Status: Passed - [x] Validator - `hipdnn_validate_descriptors <shard> --expect-engine … --json` - success, no warnings - Status: Passed - [x] Packaging probe - gfx942 configure with `HIPKERNELPROVIDER_ENABLE_PACKAGING_PROBES=ON`, `ctest -R '^hkp-probe-gfx950'` - 12/12; `hkp-probe-gfx950` packs 51 UKDs of the current descriptors in a gfx942-only build - Status: Passed - [x] Device suites, MI355X - engine-pinned SdpaFwd quick + standard with GPU reference (`hip_kernel_provider_gfx950_attention_*_gpu_ref_integration_tests`) and CPU reference (`*-external-integration_{quick,standard}_suite`) - 0 failures: Unified2D 43 served / 6 declined, Unified3D 4 / 45, Dense 45 / 4 - Status: Passed - [x] Device knobs suite, MI355X - `hip_kernel_provider_integration_tests --gtest_filter='*Gfx950AttentionUnified*'` - 30/30; Dense knobs 62/62 - Status: Passed - [x] Fault re-check, MI355X - the whole `u2d_rocke` corpus in one autotuned process, twice and then with serialized launches - 93/93 each, 0 errors - Status: Passed - [x] Reduction sweep + `variant_reduce.py` - 3 rounds, correctness on; 2D `SWEEP_DONE`, 3D `SWEEP_DONE`; 2D 792 → 222, 3D 348 → 123 - Status: Passed - [x] Final sweep, reduced against full (RUNBOOK §7) - same coverage, correctness passed on both arms, ratios above - Status: Passed - [x] Unified3D catalog regenerated for #14367's LLC-aware K/V cache policy - all 123 kernels now `kv_cache_policy: "all"` (was `"stream"`), same kernel classes; MI355X re-run: suites and knobs unchanged, 246/246 decode rows correct, kernel time against the `"stream"` catalog geomean 0.998, time-weighted 0.999 - Status: Passed - [x] Device evidence at this head (rebased onto develop 2026-11-03, after the review-round refactors) - Dense's and Unified3D's code objects are byte-identical to the previous device runs; develop's #13652 recompiled 201 of the 222 Unified2D code objects (same kernels and ids), and the Unified2D/3D matchers moved onto a shared header, so both engines' device runs were repeated at this head: suites Unified2D 43 / 6 / 0, Unified3D 4 / 45 / 0, Dense 45 / 4 / 0 (passed / skipped / failed), knobs 30/30 and 62/62; compare 188/188 Unified2D graphs (564/564 rows) and 82 Unified3D graphs (246/246 rows) correct against PyTorch in all 3 rounds. Kernel time vs the previous runs: Unified2D geomean 0.985 (none of 188 graphs >2% slower), Unified3D 1.003 (ASM_SDPA control 1.006). A final rebase picked up #14059, which recompiled Dense's 840 code objects only (Unified2D/3D byte-identical to the run above); the suites and knobs were re-run at this head with the same results, `DenseStillRanksAheadWhenUnpinned` included - Status: Passed - [x] DVC - both reduced KDPs on the `ingestor` remote; `dvc status -c` in sync - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - **Unified2D pack** (`Gfx950AttentionUnified2DNative.cpp`): matches BSHD SDPA forward on the identity paged-KV layout and builds block_tables, seq_lens and query_start_len in the workspace. - It declines: batch > 511 (the baked 9-iteration binary search), Skv not a multiple of 16, i32 KV over 0x7FFF0000, any window, and the decode cells Unified3D covers (Sq = 1 with the GQA group dividing 16). - Engine knobs: tile_size, num_warps, block_m_per_warp, use_mfma_32x32 and use_transposed_qk_32x32. - **Unified3D pack** (`Gfx950AttentionUnified3DNative.cpp`): decode only. - `launchNames()` is `{segment, reduce}`, and both definitions are resolved through `launchKernelDefinition` with a signature check before any load. - The segment launch writes `[token, Hq, seg, D]` scratch and the reduce always runs, so the workspace needs no zeroing. - Knob: num_segments. No batch limit. - **Arch-neutral geometry** (`UnifiedAttentionGeometry.hpp`): the launch geometry and workspace layouts take a `UnifiedAttentionArchTraits`, and the gfx950 headers are thin wrappers, so a gfx942 pack (AICK-2544) only supplies traits. - **rocKE**: `build_unified_attention_3d_tiled_module` packs the segment and its derived reduce (`reduce_spec_for`) as one `KernelModule`. The module predicate gates both halves. The 3D pack compiles develop's running-max fix for split-KV (#14276). - **Generator**: - `enumerate.also` sections may use `space: full` and carry manifest `launches`. The enumerator refuses a registry that pairs a reduce when the manifest declares no launches, and keeps one default per cell. - The gates are launch-aware: g1 keys on launches, g3 uses `--subset` for full spaces, g4 compares launch sets, and g6 walks every launch and checks module roles. - g2 and g3 run per engine with the shipped ids reused (#14109), comparing every descriptor byte for byte. - `generate.py --stamp-provenance` stamps `provenance.generated_from` (config, manifest, generator and rocKE commits). It is opt-in, so a default regeneration stays byte-identical. - `variant_reduce` fixes: - it reads the format-2 winner lines the runtime has written since #13028; - it keys a graph on its content with `Graph.id` (minted fresh at each finalization) zeroed, so sweep rounds match. - **Catalogs** (the manifests enumerate the full spaces; the selections are `variant_reduce`'s): - Unified2D: 222 kernels over 198 cells (33 head configs × bf16/fp16 × 3 masks): the narrow nw4/mw16/t64 geometry in every cell, transposed 32x32 nw4/mw32/t64 in 21 and narrow nw2/mw16/t32 in 3. The wide 32x32 geometry won no graph and is not shipped. - Unified3D: 123 kernels: num_segments 16 in all 116 decode cells, 32 or 64 in 7 bf16 bottom-right D128 variants. The K/V cache policy follows #14367's LLC-aware rule, which picks `"all"` for every shipped decode cell. - **Descriptors**: both are generated with develop's generator (post-#13028 UED/UHD, UKD `sources` for 3D), every kernel id reused from the full catalog. A regeneration with the shipped ids reused reproduces them byte for byte. The KDPs live in DVC (`ingestor` remote). - **Shared code** (one home per fact): - `packs/UnifiedAttentionGraphMatch.hpp`: the graph-match helpers both unified packs use (operand/BSHD checks, SDPA node, mask classification, dtype, metadata readers); the tests' graph builder is `UnifiedAttentionTestGraphs.hpp`. - rocKE `kernels/common/attention_mask_type.py` (mask ordinals and band -> kind, used by both tiled specs and re-exported by dispatch) and `attention_split_kv_reduce.py` (the segment -> reduce derivation the dispatcher and the module export share). - CMake registers Dense, Unified2D and Unified3D by one loop per lane (external integration, gpu_ref, census). - **Tests**: - The device knobs suite (`IntegrationGpuGfx950AttentionUnifiedKnobs.cpp`) adds a harness `prepareWorkspace` hook. It forces the default geometry/segment count across the head configurations and each other shipped class on a cell that carries it, and `CoversEveryCatalogKernelClass` holds those tables to the catalog. - The engine-pinned `gpu_ref` targets and external-integration lanes cover both engines. - The ingestor save census names all seven pack dispatch symbols. - The Unified2D/3D integration TOMLs are reserved with the gfx950 census shard (`hkp_reserve_census_shard(... FILES ...)`, as #14573 requires of every shard entry). **Constraints and follow-ups** - #13372 and #13889 are both in develop; this branch no longer carries a copy of #13889. - Heads are catalog cells (baked), window = 0, 2D batch ≤ 511, i32 KV ≤ 0x7FFF0000. D256 is deferred for both engines. The kernarg work is tracked in AICK-2675. - There is no 2D/3D `select_path` crossover replay: the split is the static decode rule above. - Runtimes older than UKD 1.1 skip the 3D descriptors. Strided-KV 3D is a follow-up. - Follow-up: execution-plan save/restore (#13315) for both handlers. Today they report `supportsSaving() == false`, as the RDNA handler does. - One earlier MI355X sweep hit a single GPU fault in Unified2D on `u2d_rocke/rocke_sdpa_035` (bf16, D 128, Hq 64 / Hkv 8, causal, S 512). It has not reproduced since: not in the forced geometries on that cell, 15 single-graph runs, 4 full-corpus runs on MI350X, or the MI355X full-corpus re-check above. - Unified2D kernel time moved with develop's tiled-2D changes (#14542, #14616), not with this PR: at this head it is 1.035× (geomean) its pre-rebase time on the same 188 graphs, up to 1.36× on long causal bf16 D128 prefill (S 8192), while Dense in the same runs is flat (0.997). Dense still ranks first wherever it serves; the table above is at this head. - Measurement note: the pinned dnn-benchmarking (v1 result schema, which `sweep.py` reads) has no bottom-right causal PyTorch reference. The sweeps used upstream dnn-benchmarking #84's reference, with the 9 bottom-right corpus graphs written as `causal_mask: false, right_bound: 0, diagonal_alignment: BOTTOM_RIGHT` (the same band under hipDNN's `resolveDiagonalBand`): #84 reads `causal_mask: true` as top-left regardless of the alignment. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Christopher Millette <63608002+cgmillette@users.noreply.github.com>
cycheng
pushed a commit
that referenced
this pull request
Nov 6, 2026
…r engines (#14242) ## Summary Packages rocKE's gfx950 unified attention as two hipDNN kernel-ingestor engines: `hipkernel:Gfx950AttentionUnified2D` (prefill and the decode shapes 3D does not cover) and `hipkernel:Gfx950AttentionUnified3D` (split-KV decode, a segment and a reduce launch per kernel through #14158's multi-launch `sources`). Both ship AOT, so an SDPA-forward graph on gfx950 can be served by unified attention with no JIT. Both catalogs are reduced to what a correctness-checked MI355X sweep measured as needed: Unified2D ships 222 of its 792-kernel full space and Unified3D 123 of 348, with no graph losing coverage or more than the measured noise (the final reduced-vs-full sweep is below). JIRA ID : AICK-2542 JIRA ID : AICK-2682 ## Risk Assessment Risk 2 of 5. The PR adds two engines and their descriptors. No existing engine's code, catalog or ranking changes: an unpinned SdpaFwd graph still ranks `hipkernel:Gfx950AttentionDense` first (device test `DenseStillRanksAheadWhenUnpinned`). The shared pieces it touches: - generator tooling: `registry_enumerator` full-space and multi-launch sections, the regeneration gates, provenance stamping, and two fixes in `variant_reduce`. Only `provenance.generated_from` is opt-in. The expanded registry output now always writes per-kernel `provenance.rocke_spec_id` / `rocke_tuning_id`, including Dense's 840 catalog kernels (its manifest pins a `spec_id`). The compact `variants:` selections drop it, so no shipped descriptor changes, but any other consumer of the expanded output sees the new field; - one rocKE builder, `build_unified_attention_3d_tiled_module`; - the ingestor save census's list of dispatch symbols. These are covered by their own tests and by the existing dense regeneration gates, which still pass. 3D is the first production user of `launchNames()` (#14158). ## ASIC Coverage gfx950 only: both engines ship only for gfx950, and other architectures' shards carry neither engine. PR CI has no gfx950 device lane (`hkp-probe-gfx950` packs and asserts but launches nothing), so the device evidence comes from MI355X (and, before the rebase, MI350X) runs. ## Testing Summary - Host: unit and census tests for both packs, the arch-neutral geometry, and the generator suite under `hipdnn_strict_skips`, including regeneration gates g1–g7 for all three gfx950 manifests; static descriptor gates on the installed shards (variant sets, coverage rungs 1–2, validator, launch surface). - Device, MI355X (gfx950), on the shipped (reduced) catalogs at this head's code: - engine-pinned SdpaFwd quick and standard suites against the CPU and GPU references, both engines plus a Dense regression run; - the knobs suite: every shipped kernel class forced and checked against the CPU reference, plus declines, precedence and a NaN-prefilled workspace; - a timing comparison of every engine that serves each corpus graph, and PyTorch; - RUNBOOK stage 7's final sweep: the full catalog against the reduced one in one session. - How the catalogs were reduced: a 3-round dnn-benchmark sweep with correctness on, over 188 prefill graphs (2D) and 88 decode graphs (3D), then `variant_reduce`. ## Packaging and performance statistics All numbers below come from tool output (packaging logs, device suite runs, dnn-benchmark sweeps). **Packaging** (gfx950 shard, shipped catalogs; pack time is the build's own `hkp_pack` command per engine, median of 3, relative to Dense's on the same job) | Engine | Kernels (full space) | AOT budget | Code objects, zstd | KDP | Pack time vs Dense, build default workers | Serial time per kernel vs Dense | |---|---:|---:|---:|---:|---:|---:| | Unified2D | 222 (792) | 1000 | 1.55 MiB | 0.29 MB | 0.11× | 1.58× | | Unified3D | 123 (348), 246 launches | 400 | 0.64 MiB | 0.26 MB | 0.41× | 1.06× | | Dense (catalog unchanged, for scale) | 840 (52 distinct binaries since #14059) | 1000 | 6.16 MiB | 1.26 MB | 1 | 1 | - The gfx950 kpack is 8.49 MiB at this head; the two engines are 2.19 MiB, 26% of it. Since #14059 (runtime head counts for Dense), Dense's 840 catalog kernels compile to 52 distinct binaries, but the kpack still stores one frame per kernel, so Dense grew from 5.79 to 6.16 MiB. Collapsing those entries is #14059's stated follow-up. The pack-time ratios and the full-space figure (14.45 MiB) were measured against Dense before #14059. - Provider load (`hipdnn_list_engines`, median of 5) is 1.24× its time without these descriptors. **Correctness and coverage** (MI355X, shipped catalogs) | Suite | Unified2D | Unified3D | Dense | |---|---|---|---| | SdpaFwd quick + standard, CPU and GPU reference (served / declined / failed) | 43 / 6 / 0 | 4 / 45 / 0 | 45 / 4 / 0 | | Knobs suite (forced kernel classes, declines, precedence) | 30 / 30 (both engines) | | 62 / 62 | | Corpus graphs served and validated against PyTorch | 188 / 188 | 82 / 88 | | Every decline is intentional: Unified2D leaves decode to Unified3D and declines Hkv > Hq; Unified3D serves decode only, and its 6 corpus declines are batch-256 decode graphs whose K/V is over the 32-bit addressing cap (ASM SDPA serves those). **Kernel time** (MI355X, autotuned, median of 3 rounds; ratio = Unified / other, below 1 = Unified faster; every engine's round-to-round drift ≤ 0.4%). The Unified2D rows were re-measured at this head, on the kernels develop's tiled-2D changes recompiled (see the checklist). | | Graphs both serve | Geomean ratio | Time-weighted (other / Unified) | Unified faster on | |---|---:|---:|---:|---:| | Unified3D vs ASM SDPA (decode) | 60 | 0.200 | 8.17 | 60 / 60 | | Unified3D vs PyTorch | 82 | 0.062 | 34.2 | 82 / 82 | | Unified2D vs Dense (prefill) | 179 | 2.37 | 0.38 | 0 / 179 | | Unified2D vs PyTorch | 188 | 1.38 | 0.55 | 39 / 188 | - Unified3D serves 22 decode graphs no existing hipDNN engine serves. - Unified2D is slower than Dense wherever Dense serves, and Dense keeps ranking first there; it adds coverage Dense lacks (the 9 bottom-right causal graphs in these corpora). **Variant reduction** | Engine | Kernels | Chosen eps | Measured drift | Final sweep, reduced / full kernel time (geomean, rounds 1–3) | |---|---|---:|---:|---| | Unified2D | 792 → 222 | 0.02 | 0.62% | 1.000, 0.999, 1.000 (u2d_attn); 1.000, 0.999, 1.000 (u2d_rocke) | | Unified3D | 348 → 123 | 0.03 | 1.08% | 1.004, 1.002, 0.999 (u3d_attn); 1.002, 1.007, 1.004 (u3d_rocke) | Both arms of the final sweep served the same graphs and passed correctness on every corpus. ## Testing Checklist - [x] Unit (host, no GPU) - `HIP_VISIBLE_DEVICES=-1 hip_kernel_provider_tests` - 1454 passed, including `TestKernelIngestorEngine.EveryHandlerReportsSupportsSavingExactlyWhenItSavesAndRestores`. Our build host also fails 11 Dense signature/dispatch tests: it has a `/nonexistent` directory, which those tests assume is absent. They pass in CI - Status: Passed - [x] Census - `ctest -R 'hkp-census|rocke_unified'` - 36/36 (the census inventories are the reduced catalogs') - Status: Passed - [x] Generator suite (strict skips) - `pytest -p hipdnn_strict_skips IngestorGenerator/tests` - 1646 passed, 0 skipped - Status: Passed - [x] Regeneration gates on the reduced selections - `tests/test_registry_regeneration.py` (g1–g7) 39 passed; AOT budget tests 12 passed - Status: Passed - [x] rocKE attention CPU tests and packaging guards - 2D/3D mask_type, 3D module, decode split, KV cache policy, spec-predicate parity: 145 passed, plus 6 develop KV-cache tests that need the compiled `rocke_engine`; `test_hkp_pack_producer_guards.py` and `test_gfx950_attention_unified_3d_bundle_coverage.py`: 26 passed - Status: Passed - [x] Golden IR, before the reduction - `ctest -R rocke_golden_static` - the reduction changes no rocKE code - Status: Passed - [x] Launch surface - `launch_surface.py configs/gfx950_attention_unified_{2d,3d}.engine.yaml --check` - Status: Passed - [x] Variant sets - `verify_variant_sets.py --mode structural` and `--mode full --arch gfx950 --manifest …` - 2D 222/222 distinct binaries, 3D 123/123 - Status: Passed - [x] Coverage gate rungs 1–2 - `coverage_gate.py --mode full --expect-engine …` - Status: Passed - [x] Validator - `hipdnn_validate_descriptors <shard> --expect-engine … --json` - success, no warnings - Status: Passed - [x] Packaging probe - gfx942 configure with `HIPKERNELPROVIDER_ENABLE_PACKAGING_PROBES=ON`, `ctest -R '^hkp-probe-gfx950'` - 12/12; `hkp-probe-gfx950` packs 51 UKDs of the current descriptors in a gfx942-only build - Status: Passed - [x] Device suites, MI355X - engine-pinned SdpaFwd quick + standard with GPU reference (`hip_kernel_provider_gfx950_attention_*_gpu_ref_integration_tests`) and CPU reference (`*-external-integration_{quick,standard}_suite`) - 0 failures: Unified2D 43 served / 6 declined, Unified3D 4 / 45, Dense 45 / 4 - Status: Passed - [x] Device knobs suite, MI355X - `hip_kernel_provider_integration_tests --gtest_filter='*Gfx950AttentionUnified*'` - 30/30; Dense knobs 62/62 - Status: Passed - [x] Fault re-check, MI355X - the whole `u2d_rocke` corpus in one autotuned process, twice and then with serialized launches - 93/93 each, 0 errors - Status: Passed - [x] Reduction sweep + `variant_reduce.py` - 3 rounds, correctness on; 2D `SWEEP_DONE`, 3D `SWEEP_DONE`; 2D 792 → 222, 3D 348 → 123 - Status: Passed - [x] Final sweep, reduced against full (RUNBOOK §7) - same coverage, correctness passed on both arms, ratios above - Status: Passed - [x] Unified3D catalog regenerated for #14367's LLC-aware K/V cache policy - all 123 kernels now `kv_cache_policy: "all"` (was `"stream"`), same kernel classes; MI355X re-run: suites and knobs unchanged, 246/246 decode rows correct, kernel time against the `"stream"` catalog geomean 0.998, time-weighted 0.999 - Status: Passed - [x] Device evidence at this head (rebased onto develop 2026-11-03, after the review-round refactors) - Dense's and Unified3D's code objects are byte-identical to the previous device runs; develop's #13652 recompiled 201 of the 222 Unified2D code objects (same kernels and ids), and the Unified2D/3D matchers moved onto a shared header, so both engines' device runs were repeated at this head: suites Unified2D 43 / 6 / 0, Unified3D 4 / 45 / 0, Dense 45 / 4 / 0 (passed / skipped / failed), knobs 30/30 and 62/62; compare 188/188 Unified2D graphs (564/564 rows) and 82 Unified3D graphs (246/246 rows) correct against PyTorch in all 3 rounds. Kernel time vs the previous runs: Unified2D geomean 0.985 (none of 188 graphs >2% slower), Unified3D 1.003 (ASM_SDPA control 1.006). A final rebase picked up #14059, which recompiled Dense's 840 code objects only (Unified2D/3D byte-identical to the run above); the suites and knobs were re-run at this head with the same results, `DenseStillRanksAheadWhenUnpinned` included - Status: Passed - [x] DVC - both reduced KDPs on the `ingestor` remote; `dvc status -c` in sync - Status: Passed - [ ] PR CI - GitHub PR checks - Status: Pending ## Technical Changes - **Unified2D pack** (`Gfx950AttentionUnified2DNative.cpp`): matches BSHD SDPA forward on the identity paged-KV layout and builds block_tables, seq_lens and query_start_len in the workspace. - It declines: batch > 511 (the baked 9-iteration binary search), Skv not a multiple of 16, i32 KV over 0x7FFF0000, any window, and the decode cells Unified3D covers (Sq = 1 with the GQA group dividing 16). - Engine knobs: tile_size, num_warps, block_m_per_warp, use_mfma_32x32 and use_transposed_qk_32x32. - **Unified3D pack** (`Gfx950AttentionUnified3DNative.cpp`): decode only. - `launchNames()` is `{segment, reduce}`, and both definitions are resolved through `launchKernelDefinition` with a signature check before any load. - The segment launch writes `[token, Hq, seg, D]` scratch and the reduce always runs, so the workspace needs no zeroing. - Knob: num_segments. No batch limit. - **Arch-neutral geometry** (`UnifiedAttentionGeometry.hpp`): the launch geometry and workspace layouts take a `UnifiedAttentionArchTraits`, and the gfx950 headers are thin wrappers, so a gfx942 pack (AICK-2544) only supplies traits. - **rocKE**: `build_unified_attention_3d_tiled_module` packs the segment and its derived reduce (`reduce_spec_for`) as one `KernelModule`. The module predicate gates both halves. The 3D pack compiles develop's running-max fix for split-KV (#14276). - **Generator**: - `enumerate.also` sections may use `space: full` and carry manifest `launches`. The enumerator refuses a registry that pairs a reduce when the manifest declares no launches, and keeps one default per cell. - The gates are launch-aware: g1 keys on launches, g3 uses `--subset` for full spaces, g4 compares launch sets, and g6 walks every launch and checks module roles. - g2 and g3 run per engine with the shipped ids reused (#14109), comparing every descriptor byte for byte. - `generate.py --stamp-provenance` stamps `provenance.generated_from` (config, manifest, generator and rocKE commits). It is opt-in, so a default regeneration stays byte-identical. - `variant_reduce` fixes: - it reads the format-2 winner lines the runtime has written since #13028; - it keys a graph on its content with `Graph.id` (minted fresh at each finalization) zeroed, so sweep rounds match. - **Catalogs** (the manifests enumerate the full spaces; the selections are `variant_reduce`'s): - Unified2D: 222 kernels over 198 cells (33 head configs × bf16/fp16 × 3 masks): the narrow nw4/mw16/t64 geometry in every cell, transposed 32x32 nw4/mw32/t64 in 21 and narrow nw2/mw16/t32 in 3. The wide 32x32 geometry won no graph and is not shipped. - Unified3D: 123 kernels: num_segments 16 in all 116 decode cells, 32 or 64 in 7 bf16 bottom-right D128 variants. The K/V cache policy follows #14367's LLC-aware rule, which picks `"all"` for every shipped decode cell. - **Descriptors**: both are generated with develop's generator (post-#13028 UED/UHD, UKD `sources` for 3D), every kernel id reused from the full catalog. A regeneration with the shipped ids reused reproduces them byte for byte. The KDPs live in DVC (`ingestor` remote). - **Shared code** (one home per fact): - `packs/UnifiedAttentionGraphMatch.hpp`: the graph-match helpers both unified packs use (operand/BSHD checks, SDPA node, mask classification, dtype, metadata readers); the tests' graph builder is `UnifiedAttentionTestGraphs.hpp`. - rocKE `kernels/common/attention_mask_type.py` (mask ordinals and band -> kind, used by both tiled specs and re-exported by dispatch) and `attention_split_kv_reduce.py` (the segment -> reduce derivation the dispatcher and the module export share). - CMake registers Dense, Unified2D and Unified3D by one loop per lane (external integration, gpu_ref, census). - **Tests**: - The device knobs suite (`IntegrationGpuGfx950AttentionUnifiedKnobs.cpp`) adds a harness `prepareWorkspace` hook. It forces the default geometry/segment count across the head configurations and each other shipped class on a cell that carries it, and `CoversEveryCatalogKernelClass` holds those tables to the catalog. - The engine-pinned `gpu_ref` targets and external-integration lanes cover both engines. - The ingestor save census names all seven pack dispatch symbols. - The Unified2D/3D integration TOMLs are reserved with the gfx950 census shard (`hkp_reserve_census_shard(... FILES ...)`, as #14573 requires of every shard entry). **Constraints and follow-ups** - #13372 and #13889 are both in develop; this branch no longer carries a copy of #13889. - Heads are catalog cells (baked), window = 0, 2D batch ≤ 511, i32 KV ≤ 0x7FFF0000. D256 is deferred for both engines. The kernarg work is tracked in AICK-2675. - There is no 2D/3D `select_path` crossover replay: the split is the static decode rule above. - Runtimes older than UKD 1.1 skip the 3D descriptors. Strided-KV 3D is a follow-up. - Follow-up: execution-plan save/restore (#13315) for both handlers. Today they report `supportsSaving() == false`, as the RDNA handler does. - One earlier MI355X sweep hit a single GPU fault in Unified2D on `u2d_rocke/rocke_sdpa_035` (bf16, D 128, Hq 64 / Hkv 8, causal, S 512). It has not reproduced since: not in the forced geometries on that cell, 15 single-graph runs, 4 full-corpus runs on MI350X, or the MI355X full-corpus re-check above. - Unified2D kernel time moved with develop's tiled-2D changes (#14542, #14616), not with this PR: at this head it is 1.035× (geomean) its pre-rebase time on the same 188 graphs, up to 1.36× on long causal bf16 D128 prefill (S 8192), while Dense in the same runs is flat (0.997). Dense still ranks first wherever it serves; the table above is at this head. - Measurement note: the pinned dnn-benchmarking (v1 result schema, which `sweep.py` reads) has no bottom-right causal PyTorch reference. The sweeps used upstream dnn-benchmarking #84's reference, with the 9 bottom-right corpus graphs written as `causal_mask: false, right_bound: 0, diagonal_alignment: BOTTOM_RIGHT` (the same band under hipDNN's `resolveDiagonalBand`): #84 reads `causal_mask: true` as top-left regardless of the alignment. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Christopher Millette <63608002+cgmillette@users.noreply.github.com>
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
The gfx950 census shard is written into the build tree as a real
CTestTestfile.cmake, but its Gfx950AttentionDense tiered suites pass--test-config hipkernel_Gfx950AttentionDense.tomlby bare name, and onlycmake --installput that TOML beside the shard. Running the shard from the build tree therefore failed withError: Config path does not existbefore any case ran. This PR copies the files a shard reads by bare name beside its build-tree copy, adds a configure-time guard so a later shard-routed engine cannot reintroduce the gap, and keeps CTest'sTesting/scratch directory out of the install.JIRA ID : ALMIOPEN-2879
Risk Assessment
Risk 3 of 5. Build and test infrastructure only, with a new check that can fail configuration; no library, kernel, dispatch or default-selection code changes. The guard fails only for a shard that names a
*.tomlnothing places beside it, which is the defect itself. The one shipped-content rule touched, the wholesaleinstall(DIRECTORY)of the test descriptor tree, gains one exclusion.ASIC Coverage
Architecture-independent build wiring; no multi-arch sweep is needed. The superbuild lanes (gfx942, gfx1151) do not generate the gfx950 shard, so in PR CI only the new host-only test exercises the copy and the guard. The gfx950 path was verified on gfx950 hardware.
Testing Summary
test_hkp_census_shard_files.pyruns the realhkp_reserve_census_shard()andhkp_finalize_census_install()undercmake -P. It fails if the copy, the guard or the argument check is removed.census/<arch>/Testingdirectory is installed after a build-tree shard run; otherwise the installed file set is unchanged.HIP_MLOPS_ENGINE.tomlandASM_SDPA_ENGINE.tomlsuites also use bare names, but in staging files (install_CTestTestfile.cmake,CTestTestfile.cmake.install) that CTest does not read in the build tree.Testing Checklist
ctest --test-dir <build>/lib/hipdnn_plugins/engines/test_arch_content/hip-kernel-provider/census/gfx950 --no-tests=error -V -R 'gfx950-attention-dense-external-integration_quick_suite$'- Devices: gfx950 - Status: Passed (8 ran, 0 skipped; fails on develop withConfig path does not exist)_standard_suite$- Devices: gfx950 - Status: Passed (45 ran, 0 skipped)ctest --test-dir <build>/lib/.../census/gfx950 -N -L hipkernel_Gfx950AttentionDense- Status: Passed (5 tiered suites andgpu_ref; TOML byte-identical toconfig/hipkernel_Gfx950AttentionDense.toml)pytest descriptor-packaging/tests -m quick -k test_hkp_census_shard_files- Status: Passed (4/4)FILESfromhkp_reserve_census_shard()and reconfigure - Status: Passed (configure fails with the guard message)cmake --build- Status: Passedcmake --installbefore and after the fix - Status: Passed (identical, except that develop also installscensus/gfx950/Testing/after an in-place shard run)hip-kernel-provider,GPU_TARGETS=gfx950, kernel ingestor, SDPA and rocKE ON - Status: PassedTechnical Changes
HkpPackaging.cmake:hkp_reserve_census_shard(<out> <arch> [FILES <file>...])records files that entries routed to the shard read by bare name, and rejects unrecognised arguments. A new_hkp_write_census_shard()copies them beside the shard's build-treeCTestTestfile.cmakeand fails configuration if the shard names a*.tomlthat is not there.src/CMakeLists.txt: the gfx950 dense TOML path is one variable, passed both asTEST_CONFIGand to the shard reservation. The existinginstall(FILES)for the TOML stays, since the integration-tests helper that issues it is shared with other providers.CMakeLists.txt: the wholesale install of the test descriptor tree excludescensus/<arch>/Testing.test_hkp_census_shard_files.py: new quick test of the copy, the guard and the argument check.hipdnn-superbuild-testskill: the build-tree registration of the engine-pinned suites, the build-tree path of the census shard, the engine-label rule, and which CTest files are install-only.Related: #12958.