Repository navigation
feat(hipdnn): allow zero (broadcast) strides in data SDK Tensor - #14709
Merged
CMiservaAMD merged 4 commits intoNov 2, 2026
Merged
Conversation
A zero stride on a dimension of size greater than 1 is a broadcast: every index along that dimension reads the same memory. Tensor now accepts non-negative strides, sizes its allocation from the largest reachable offset, and reports a tensor with a broadcast dimension as not packed. isTensorPacked applies the same rule. Negative strides remain unsupported. ALMIOPEN-3109 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ro-stride tests ALMIOPEN-3109 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
✅ 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 #14709 +/- ##
===========================================
+ Coverage 77.34% 77.36% +0.02%
===========================================
Files 2894 2893 -1
Lines 487274 485987 -1287
Branches 73752 73472 -280
===========================================
- Hits 376860 375943 -917
+ Misses 85525 85222 -303
+ Partials 24889 24822 -67
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Contributor
|
Review: approve, three optional suggestions. Tracking (ALMIOPEN-3109), regression proof and CI coverage all check out. Linux gfx94X and Windows gfx110X ran green for hipdnn, the integration tests, samples and the three providers. I found no correctness or ownership issues. I read Suggestions (non-blocking)
Follow-ups (out of scope here)
|
…d update isTensorPacked comment Remove IsTensorPackedZeroStrideBroadcast, whose cases were already rejected by the span check, and the identity loop in the zero-stride fill test. State the zero-stride rule in the isTensorPacked comment instead of the claim that negative strides are supported. ALMIOPEN-3109 Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Contributor
Author
|
@adickin-amd thanks for the review. Responses to each point (changes are in 86bb421):
Follow-ups:
|
CMiservaAMD
enabled auto-merge (squash)
November 2, 2026 11:13
CMiservaAMD
deleted the
users/cmiserva/almiopen-3109/data-sdk-zero-stride
branch
November 2, 2026 14:37
1 task
Contributor
TheRock Submodule Bump ActivityNewest first
|
10 of 13 tasks
15 of 16 tasks
brockhargreaves-amd
added a commit
that referenced
this pull request
Nov 4, 2026
…rged RFC 0022 and develop - The status note, the level and golden-criteria links point at the merged RFC 0022 (#13937); its methodology guide and test-plans index do not exist yet. The baseline is develop 961dd76, with #14358, #14346 and #14709. - Deviations: the merged RFC adopted DV-1 (§7 exception) and DV-2 (§12.3 uplift default) and lists DV-3 as an oracle gap (§11); only DV-4 still needs approval. - Gates that named #14358 now name what is left (I6, I4, HG4, HG20, GPU forward, timing run); the reference-status table, the oracle columns and the §6.4 gap table show #14358's references as in place. - §2.8 to §4: F_active is F_R13R14; PRB-M02 and PRB-M02f are frontend unit tests; rejection probes carry expected_rejection, and the six ASM accepts are quarantined in the tier YAML; the dBias probes and GL-14 are written; the second trim (M07.pr, M26.n); counts, budgets and topology rows from the generator. - §4.5 records the 2026-11-04 measured run on gfx942 and the gfx950 run. - §5: #14358's golden cases stay under quick/SdpaBackward/Golden; GT1 and the SEM-010 −inf exception are in place, so GL-09 and GL-16b can join the first golden data. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
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 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 data SDK
Tensorrejected any stride of 0. A zero stride on a dimension of size > 1 is a broadcast: every index along that dimension reads the same element (the layout PyTorchexpand()produces). Tests could therefore not allocate broadcast inputs, for example K, V or a bias expanded across batch or heads, so the CPU/GPU references and integration tests could not exercise those layouts. This PR letsTensoraccept non-negative strides, asRaggedTensoralready does. Whether an operation accepts a broadcast tensor stays the frontend's and the engines' decision.JIRA ID : ALMIOPEN-3109
Risk Assessment
Risk level 2. The change is host-only, header-only code in the data SDK (
Tensor.hpp,ShapeUtilities.hpp) with no new option or flag. It widens what theTensorconstructor accepts (stride 0 used to throw) and narrows what counts as packed: a tensor with a zero stride on a dimension of size > 1 is now never packed, in bothTensor::isPacked()andisTensorPacked(). Before, a layout such as dims {2,2} / strides {3,0} was wrongly reported packed.isTensorPacked()is used by the hip-kernel-provider and miopen-provider applicability checks and the GPU batchnorm reference, so such a layout is now correctly routed away from linear fast paths. Positive-stride layouts are unaffected.ASIC Coverage
Architecture-independent: no kernel, dispatch table or arch guard changes. Standard PR CI is sufficient (TheRock presubmit runs
hipdnn_data_sdk_testsand the dependent provider jobs on gfx94X Linux and gfx110X Windows). No multi-arch sweep is required. Locally verified on gfx942 (MI300X).Testing Summary
TestTensor.cpp(one element read along a zero-stride axis, all-zero strides allocate one element, zero stride is not packed even when element span equals item count, zero stride on a size-1 axis stays packed, random fill broadcasts along the zero-stride axis, host/device round trip) and 3 inTestShapeUtilities.cppforisTensorPacked().TestTensor.ConstructorThrowsOnZeroStrideis removed because zero strides are now valid by design;ConstructorThrowsOnNegativeStridestill covers the rejection path.isTensorPacked({2,2},{3,0})returning true); all pass at this head.isTensorPacked()) and the integration quick suites were run to catch consumer regressions.Testing Checklist
All local runs on gfx942 (MI300X), superbuild with components
hipdnn;hipdnn-integration-tests;hip-kernel-provider, Release and Debug,ctest --timeout 2000in each build tree.ctest -R hipdnn_data_sdk_tests- Status: Passed (1542 Release / 1543 Debug, all 9 new tests pass)ctest -R "hipdnn_(backend|frontend|test_sdk|plugin_sdk|public_backend|public_frontend)_tests"- Status: Passed (Release and Debug)ctest -R hip_kernel_provider_tests_quick_suite- Status: Passedctest -R "hipdnn_integration_tests_unit_tests_quick_suite|hip_kernel_provider_integration_tests_quick_suite|hip-kernel-provider-external-integration_quick_suite"- Status: Passedpre-commit run --from-ref 7e3d78a3d16 --to-ref 6b7c38da71b- Status: PassedisTensorPacked()) - TheRock PR CI - Status: PendingTechnical Changes
Tensorconstructor: strides must be >= 0 (dimensions still >= 1). The two positivity checks are merged into onevalidateAllAtLeast(values, minimum, ...)validator; the error message names the requirement ("non-negative" for strides).hasBroadcastAxis(dims, strides)inShapeUtilities.hpp: true when a dimension of size > 1 has stride 0. BothTensor::computeIsPacked()andisTensorPacked()return false when it is true, so broadcast tensors never take linear (packed) fast paths.isPacked()doc comment updated to state the zero-stride rule.🤖 Generated with Claude Code