Repository navigation
test(integration-tests): SDPA forward plan bundles (draft set) - #14999
Draft
CMiservaAMD wants to merge 12 commits into
Draft
CMiservaAMD wants to merge 12 commits into
CMiservaAMD wants to merge 12 commits into
Conversation
✅ 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. |
…set, not yet validated) Add 293 grouped sweeps (9,883 SDPA forward cases) under <tier>/SdpaFwdPlan/<Topology>/ in all four tiers: quick 23 sweeps / 749 cases, standard 102 / 4,152, comprehensive 59 / 1,582, full 109 / 3,400. The cases come from a model-driven pairwise plan for SDPA forward. Each sweep was produced by the opt-in grouped mode of import_graph.py (the import_mode marker records the grouping and the compact layout). Every support.json is empty; no engine claims are recorded yet. These bundles are not yet validated on a GPU. ASM_SDPA_ENGINE is known to fail some of them (ALMIOPEN-3041, ALMIOPEN-3042, ALMIOPEN-3043, ALMIOPEN-3123 to ALMIOPEN-3128) until engine skips are added. SUPPORT_MATRIX_OVERVIEW.md is not regenerated in this commit. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Add an opt-in grouped mode to import_graph.py so generated cases with the same graph structure land in one sweep. Default mode is unchanged. - --op-dir, --topology, --only-op-dir, --require-marked-sweep and --case-id choose where a case goes and refuse unintended appends. - --template-attributes, --template-alignment and --skeleton-value-form turn node attributes, tensor alignment and value forms into per-case values. Grouped sweeps also template min_required_engine_api_version and is_override_shape_enabled, which hipDNN derives from those features. - Grouped sweeps record an import_mode marker. Default mode never appends into a marked sweep; it still reports an exact duplicate found there. - Grouped sweeps use a compact sweep.json layout (--sweep-format), stored in the marker. One writer, dump_sweep, keeps the layout on append and in retier_bundles.py, which also refuses to merge sweeps whose grouping or layout differ. test_migration.py covers each flag, round trips through both Python expanders, determinism and the unchanged default mode (49 tests). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…derived Add docs/test-plans/sdpa-fwd.md, which explains how the SDPA forward plan cases were derived: the theoretical axes, the narrowing to meaningful and expressible combinations, the model and its rules, how pairwise cases and tiers are generated, and how rows become grouped bundles. Add the model files the cases are generated from (test-plans/sdpa-fwd/model.yaml, rules.yaml, injections.yaml) and a row in docs/README.md. The case list (test-plans/sdpa-fwd/cases.csv) that the document describes follows in a later commit. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…y per-tensor placeholders Register migration-scripts/test_migration.py as the CTest hipdnn_migration_scripts_python_tests, labelled quick and ffm-quick, so the import_graph.py tests run in CI. verify_golden_bundles.py now collects placeholders from every tensor field, not only dims, strides and data_type, and matches values entries by the first token of the placeholder path. Grouped sweeps resolve per-tensor alignment from values.alignment, which the verifier used to report as an unused entry once per case. Output for the existing tree is unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ues in the support matrix
The SDPA forward plan bundles use values the SdpaAttributes feature catalogue
did not list, so feature_matrix.py refused to render. Add them to
scripts/support_features/sdpa.py:
- dtype fp32 ("float"); an fp16 q with an fp32 o counts as mixed output
- head dims 1 and 256, and unequal q/k x v pairs (for example 128x64,
64x256, 256x128)
- layout "sbhd", q stride order (2, 0, 1, 3)
A graph outside the catalogue still raises. Update test_support_features.py
and regenerate SUPPORT_MATRIX_OVERVIEW.md. The SDPA percentages drop because
the catalogue grew; sdpa_backward.py shares these head-dim and layout lists.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…e lanes Add "*_SdpaFwdPlan_*" to the shared exclude list of the HIP_MLOPS_ENGINE, MIOpen and hipBLASLt integration lanes. These engines decline every SDPA graph; the exclude stops them from loading and querying 9,883 cases per run. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…attributes visitNativeType() in FlatbufferTensorAttributesUtils.hpp had no case for FP8_E4M3_FNUZ or FP8_E5M2_FNUZ, so createTensor and createRaggedTensor threw "Unsupported data type for tensor" for those types. Every fp8 FNUZ bundle stopped before its output comparison. Map both types to their native fp8 types, as SdkFrontendTypeConversions.hpp does, and test that every DataType except UNSET builds a tensor. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… in the plan bundles Each skip reason starts with the Jira ticket of the defect. Filters are exact test names, or a glob only where every matching case the engine accepted failed with that defect. archs list only the GPUs where the failure was seen. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…e within the GPU reference's launch limit Targeted case X10 moves from 2 batches x 128 heads to 2 batches x 96 heads. Its fp8 K/V byte offsets still pass 2^31, and the GPU reference can now check it (ALMIOPEN-3134). One case in full/SdpaFwdPlan/DqBv_DkBv_DvBv changes; no other case changes. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ses were made without the tooling The tools that generated the cases stay out of this change; the document describes them without paths or commands and notes they may be published later. Section 10 describes how the engine skips are made, section 11 states that no support claims are recorded yet, and the rule counts match rules.yaml. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… the model's The 33,314 valid pairs belong to the Layer 2 research space; the model keeps fewer settings and has 30,559. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…rd seed case S9 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
CMiservaAMD
force-pushed
the
users/cmiserva/almiopen-2893/sdpa-fwd-integration-tests-plan-bundles
branch
from
November 4, 2026 01:29
5e2c4bf to
a06c872
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #14999 +/- ##
===========================================
+ Coverage 77.42% 77.46% +0.04%
===========================================
Files 2903 2903
Lines 490638 490411 -227
Branches 74225 74227 +2
===========================================
+ Hits 379863 379867 +4
+ Misses 85771 85541 -230
+ Partials 25004 25003 -1
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
15 of 16 tasks
brockhargreaves-amd
added a commit
that referenced
this pull request
Nov 5, 2026
…rated bundles (#14622) ## Summary This PR adds the SDPA backward test plan required by RFC 0022 (Per-Op Test Plans, merged in #13937). It also adds a generator that writes the SdpaBackward bundles from the plan, and the bundles it writes. The 16 inherited SdpaBackward cases (all B = H = 1, packed BHSD, fixed scale) are replaced by 503 generated cases in four levels. Each case carries a purpose and traces to the plan items it covers. > Written with AI assistance (Claude Code). The plan follows the merged RFC 0022; it records one remaining deviation (DV-4) for reviewers to approve. JIRA ID : ALMIOPEN-2894 ## Risk Assessment Risk 2 of 5. The PR changes test data, test tooling, test configuration and documentation only; no library or provider code changes. The changes that reach beyond the SdpaBackward bundles: - `support_matrix_units.collect_sweep_units()` leaves out cases whose `metadata.purpose` is `probe`, as RFC 0022 §9 requires. Today only SdpaBackward cases carry a `purpose`, so other ops' matrices are unchanged. The committed `SUPPORT_MATRIX_OVERVIEW.md` is re-rendered. - `ASM_SDPA_ENGINE.toml`: two dtype-only SdpaBackward tolerance rows are replaced by 12 generated rows (seven tolerance classes), plus temporary skips for the open engine defects, each naming its ticket. - `ASM_SDPA_ENGINE_test_categories_integration.yaml`: six rejection probes that ASM_SDPA_ENGINE wrongly accepts are quarantined with their tickets (RFC 0022 §6). Both files affect the ASM lane only. - #14358's eight PyTorch golden cases move, unchanged, from `quick/SdpaBackward/{Default,AdditiveMask}/` to `quick/SdpaBackward/Golden/{Default,AdditiveMask}/`, so the generator owns the `Default` sweeps. They keep their fp16 tolerance row. `bhsd/` and the `*_fwd_pair` bundles are untouched. The main risk is CI noise from new cases. Every level has been run on gfx942 and gfx950 with no failure (see Testing Checklist). ## ASIC Coverage The bundles run in every provider lane whose tier YAML matches `*Sdpa*`. Only `ASM_SDPA_ENGINE` accepts SdpaBackward graphs today, and only on gfx942: its backward plan builder declines every other device ("Device string does not match gfx942"). Claims are harvested for gfx942 only. PR CI tests gfx942. On gfx950 every generated backward case is declined, so a gfx950 run only shows that nothing fails; one is recorded below. PR CI on gfx942 is sufficient. ## Testing Summary - **Generator tests** (`hipdnn_sdpa_bwd_bundle_generator_python_tests`, label `quick`, 19 tests). They fail when the checked-in sweeps or the TOML block drift from the generator. They also check legality, all-pairs coverage up to nightly, the per-case ceilings and cumulative budgets of each level, two-way traceability between plan items and cases, that every TOML skip still matches a written case, and that the tier-YAML quarantine holds only rejection probes. - **`generate_sdpa_bwd_bundles.py --dry-run --report`.** Reports counts, budgets, the all-pairs proof, legality, and each declared-but-gated case with its gate. - **Claims run** (2026-10-26, MI300X, gfx942). All generated bundles ran on `ASM_SDPA_ENGINE`, configured like CI. `harvest_support_claims.py` turned the passing cases into the 272 claims. Every claimed case passes its tolerance class with a margin of at least 2.4 (median 4.2). - **Measured run, every level on GPU** (2026-11-04, plan §4.5). `hipdnn_integration_tests` ran on `ASM_SDPA_ENGINE` configured like the CI lane: `--test-engine` (claims enforced), the PR's `ASM_SDPA_ENGINE.toml`, the tier YAML's quarantine as a GTest filter, device input fill and the GPU reference as oracle. It ran all four levels, on a fresh MI300X (gfx942) and an MI355X (gfx950) node, built from this branch. The bundles, TOML and YAML are identical to `6e014ad50e0`. - **PR CI** runs the quick and standard levels on gfx942. ## Testing Checklist - [x] Generator tests - CTest `hipdnn_sdpa_bwd_bundle_generator_python_tests` - 19 of 19 - Status: Passed - [x] Plan report - `python3 reference-data-scripts/generate_sdpa_bwd_bundles.py --dry-run --report` - 503 written, 77 gated, 0 legality violations, all-pairs 669 of 669 - Status: Passed - [x] Bundle and support-matrix checks - `verify_support_claims.py`, `render_support_matrix.py --output -`, `feature_matrix.py --check`; the bundle, claim, render and support-feature pytest suites - Status: Passed (`verify_golden_bundles.py` reports no SdpaBackward error; its 322 errors in other ops' bundles are on develop too) - [x] pre-commit on the changed files - Status: Passed - [x] Claims run - `hipdnn_integration_tests`, `ASM_SDPA_ENGINE`, MI300X (gfx942), 2026-10-26 - Status: Passed - [x] Measured run, gfx942 - `hipdnn_integration_tests --test-engine ASM_SDPA_ENGINE --test-config ASM_SDPA_ENGINE.toml '--gtest_filter=*SdpaBackward*-<the tier YAML quarantine>'` - MI300X, Slurm job 68275259 - Status: Passed. 553 tests in 43.3 s, 0 failed, no claim broken: - 289 executed and passed against the GPU reference, and 15 expected rejections rejected; - 226 declined by ASM, and 23 hit the DEF skips; - per level: quick 2.4 s, PR 2.8 s, nightly 16.5 s, full 21.4 s. - [x] Quarantine check, gfx942 - the six quarantined cases run alone - Status: all six fail as expected ("Engine ASM_SDPA_ENGINE accepted expected-rejection bundle"), which shows the quarantine is needed. - [x] Every level, gfx950 - same command - MI355X, Slurm job 68275260 - Status: Passed. 0 failed: ASM declines every backward graph on gfx950, and the rejection probes, quarantined ones included, are correctly rejected. - [x] The other SDPA bundles, same lane - `--gtest_filter=*Sdpa*:-*SdpaBackward*` - Status: Passed. gfx942: 47 passed, 84 skipped. gfx950: 16 passed, 115 skipped. 0 failed on either. - [ ] PR CI on gfx942 - TheRock multi-arch CI - Status: Pending on `6e014ad50e0`. The 2026-11-02 run failed one case, `prbr01`, which is now gated. Later runs failed only on the hipCCL configure break from #14976, which develop reverted in #15089 (merged into this branch). ## Technical Changes ### What this adds - **The plan:** `dnn-providers/integration-tests/docs/test-plans/sdpa-bwd.md`. It follows the RFC 0022 template: the operation, the feature inventory, the shape catalog, building the sweep (levels, budgets, probes, regressions, trade-off register), the golden reference, pass criteria, and exhaustive mode. - **The generator:** `reference-data-scripts/sdpa_bwd_plan.py` holds the declarations: axes, legality, the shape catalog, level rules and the cost model, one record per case. `generate_sdpa_bwd_bundles.py` writes the sweeps, the support-claim files and the engine-TOML tolerance block from them. Nobody edits the bundles by hand. RFC 0022 open question 1 cites this module as the per-op Python option. - **The bundles:** 503 written cases, spread across the levels: | Level | Directory | Cases | |---|---|---| | quick | `quick/` | 56 | | PR | `standard/` | 152 | | nightly | `comprehensive/` | 220 | | full | `full/` | 75 | - Each case carries `purpose` (`claim`, `probe` or `regression`) and its plan items. - 77 more cases are declared but not written yet, each waiting on a named gate: an oracle extension, a harness feature or golden data. - One of the 77 is PRB-R01 (`dropout_probability = 0`). Since #14358 the references decline any graph that sets `dropout_probability`, so it waits on ALMIOPEN-3149 or ALMIOPEN-3150. - **Levels:** the cost model assigns each case to the lowest level whose per-case limits it fits: | Level | Per-case time | Per-case bytes | Cumulative budget | |---|---|---|---| | quick | 1.5 s | 4 MiB | 20 s | | PR | 5 s | 16 MiB | 120 s | | nightly | 75 s | 64 MiB | 480 s | | full | 300 s | no limit | 1,800 s | Each cumulative budget covers that level and every level below it, and keeps 5 % in reserve. The budgets come from a CPU-reference cost model; with the GPU reference, the measured run above is far below them. The ASM lane YAML picks the cases up by GTest prefix (`quick_*Sdpa*`, `standard_*Sdpa*` and so on), and each level also runs every level below it. - **Rejection probes** (21) carry `metadata.expected_rejection` (#14346): the harness never executes them, and an engine that accepts one fails the lane. PRB-M02 and PRB-M02f (`right_bound = -2`) are not written, because the frontend rejects them since #14358; `TestSdpaBwdNode.PreValidateChecksDiagonalBandBounds` covers that rule. - **dBias:** the five dBias probes and the live GL-14 case are written, now that #14358's references model dBias. To keep PR and nightly inside their reserve, the plan's trim rule moved M07.pr to nightly and M26.n to full. - **Tolerance classes:** seven classes, written as 12 `[[tolerance_overrides]]` rows between generated BEGIN and END markers in `ASM_SDPA_ENGINE.toml`. - **Support claims:** 272 claims for `ASM_SDPA_ENGINE` on gfx942. - **Probes left out of the support matrices** (RFC 0022 §9). This is the `support_matrix_units` change above. - **Generator file I/O is UTF-8**, so its tests pass on Windows. ### Temporary ASM_SDPA_ENGINE exclusions These stay until each fix lands; the PR with the fix deletes its entries. **Open engine defects** (`[[test_skips]]` in `ASM_SDPA_ENGINE.toml`). These are valid graphs that ASM accepts and computes wrong. Other engines still run them. | Cases | What goes wrong | Ticket | Fix | |---|---|---|---| | DEF-01 | head-dim stride other than 1 | ALMIOPEN-3036 | #14684 | | DEF-02 | non-packed Stats | ALMIOPEN-3037 | #14684 | | DEF-04 | Hk = Hq > Hv | ALMIOPEN-3039 | #14684 | | DEF-05 and H61 | odd element strides | ALMIOPEN-3038 | #14684 | | DEF-06 | scale of 0 or below | ALMIOPEN-3043 | #14688 | | DEF-03 | fp16 dV precision in the AITER kernel | ALMIOPEN-3046 | none in hipDNN (kernel) | **Quarantined rejection probes** (`_common_excludes` in `ASM_SDPA_ENGINE_test_categories_integration.yaml`, as RFC 0022 §6 asks). These are undefined graphs ASM should decline but accepts. | Cases | What ASM accepts | Ticket | Fix | |---|---|---|---| | prbu01, prbu03 | O or dQ dims that differ from Q's | ALMIOPEN-3039 | #14684 | | prbu13 | a dQ sequence stride below D (overlapping rows) | ALMIOPEN-3038 | #14684 | | prbt04, prbt04n | an O dtype that differs from Q's | ALMIOPEN-3050 | none yet | | prbr05 | `dropout_scale_inv` without a dropout mask | ALMIOPEN-3051 | none yet | ### Dependencies (all merged) - #13937, RFC 0022 (2026-10-30). Its first-phase documents, the methodology guide and the `docs/test-plans/` index, are not on develop yet, so the plan links the RFC's sections directly. - #14312: O and Stats derived from the forward of the same inputs. - #14483: `left_bound` uses cuDNN's meaning. The bundles write `left_bound = W`. - #14358 (2026-11-02): SDPA references cover the full graph API, and the GPU reference is the first oracle in `auto` mode. It also adds frontend rule R14 (`right_bound >= -1`). - #14346 (2026-11-03): the `expected_rejection` and `no_*_reference` bundle flags. - #14709 (2026-11-02): zero strides in the data SDK `Tensor`. ### Interaction with other open PRs - **#14684, #14686, #14687, #14688** (ASM declines, ALMIOPEN-3035). Each deletes its own entries: TOML skips for the DEF cases, and tier-YAML lines for the quarantined probes it fixes (#14684: prbu01, prbu03, prbu13). Whichever of these and this PR lands second removes the entries. - **#14999** (SDPA forward plan bundles, ALMIOPEN-2893). Both PRs edit two files the same way: - `ASM_SDPA_ENGINE.toml`: keep this PR's generated block and put #14999's `SdpaFwdPlan` skips after the `# END SdpaBackward tolerance classes` line. The generator test only checks the BEGIN to END span. The filters do not overlap (`SdpaBackward` vs `SdpaFwdPlan`). - `SUPPORT_MATRIX_OVERVIEW.md`: re-render it; do not hand-merge. The two PRs also add lines near each other in `integration-tests/CMakeLists.txt` and `test_categories_external.yaml`; these auto-merge. ### Before review - [x] `prbr01` (`dropout_probability = 0`) gated; ALMIOPEN-3149 and ALMIOPEN-3150 track the decision. - [x] Rejection probes carry `expected_rejection`; the six ASM accepts are quarantined in the ASM tier YAML with their tickets. - [x] `prbm02` and `prbm02f` removed; the frontend unit test of R14 covers them. - [x] The five dBias probes and GL-14 written; REG-08 (rank-3 Stats) still waits for the derivation extension. - [x] The plan refreshed against the merged RFC: DV-1 and DV-2 now agree with it, DV-3 is an RFC §11 gap, and only DV-4 remains a deviation. - [x] The measured run (plan §4.5) recorded above. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
This branch has not been deployed
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
This PR adds 9,883 SDPA forward test cases to the cross-provider integration suite, in 293 bundles
across all four tiers, together with the model files they were derived from and a document that
explains the derivation.
Some words first:
The suite builds each case, asks each engine whether it supports it, runs it, and compares the
output with a reference.
GPU) whose output the engine's output is checked against.
The cases come from a written model of every setting SDPA forward can vary, plus rules that remove
meaningless combinations. A case generator turns the model into a case list with one CSV row per
case, and further tools turn the rows into bundles. A document explains every step, so a reviewer
can check it.
Two things are deliberately not in this PR:
new
support.jsonis empty. Claims come in a later PR.the bundle tools are proof-of-concept programs that are still being worked on. They are left out
of this PR and may be shared once they settle. The bundles are committed as they are.
JIRA ID : ALMIOPEN-2893
Risk Assessment
Risk 2 of 5. Almost everything is test data, test tooling and documentation. No library code
changes. The changes that reach beyond the added bundles:
test_sdkgains FP8 E4M3 FNUZ / E5M2 FNUZ cases invisitNativeType(). Without them, a test thatbuilds an FNUZ tensor from flatbuffer attributes fails with "Unsupported data type for tensor". It
only adds cases to a switch.
*_SdpaFwdPlan_*. Those engines donot run SDPA forward, so they would only load and decline about 4,900 cases per PR run.
ASM_SDPA_ENGINE.tomlgains skip entries for known engine defects (see below). They only stopcases from running; they change no engine behaviour.
test_migration.pyjoins CTest with thequicklabel, so it runs in every lane that runs quickPython suites, including Windows (not run on Windows here).
import_graph.pygains an opt-in grouped mode. Its default output is byte-identical to before: alocal harness re-imported all 11,390 existing graphs with the old and the changed scripts and got
identical results, and a command-line check over 649 graphs gave identical trees and logs.
ASIC Coverage
These cases only run on engines that accept them. Today that is
ASM_SDPA_ENGINE, which runs ongfx942 and gfx950; both were run on Linux (see Validation evidence). Other lanes exclude the cases.
On other architectures every engine declines the cases, which costs about 3 ms per case (measured on
HIP_MLOPS). The reference check ran on gfx942 and gfx950. PR CI is sufficient for the rest: no other
architecture can run these cases.
Testing Summary
migration-scripts/):test_migration.py, 49 tests, all pass; it runs inCTest as
hipdnn_migration_scripts_python_tests(labelquick).command-line check over 649 graphs, show byte-identical output between the old and changed scripts.
verify_golden_bundles.py,verify_support_claims.py,render_support_matrix.pyandfeature_matrix.py --checkall return 0.test_support_features.py, 1,827 tests pass.hipdnn_test_sdk_testspass with the FNUZ change.hipdnn_integration_tests,ASM_SDPA_ENGINE, gfx942 and gfx950, Linux, beforethe skips): every case registers (749 / 4,152 / 1,582 / 3,400 per tier). Wall time per tier is
10–122 s, far under the lane limits. These runs found the failures that the skips cover.
passing case: gfx942 535 of 535 failures skipped, 0 of 708 passing cases skipped; gfx950 354 of
354 and 0 of 670.
their outputs agree? Phase 1: 43 of 43 accepted and agreeing. Phase 2 (all 9,883 cases, on gfx942 and gfx950): every case has a result; both references accept 9,823 and agree on all of them; both decline the same 60 full-tier cases (explained below).
Images:
hipdnn_20261031_gfx942.sqsh,hipdnn_20261031_gfx950.sqsh.Testing Checklist
python3 migration-scripts/test_migration.py(49 tests) and CTesthipdnn_migration_scripts_python_tests- Status: Passedverify_golden_bundles.py,verify_support_claims.py,render_support_matrix.py,feature_matrix.py --check- Status: Passedtest_support_features.py(1,827) - Status: Passedhipdnn_test_sdk_tests- Status: Passedhipdnn_integration_tests --test-engine ASM_SDPA_ENGINE, four tiers, before the skips - Status: Ran; failures traced to tickets and skippedpre-commit run --files ...- Status: PassedTechnical Changes
What is in this PR
Bundles:
integration-test-bundles/{quick,standard,comprehensive,full}/SdpaFwdPlan/<Topology>/(
graph.template.json,sweep.json, and an emptysupport.json). A topology is the set oftensors and options a graph uses (for example "with a causal mask and a Stats output"); each
topology in each tier is one sweep. Each case carries
case_key(the key of its row in the caselist) and a seed derived from it.
About 21 MB in 879 files. The sweeps use a compact layout (one case per line) to keep the size
and line count down.
One full-tier case, targeted case X10 (fp8, Sq 1, Skv 131072), checks fp8 K/V byte offsets past
2^31. With 2 batches of 128 heads its K tensor has 2^32 elements, and the GPU reference cannot
launch a kernel that large (ALMIOPEN-3134), so only the slow CPU reference could check it. X10 now
uses 2 batches of 96 heads: its byte offsets still pass 2^31, and both references accept it and
agree on gfx942 and gfx950. When ALMIOPEN-3134 is fixed, X10 goes back to 128 heads. The change
touches one case in
full/SdpaFwdPlan/DqBv_DkBv_DvBv/sweep.jsonandmodel.yaml.Model files
test-plans/sdpa-fwd/model.yaml,rules.yaml,injections.yaml: the settings,their values, the tier recipe, the rules that remove combinations, and the cases added on purpose.
Suite tooling: an opt-in grouped import mode in
migration-scripts/import_graph.py(helpersin
bundle_utils.py,retier_bundles.py), tests intest_migration.py, and a section inmigration-scripts/README.md. Grouped mode puts every case of one topology in one sweep andturns the per-case attributes into template placeholders.
verify_golden_bundles.pylearnsper-tensor placeholders.
CTest wiring for
test_migration.py(CMakeLists.txt,test_categories_external.yaml).Support-feature catalogue (
scripts/support_features/sdpa.py): it lists the head dimensions,fp32 and the (2,0,1,3) layout that the cases use, so the coverage overview can count them.
SUPPORT_MATRIX_OVERVIEW.mdis re-rendered.Lane excludes:
*_SdpaFwdPlan_*in every category of the HIP_MLOPS, MIOpen and hipBLASLtintegration tier YAMLs.
test_sdk: FP8 FNUZ cases in
visitNativeType(), with tests.Engine skips in
ASM_SDPA_ENGINE.toml(next section).Derivation document
docs/test-plans/sdpa-fwd.md, written for a reader who does not know thetopic, and a row for it in
docs/README.md. It covers the model, how the cases were generated,the tiers, how rows become bundles, and how the skips are made and removed.
What is not in this PR
engines have been run with the skips in place.
are still in development and are not part of this PR; they may be published later. The bundles
are committed as they are.
the document describes its columns, because each case's
generated_from_sha256is the SHA-256 ofthat list and its
case_keynames its row. The suite is expected to read such a list directly ina later step; the
SdpaFwdPlanbundles will then be removed.Known engine defects and their skips
A skip tells the suite not to run a case on one engine, with a reason. The case stays in its tier,
so every other engine still runs it, and it runs again on
ASM_SDPA_ENGINEwhen its skip isdeleted.
How the skips were made: every tier was run on
ASM_SDPA_ENGINEon gfx942 and gfx950, and everyfailure was run again with the other input fill (the suite can make input data on the GPU or on
the CPU). Every failure was traced to a Jira ticket for a known engine defect (no failure was
skipped without one), and each skip entry lists only the GPUs where the failure was seen. The
entries are in
dnn-providers/hip-kernel-provider/config/ASM_SDPA_ENGINE.toml, after the comment-- SdpaFwdPlan known engine defects --; section 10 of the document describes the process in full.Each
reasonstarts with the Jira key, then a colon, then the defect in one sentence, for example:When a fix lands, the PR with the fix deletes every entry whose reason starts with its key. Fixes for
ALMIOPEN-3041 and 3042 (#14686) and for ALMIOPEN-3043 (#14688) are in separate PRs.
A case can hit more than one defect. A skip covers the whole test and the first matching entry
wins, so such a case is skipped under one ticket only. 101 cases have more than one known
defect. When the ticket they are skipped under is fixed, they will still fail until their other
ticket is fixed too, so they must move to the entry for that ticket rather than be un-skipped.
The 101 cases with more than one known defect (GTest name: skipped under; other ticket)
ALMIOPEN-3042 has two parts, marked "(Q/K/V strides)" and "(Stats layout)".
comprehensive_SdpaFwdPlan_Base.bf16_e2a5c8a2105ee079: ALMIOPEN-3043; also ALMIOPEN-3042 (Q/K/V strides)comprehensive_SdpaFwdPlan_DSeedBv_DOffBv_Rng.bf16_4de3ea6cf356794b: ALMIOPEN-3126; also ALMIOPEN-3042 (Q/K/V strides)comprehensive_SdpaFwdPlan_DSeedBv_DOffBv_Rng.bf16_da52eced58d3635f: ALMIOPEN-3126; also ALMIOPEN-3042 (Q/K/V strides)comprehensive_SdpaFwdPlan_DSeedBv_DOffBv_Rng_Max.bf16_18ccbe1caeb3217c: ALMIOPEN-3126; also ALMIOPEN-3124comprehensive_SdpaFwdPlan_DSeedBv_DOffBv_Rng_SumExp.bf16_2c1ed2d20e8536ad: ALMIOPEN-3126; also ALMIOPEN-3124comprehensive_SdpaFwdPlan_DSeedBv_DOffBv_Rng_SumExp.bf16_cddab8accb3370aa: ALMIOPEN-3126; also ALMIOPEN-3124comprehensive_SdpaFwdPlan_ScRt_DSeedBv_DOffBv_Rng_Stats.bf16_e50c50faea05a145: ALMIOPEN-3126; also ALMIOPEN-3042 (Stats layout)comprehensive_SdpaFwdPlan_Stats.bf16_616818f26c39eb6f: ALMIOPEN-3042 (Stats layout); also ALMIOPEN-3043full_SdpaFwdPlan_Base.bf16_39bd7cbbfa5fe92b: ALMIOPEN-3043; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_DSeedBv_DOffBv_Rng.bf16_6ed083591959defe: ALMIOPEN-3126; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_DSeedBv_DOffBv_Rng.bf16_f040f887ddad5cb6: ALMIOPEN-3126; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_DSeedBv_DOffBv_Rng_Max.bf16_385f8dd6e2aa9f17: ALMIOPEN-3126; also ALMIOPEN-3124full_SdpaFwdPlan_DSeedBv_DOffBv_Rng_SumExp.bf16_8eab9fe6024794d5: ALMIOPEN-3126; also ALMIOPEN-3124full_SdpaFwdPlan_DSeedBv_DOffBv_Rng_SumExp.bf16_f35a79d8051ac65e: ALMIOPEN-3126; also ALMIOPEN-3124full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_085a41450e3b061b: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_41ae1fd5088002ec: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_5758e86bc28299f4: ALMIOPEN-3125; also ALMIOPEN-3043full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_60ec4f495de4a5c8: ALMIOPEN-3125; also ALMIOPEN-3043full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_7f4b8e1f97baf09a: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_9439bbc4e657a75a: ALMIOPEN-3125; also ALMIOPEN-3043full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_a11a311467a9d3b8: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_a8f26e16df647595: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_cd68d00fcae399b1: ALMIOPEN-3125; also ALMIOPEN-3043full_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_d8eec7df2382a56b: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_Max.bf16_6685cd705112febb: ALMIOPEN-3124; also ALMIOPEN-3043full_SdpaFwdPlan_Max.bf16_b9490d539f410f21: ALMIOPEN-3124; also ALMIOPEN-3043full_SdpaFwdPlan_Max_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_02a5051cdab5d461: ALMIOPEN-3125; also ALMIOPEN-3124full_SdpaFwdPlan_Max_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_8c1f2f7de591850f: ALMIOPEN-3125; also ALMIOPEN-3124full_SdpaFwdPlan_Max_SumExp.bf16_0b8710633ad92461: ALMIOPEN-3124; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_Max_SumExp.bf16_84da782fbe2c1acb: ALMIOPEN-3124; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_Max_SumExp_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_1f34f7cce3a63523: ALMIOPEN-3125; also ALMIOPEN-3124full_SdpaFwdPlan_Max_SumExp_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_919d426d3329825f: ALMIOPEN-3125; also ALMIOPEN-3124full_SdpaFwdPlan_ScBv_SumExp.bf16_6af13bf3c50e59ce: ALMIOPEN-3124; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_ScBv_SumExp.bf16_ebba327af827c7af: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3124full_SdpaFwdPlan_ScRt_DSeedBv_DOffBv_Rng_Stats.bf16_8c17cae75b6a6202: ALMIOPEN-3126; also ALMIOPEN-3042 (Stats layout)full_SdpaFwdPlan_Stats.bf16_055b78f5119b9800: ALMIOPEN-3042 (Stats layout); also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_Stats.bf16_2a71d620e8c9942a: ALMIOPEN-3042 (Stats layout); also ALMIOPEN-3043full_SdpaFwdPlan_Stats.bf16_2c43229d8220fbd2: ALMIOPEN-3042 (Stats layout); also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_Stats.bf16_2e62a03e6f84a7c8: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)full_SdpaFwdPlan_Stats.bf16_345a0d1d76346b29: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)full_SdpaFwdPlan_Stats.bf16_8d53fe10f0c97305: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)full_SdpaFwdPlan_Stats.bf16_b44f0167af7efce5: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)full_SdpaFwdPlan_Stats_SumExp.bf16_17e9b85d53154b92: ALMIOPEN-3042 (Stats layout); also ALMIOPEN-3124full_SdpaFwdPlan_Stats_SumExp.bf16_da3beac017df4621: ALMIOPEN-3042 (Stats layout); also ALMIOPEN-3124full_SdpaFwdPlan_SumExp.bf16_377b02a9d6c01de5: ALMIOPEN-3124; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_SumExp.bf16_653fd1b493db0c41: ALMIOPEN-3124; also ALMIOPEN-3043full_SdpaFwdPlan_SumExp.bf16_6e803f11b0e67494: ALMIOPEN-3124; also ALMIOPEN-3043full_SdpaFwdPlan_SumExp.bf16_7cb531c14227b254: ALMIOPEN-3124; also ALMIOPEN-3042 (Q/K/V strides)full_SdpaFwdPlan_SumExp_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_3e0c54d549e1c4c8: ALMIOPEN-3125; also ALMIOPEN-3124full_SdpaFwdPlan_SumExp_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_64c2d041420d0359: ALMIOPEN-3125; also ALMIOPEN-3124standard_SdpaFwdPlan_Base.bf16_47d623f715e3e1df: ALMIOPEN-3043; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_Base.bf16_a979e3bb9d1ce8ea: ALMIOPEN-3043; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_Base.bf16_cd939cca94979445: ALMIOPEN-3043; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng.bf16_367d7816ad4f2217: ALMIOPEN-3126; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng.bf16_38fca711ba17b619: ALMIOPEN-3126; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng.bf16_ed0d2bb571ff2d47: ALMIOPEN-3126; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng.bf16_f98e8782161d15d5: ALMIOPEN-3126; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng_Max.bf16_70953c6ac1abf277: ALMIOPEN-3126; also ALMIOPEN-3124standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng_Max.bf16_e1e057e185e89a58: ALMIOPEN-3126; also ALMIOPEN-3124standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng_SumExp.bf16_293b42f60aeb3612: ALMIOPEN-3126; also ALMIOPEN-3124standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng_SumExp.bf16_57f52fa78b15df1d: ALMIOPEN-3126; also ALMIOPEN-3124standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng_SumExp.bf16_81fccfd39b8bae4a: ALMIOPEN-3126; also ALMIOPEN-3124standard_SdpaFwdPlan_DSeedBv_DOffBv_Rng_SumExp.bf16_bf260f91099c1df2: ALMIOPEN-3126; also ALMIOPEN-3124standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_21929cfdb60c920d: ALMIOPEN-3125; also ALMIOPEN-3043standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_23b56b69ee42c7a0: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_28cf9087438f4a57: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_2f6ef475bb9955a1: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_434f020145dc78d2: ALMIOPEN-3125; also ALMIOPEN-3043standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_60112998c8449209: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_75c01d935951ca58: ALMIOPEN-3125; also ALMIOPEN-3043standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_8cc1ff8c3795fe62: ALMIOPEN-3125; also ALMIOPEN-3043standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_9c9838241d77ba6a: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_de645541a2648f50: ALMIOPEN-3125; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_Max.bf16_f0b7bfae99e5bf76: ALMIOPEN-3124; also ALMIOPEN-3043standard_SdpaFwdPlan_Max.bf16_fe65db707b806474: ALMIOPEN-3124; also ALMIOPEN-3043standard_SdpaFwdPlan_Max_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_3dcf39c2a94d8dc1: ALMIOPEN-3125; also ALMIOPEN-3124standard_SdpaFwdPlan_Max_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_78ebd5f35e223285: ALMIOPEN-3125; also ALMIOPEN-3124standard_SdpaFwdPlan_Max_SumExp.bf16_a05037da91bcc3f5: ALMIOPEN-3124; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_Max_SumExp.bf16_a9ec3eba57ed7994: ALMIOPEN-3124; also ALMIOPEN-3042 (Q/K/V strides)standard_SdpaFwdPlan_Max_SumExp_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_a3b09d2163dd69c7: ALMIOPEN-3125; also ALMIOPEN-3124standard_SdpaFwdPlan_Max_SumExp_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_c715c785ef646d06: ALMIOPEN-3125; also ALMIOPEN-3124standard_SdpaFwdPlan_ScBv_SumExp.bf16_07c21ab7eba17861: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3124standard_SdpaFwdPlan_ScBv_SumExp.bf16_770e215de83713e2: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3124standard_SdpaFwdPlan_ScRt_DSeedBv_DOffBv_Rng_Stats.bf16_2653f87bac2f1568: ALMIOPEN-3126; also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_ScRt_DSeedBv_DOffBv_Rng_Stats.bf16_7c12b5aa3e954ee0: ALMIOPEN-3126; also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats.bf16_0bcd85e3eb63a103: ALMIOPEN-3043; also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats.bf16_2544ff6d8943aa9b: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats.bf16_2557c7a0fc74e7f9: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats.bf16_5d9aaab3a2afa285: ALMIOPEN-3043; also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats.bf16_736f1adc2e7d0026: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats.bf16_7a7ee063a4b52b80: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats.bf16_8bec946bd80e3cc4: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats.bf16_98e96a206233c714: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3042 (Stats layout)standard_SdpaFwdPlan_Stats_SumExp.bf16_1dac8b5e1b60fd9e: ALMIOPEN-3042 (Stats layout); also ALMIOPEN-3124standard_SdpaFwdPlan_Stats_SumExp.bf16_c868341d9fbfbcfe: ALMIOPEN-3042 (Stats layout); also ALMIOPEN-3124standard_SdpaFwdPlan_SumExp.bf16_39039c0e809c7cb8: ALMIOPEN-3043; also ALMIOPEN-3124standard_SdpaFwdPlan_SumExp.bf16_c3939c2fad57cc9f: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3124standard_SdpaFwdPlan_SumExp.bf16_d463f9eca753eb02: ALMIOPEN-3124; also ALMIOPEN-3043standard_SdpaFwdPlan_SumExp.bf16_e7eb3151c08cdda6: ALMIOPEN-3042 (Q/K/V strides); also ALMIOPEN-3124standard_SdpaFwdPlan_SumExp_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_2950328307ae6e7a: ALMIOPEN-3125; also ALMIOPEN-3124standard_SdpaFwdPlan_SumExp_DqBv_DkBv_DvBv.fp8_e4m3_fnuz_7a4f6778868574a3: ALMIOPEN-3125; also ALMIOPEN-3124Cases that are not bundles
6,846 of the 16,729 rows of the case list are not bundles. A bundle can describe a graph and its sizes, but not
the exact contents of an input tensor or tensors that share memory. Reasons (a row can have more
than one):
seq_len_contentspage_table_contentsblock_mask_bitsdropout_mask_cycleragged_offsetsdescale_liststensor_overlaparch_dtypezero_stride_needs_pr14709dropout_without_rng_dumpValidation evidence
Engine runs (
ASM_SDPA_ENGINE, before the skips):Passed / failed / declined: gfx942 quick 134/44/571, standard 261/247/3,644, comprehensive 134/49/1,399,
full 179/195/3,026; gfx950 quick 124/12/613, standard 251/186/3,715, comprehensive 125/39/1,418,
full 170/117/3,113. The FP8 FNUZ cases were re-run on gfx942 after the test_sdk fix (job 68252469).
Reference check. It runs every bundled case through both hipDNN references, the GPU one and the
CPU one, without an engine, using a local tool that is not part of this PR. A case passes when both
references accept it and their outputs agree: fp8 outputs within one step of the 8-bit format
(1 ULP) with any NaN or infinity in the same places, other outputs within the suite's tolerance.
Phase 1 (calibration, gfx942, job 68253911):
43 cases chosen to be the largest and the hardest (every fp8 output format, fp32 outputs, the
largest case of each tier). 43 of 43 accepted by both references, 43 of 43 agree.
Phase 2 (every bundled case, gfx942 jobs 68254208, 68254209, 68254210; gfx950 jobs 68254214,
68254215, 68254218): on both GPUs, every case has a result.
fp8 outputs are compared within one fp8 rounding step: at most about 1 element in 100,000 differs by
one step, and no case is above 1 %. Other outputs use the suite's tolerances.
Both references decline the same 60 full-tier cases on both GPUs. Each changes one setting from a base
case. Every one is a valid combination by the plan's definition rules and the frontend accepts the
graph, so these are gaps in the references, not invalid cases:
with fp16 out, bf16 in with fp8 out).
inputs with one); the bf16 part is ALMIOPEN-3132.
They stay in the full tier. In the engine runs,
ASM_SDPA_ENGINEaccepted only the 6 bf16-descalecases (skipped, see above); every other engine declined all 60. Until the references support them, an
engine that accepts one of these cases reports it as Unverifiable.
Notes for reviewers
(ALMIOPEN-3132), and
ASM_SDPA_ENGINEaccepts them and ignores the descale (ALMIOPEN-3133). Theystay in the bundles and are skipped on
ASM_SDPA_ENGINEuntil the reference fix lands.reference_source, the case list that the bundles came from(the case list with five columns appended: topology, bundle status, exclusion reasons, bundle
path, GTest name), and, as its
generator, the program that exported the graph. Neither isincluded in this PR. Section 9.4 of the document explains these fields.
import path fails "round-trip verify failed after extraction". The old and the changed scripts give
the same result, so this PR does not cause it.
ASM_SDPA_ENGINE.toml,SUPPORT_MATRIX_OVERVIEW.md(re-render, do not hand-merge), and add lines near each other inintegration-tests/CMakeLists.txtandtest_categories_external.yaml. Both add a file underdocs/test-plans/(different files).🤖 Generated with Claude Code