Repository navigation
fix(hipdnn): Set Sdpa diagonal alignment and left/right_bounds correctly - #13065
Merged
Merged
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. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13065 +/- ##
========================================
Coverage 76.02% 76.02%
========================================
Files 2869 2869
Lines 475305 475316 +11
Branches 71207 71211 +4
========================================
+ Hits 361319 361335 +16
+ Misses 89410 89404 -6
- Partials 24576 24577 +1
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
1 task done
jerehartAMD
force-pushed
the
users/jerehart/causal-mask-override
branch
from
October 7, 2026 15:22
92a1fc0 to
da64482
Compare
jerehartAMD
marked this pull request as ready for review
October 7, 2026 15:32
DarylHawkinsAMD
approved these changes
Oct 8, 2026
DarylHawkinsAMD
left a comment
DarylHawkinsAMD
left a comment
Contributor
There was a problem hiding this comment.
Looks good, just one question inline. Good to merge if we don't need to make a change there.
BradPepersAMD
added a commit
that referenced
this pull request
Oct 14, 2026
945c10b declined graphs that set a deprecated causal boolean together with a left_bound or another right_bound, on the grounds that hipDNN's CPU reference drops those bounds. That reversed a deliberate choice: gpt-oss SWA, DeepSeek-V4 and Mistral graphs set the boolean with a sliding window, and the engines keep the window as cuDNN and #13065 do. Declining them would have cost every such graph on every arch. Restored: under a boolean, an explicit right_bound wins over its 0 and a left_bound is kept. Kept from 945c10b: both booleans together and an unnamed diagonal alignment decline. Tests now accept the boolean with a window (RDNA, CDNA) and with a band (RDNA), and the binding test checks the window's left bound survives. FlyDSL gtests 259 passed / 44 skipped on gfx1151. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
…sal-mask-override
jerehartAMD
force-pushed
the
users/jerehart/causal-mask-override
branch
from
October 15, 2026 04:56
da64482 to
da03c7c
Compare
This was referenced Oct 15, 2026
…sal-mask-override
This was referenced Oct 21, 2026
Closed
Contributor
TheRock Submodule Bump ActivityNewest first
|
BradPepersAMD
added a commit
that referenced
this pull request
Oct 27, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 27, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 28, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 28, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 28, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 28, 2026
…ound (#14483) ## Summary hipDNN's SDPA `left_bound` (and `set_sliding_window_length`) now means what it means in cuDNN, whose API hipDNN mirrors: **`left_bound = L` keeps the L keys ending at the diagonal, diagonal included.** Until now every hipDNN component kept **L + 1** keys (FlashAttention's `window_size_left` convention), so a graph written against the cuDNN contract attended to one extra key per query, forward and backward. JIRA ID : ALMIOPEN-3024 This supersedes #14287 by Daming Feng: his commit is included unchanged (same patch, authorship kept), rebased on current `develop`, with follow-ups for code that landed since and a generator flag that names each convention explicitly. ## Description cuDNN's left bound is a `CMP_GT` diagonal-band mask that keeps `col + L > row` (cudnn-frontend `attn_score_modifiers.h` `sliding_window_mask`; reference `test/python/sdpa/fp16_ref.py` masks `rel <= diag_offset - left_bound`). Example, `L = 2`, query 5, TOP_LEFT: cuDNN keeps keys {4, 5}; hipDNN kept {3, 4, 5}. The right bound already matched. Nothing caught it because every check shared hipDNN's convention: the golden generator built its own mask and cross-checked against AITER with the same value, the engine tests compared against our own CPU reference, and no golden bundle uses a left bound. | Convention | Name | `L = 3`, query on the diagonal | |---|---|---| | hipDNN graph (now), cuDNN | `left_bound`, `set_sliding_window_length` | the 3 keys ending at the diagonal | | vLLM | `sliding_window = W` | W keys including the token (`left_bound = W`) | | FlashAttention, AITER, TransformerEngine, FlashInfer | `window_size_left` / `window_left` | the 3 keys before the diagonal plus the diagonal (`left_bound = W + 1`) | Changes: - **References:** `CpuFpReferenceSdpa` (forward and backward), `CpuFpReferenceSdpaRagged`, `GpuRefSdpaFwd` and `GpuRefSdpaRaggedFwd` mask `skv <= sq + off - L`. - **Frontend:** `detail::validateSdpaLeftBound` in `SdpaFwdNode` / `SdpaBwdNode` rejects `left_bound = 0` and values below -1. Doxygen on `set_diagonal_band_left_bound` and `set_sliding_window_length` states the cuDNN meaning and the FlashAttention mapping. - **Graph plans:** the CPU executor's `extractDiagonalBandParams` and the GPU `GpuSdpaRaggedFwdPlan` reject 0. - **hip-kernel-provider:** `plan_utils::resolveDiagonalBand` is the single conversion point; `DiagonalBand::leftBound` holds the kernels' FlashAttention window (`L - 1`) for asm/AITER forward and backward, hip_flash2 and the gfx950 dense pack (a window that exactly covers the sequence is now correctly classified as causal). `Gfx1151WmmaAttentionNative::maskFor` (prefill and decode) converts the same way and declines 0. Kernels and rocKE keep their own conventions. - **Golden generator:** `generate_sdpa_fwd_golden.py` gains `--left-bound L` (graph / cuDNN convention, written unchanged). `--window-left W` keeps FlashAttention's convention (written as `left_bound = W + 1`) for `--validate` against AITER. The two are mutually exclusive; the README records the conversion. No checked-in bundle uses a left bound, so no golden data changes. - **IngestorGenerator shape miner:** the graph's `left_bound` and rocKE's `sliding_window` count keys the same way, so the value passes through unchanged (was `left + 1`), and 0 is refused. - **Tests added since #14287:** the gfx1151 decode-engine window tests (#13897) use `left_bound = 1024` for their 1024-key window; a heuristics test no longer builds a `left_bound = 0` graph; the CPU and GPU plans' rejection of 0 is tested directly. An audit of every other file on `develop` that mentions `left_bound` found no other consumer of the raw value: the remaining plans and engines take the band already converted by `resolveDiagonalBand` / `extractDiagonalBandParams`, and the rest carry the field through (descriptors, packers, schema, Python bindings, corpus tools). ## Behavior change and migration - A graph with `left_bound = L` now attends to one fewer key than before. Callers that passed FlashAttention-style values straight into `left_bound` must pass `W + 1`. Out-of-tree providers that read the graph's `left_bound` must convert to their kernel's convention (FlashAttention-style kernels take `L - 1`). - `left_bound = 0` and values below -1 are rejected; -1 still means unbounded. (cuDNN's frontend rejects a forward `left_bound <= 0` only on backends before 9.10 and a backward one always; hipDNN rejects 0 in both directions, because a provider's `L - 1` would turn 0 into -1, which kernels read as unbounded.) - The serialized field is unchanged. SDPA was off by default in every hipDNN build in TheRock until less than a week ago, so no released build has users of the old meaning; whether to bump the `flatbuffers_sdk` / `frontend` versions is being discussed separately and is not part of this PR. ## Risk Assessment Low. It is an intentional behavior change to a public SDPA attribute, but SDPA was off by default in every hipDNN build in TheRock until less than a week ago, so no released build has users relying on the old meaning. The in-tree consumers all change together in this PR. | Risk | Mitigation | |---|---| | A provider left on the old convention silently attends to one extra key | References and every in-tree provider mapping change together in this PR, so the engine-vs-reference tests check them against each other (a provider on the old convention fails them). Conversion is centralized in `resolveDiagonalBand` (and gfx1151's `maskFor`), each with tests that pin `L -> L - 1`, `1 ->` diagonal only, `0 ->` declined. | | The reference and its tests share the misreading again | `LeftBoundKeepsLKeysIncludingDiagonalLikeCudnn` pins the cuDNN rule by hand; the generator's mask is compared against an independent mask built from cudnn-frontend's `fp16_ref.py` rule (top-left, bottom-right, `S_q != S_kv`). | | Out-of-tree providers or serialized graphs written to the old meaning | Migration note above; Doxygen and the generator README state the convention; versioning decision tracked separately. | | Tools and generators mixing the two conventions | Each generator flag is named after the convention it uses (`--left-bound` vs `--window-left`); the shape miner and rocKE agree on counting the diagonal. | | Code that landed on `develop` after #14287 still uses the old values | Audited every `left_bound` consumer on current `develop`; the gfx1151 decode tests (#13897) were the only ones affected and are updated. | | Work in flight assumes the old meaning | #14358 (SDPA reference implementations) assumes the cuDNN meaning and stacks on this PR; the FlyDSL attention engines (#14111) are validated against these references, so they land after it. | ## ASIC Coverage The semantic change is arch-independent. Affected engines: asm/AITER and hip_flash2 (gfx942/gfx950) and the gfx1151 WMMA pack (gfx11/gfx12). The gfx950 dense pack declines windows and is unaffected. Validated on gfx942 (below); gfx950 and gfx1151 have no device run in this PR: their changes are the shared band resolution and `maskFor`, which the host-side tests cover. ## Testing New and updated tests: - `TestCpuFpReferenceSdpaFp64.LeftBoundKeepsLKeysIncludingDiagonalLikeCudnn` (L = 2, query 5 keeps {4, 5}; L = 1 keeps only the diagonal); existing sliding-window tests keep their hand-checked masks with `leftBound` raised by one. - `TestSdpaDiagonalBandParams.RejectsZeroLeftBound` (CPU plans) and a `left_bound = 0` case in `TestGpuSdpaRaggedFwdPlanBuilder.BuildRejectsInvalidMaskAttributes` (GPU plan). - hip-kernel-provider: `TestSdpaFwdResolveDiagonalBand`, `LeftBoundZero` / `LeftBoundOneIsDiagonalOnly` mask cases, covering-window boundaries (backward 64, hip_flash2 256 vs 255), gfx1151 binding (7 binds as 6, 1 as 0, 0 declined) and the decode-engine window cases. - `test_generate_sdpa_fwd_golden.py`: the generator's mask against cudnn-frontend's rule; `--left-bound 3` and `--window-left 2` give the same graph; invalid flag combinations are rejected. Results, gfx942 (MI300X), this branch built with `-Werror`, `HIPDNN_ENABLE_SDPA=ON`, `HIPDNN_ENABLE_KERNEL_INGESTOR=ON`: | Suite | Passed / skipped | |---|---| | hipdnn_test_sdk_tests | 2423 / 1 | | hipdnn_frontend_tests | 2222 / 0 | | hipdnn_public_frontend_tests | 490 / 0 | | hipdnn_plugin_sdk_tests | 1293 / 2 | | hipdnn_flatbuffers_sdk_tests | 269 / 0 | | hipdnn_integration_tests_unit_tests | 1109 / 0 | | hipdnn_gpu_ref_tests | 13772 / 0 | | hipdnn_golden_data_tests | 66 / 0 | | hip_kernel_provider_tests (`*Sdpa*:*Flash2*:*Gfx950Attention*:*Gfx1151WmmaAttention*`) | 420 / 1 (gfx950-only) | | hip_kernel_provider_integration_tests (same filter; asm forward/backward incl. backward window cases) | 45 / 21 (gfx950-only, and causal+stats graphs the asm forward builder declines by design) | | integration-tests quick CTest | 15 / 15 | | golden generator pytest, with PyTorch | 68 passed, 23 subtests | The CI CTest for the generator tests has no PyTorch in its environment and skips them; the PyTorch run above is what exercises the generator. ## Related - Tracking: ALMIOPEN-3024. Related earlier fix of the deprecated causal setters: #13065 (ALMIOPEN-2586). - Supersedes #14287 (Daming Feng). - Must land before #14358 (SDPA CPU and GPU references), which assumes the cuDNN meaning, and before the FlyDSL attention engines (#14111), which are validated against these references. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Daming Feng <dmfeng8898@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 29, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
yassinsolim
added a commit
that referenced
this pull request
Oct 29, 2026
This branch is based on develop commit `e0b0ad5f6d`, which includes #13065. The mask-resolution changes in #13065 are the reason some of these checks are needed now: the sliding-window bounds reach an int32 kernel field, and the mask classification now reads the sequence lengths. ## What changed The ASM SDPA kernarg structs store every dimension and stride as `uint32_t`, the launch grids take `unsigned int`, and the kernel registries are keyed on `int`. The plan builders checked the byte strides and nothing else, so a batch size, head count, sequence length or head dim that did not fit its field was narrowed silently and the kernel ran on a different shape than the graph asked for. ### Dimension guards `plan_utils::dimFitsU32` and `plan_utils::dimFitsI32` sit next to the existing `byteStrideFitsU32` and follow its shape: log the offending field, return false, let the caller report every violation of a tensor in one pass. Both ASM plan builders call a per-pass wrapper from `isApplicable` immediately after the rank checks, which is ahead of the first narrowing cast in each builder. | value | limit | why that limit | | --- | --- | --- | | batch, nhead_q, nhead_k | `UINT32_MAX` | uint32 kernarg field or unsigned grid dimension, no other narrowing | | seqlen_q, seqlen_k | `INT32_MAX` | the backward casts both into `computeMaskCoordinates`, which is int32 throughout, and both passes compute `(extent + tile - 1) / tile` in `unsigned int`, which wraps near `UINT32_MAX` | | head_dim_q, head_dim_v | `INT32_MAX` | both builders cast them to `int` for the kernel-registry lookup before they reach the uint32 field | `INT32_MAX` on the sequence lengths is also what makes the window narrowing in `SdpaBwdPlanBuilder` exact. `makeCoveringBoundsUnbound` replaces any bound that reaches its span with -1, so a surviving bound is strictly below a sequence length, and every `1 + bound + offset` that `computeMaskCoordinates` forms then stays inside int32. The head dims are the case with teeth. `D_qk = 2^32 + 128` truncated to 128, resolved the hd128 kernel and dispatched it, in both passes. The registry sets ({64, 128, 192} for the backward) are what bound the head dims in practice; the guard is there so the cast that feeds those sets is well defined. ### byteStrideFitsU32 It formed `elements * elementBytes` before testing the product, so an element count near 2^62 overflowed int64 and landed back inside the uint32 range, and the log line on the failing path printed that wrapped product. It now compares `elements <= UINT32_MAX / elementBytes`, rejects a non-positive element size instead of dividing by it, and prints the largest element count that would fit. ### Ts `Ts` is the one kernarg that is a product of a stride and a tile size. The backward checked it but built `tsKv * stride_k` in int64 before calling the predicate, and the forward had no check at all even though `s_Ts = tileSizeQo * qStrideRow * inBytesPerElement` is computed in `unsigned int` in `SdpaFwdArgsBuilder`. Both now fold the tile size into the element size argument, so the product is never formed outside the predicate, and the forward resolves `ts_qo` from the row it already looked up. The forward element size is the resolved input size, one byte for fp8 and two for bf16, read from the same `dataTypeId` `buildPlan` uses; the per-field q/k/v stride check takes it too, and O stays two bytes. ### isApplicable and buildPlan now classify on the same values `buildPlan` re-read the sequence lengths as `unsigned int` and classified the mask on the truncated values while `isApplicable` classified on the graph dims. A sequence length above `UINT32_MAX` could therefore be approved as one mask type and dispatched as another. Both sites read the dims now, and the sequence length guard removes the case entirely. Before #13065 the classifier ignored the sequence lengths, so the two sites could not disagree. ### hip_flash2 `Flash2FwdParams` narrows every dimension to `int`, and `batch` had no check anywhere: `isApplicable` never read it, so a batch of `2^32 + 1` was dispatched as 1. The same dimension gate now runs in that engine's `isApplicable`, ahead of its casts, with `INT32_MAX` for all six dims because they are all `int`. `seq_len_kv` stops one KV tile lower, at `INT32_MAX - 64`: both kernels count KV tiles in signed `int`, and 64 is the largest KV tile the engine dispatches. Its stride handling is untouched: `checkedStride` already rejects out-of-range strides, and it throws at execute rather than declining at applicability time, which is a behaviour difference worth a separate decision. ## Inventory The full field inventory is attached to this PR rather than committed to the branch. It lists every field of all four kernarg structs with the expression that fills it, the file and line at `f86ff449`, the largest value that keeps its meaning, and the guard that covers it. Summary: 146 fields, 104 of them 32 bits wide, of which 35 were already guarded, 29 are guarded by this change, and 40 are bounded by construction (constants, hardwired zeros, values derived from an already-guarded value) or are 64-bit pointers and floats. ## Testing Rebuilt the develop tree `e0b0ad5f6d` and branch `feef30d804` on gfx942 and gfx950 with the same Release/Ninja configure. The gfx942 run used an MI325X node in partition `miopen`. The gfx950 run used an MI355X node in partition `miopen`; the rocke MI355X nodes had no free GPU at that time. Both jobs used account `miopen`. On gfx942: - Whole `hip_kernel_provider_tests`: develop, 898 ran, 891 passed, 7 skipped; branch, 928 ran, 921 passed, 7 skipped. No failures. - The forward filter: develop, 94 passed; branch, 109 passed. - The backward filter: develop, 99 ran, 98 passed and 1 skipped; branch, 103 ran, 102 passed and 1 skipped. The skip is the gfx950-only case. - The flash2 filter: develop, 34 passed; branch, 45 passed. - `hipdnn_integration_tests_unit_tests --gtest_filter=TestGpuSdpaFwdPlanBuilder.*`: develop, 15 passed; branch, 16 passed. On gfx950: - Whole `hip_kernel_provider_tests`: develop, 898 ran, 869 passed, 29 skipped; branch, 928 ran, 884 passed, 44 skipped. No failures. - The forward filter: develop, 94 ran, 92 passed and 2 skipped; branch, 109 ran, 107 passed and 2 skipped. Both skips are the existing fp8 descale cases. - The backward filter: develop, 99 ran, 87 passed and 12 skipped; branch, 103 ran, 87 passed and 16 skipped. The four additional skips are the new backward boundary test methods, each guarded to run only on gfx942. - The flash2 filter: develop, 34 ran, 25 passed and 9 skipped; branch, 45 ran, 25 passed and 20 skipped. All 11 new `TestHipFlash2DimRange` cases skip at the existing gfx942-only device guard. - `hipdnn_integration_tests_unit_tests --gtest_filter=TestGpuSdpaFwdPlanBuilder.*`: develop, 15 passed; branch, 16 passed. The new forward dimension and Ts cases, all 11 device-free `TestSdpaKernargRanges` tests, and the ported sliding-window GPU-reference test ran and passed on both GPU architectures. All four backward boundary methods and all 11 flash2 dimension-range cases ran and passed on gfx942. They are skipped on gfx950 because those engines' applicability paths currently accept only gfx942. No test failure was unique to the branch on either architecture. New cases: - `TestSdpaKernargRanges`, 11 cases, device-free: each predicate at its exact limit and one past it, the negative-dimension case, the tile-factor case, and the 2^62 case that the old multiply-first form accepted. - `TestSdpaFwdPlanBuilder` and `TestSdpaBwdPlanBuilder`: each dimension at the limit of its field (accepted) and one past it (declined), the head dims at a value whose low 32 bits are 128 (declined, accepted before this change), and the Ts limit, with the tile size read from the kernel registry so a registry change fails the test instead of weakening it. Descriptor values only, no allocation. - `TestHipFlash2DimRange`, 11 parameterized cases, same shape. - `TestGpuSdpaFwdPlanBuilder.DeprecatedCausalMaskWithLeftBoundIsSlidingWindow`, ported from #12923: `causal_mask=true` with `left_bound=2` runs as `leftBound=2, rightBound=0` against the GPU reference, with a control that the window actually changes the output. It holds under #13065's contract because `makeCoveringBoundsUnbound` leaves a bound of 2 alone at Sq = Skv = 8. Two backward cases cannot be isolated to the new guard: `nhead_q` and `seqlen_q` past their limits are also rejected by the stats stride guard, because the test graph utility derives the stats strides from the Q dims. They are still asserted as declined; `seqlen_k`, batch and both head dims are isolated, and the predicate suite covers the boundary exactly. Negative control, same node and configure line, with the three plan builders restored to `f86ff449` and `byteStrideFitsU32` restored to its multiply-first form (the two new predicates stay so the test TUs still compile): 13 of the 30 selected cases fail. The failures are exactly the cases that assert a decline, plus the two `byteStrideFitsU32` cases the old body accepted. The cases that assert acceptance pass in both builds, so the guards decline what they are meant to decline and nothing else. One case passes in the control and is worth naming: the backward `IsApplicableBoundsTheTsProductByTheResolvedTileSize`. The backward Ts guard already rejected that stride before this change; reordering the multiply only matters when `tsKv * stride_k` itself overflows int64, which the per-field stride guard reaches first. That case is characterization, not a regression test. The forward Ts case does fail in the control, because the forward had no Ts guard at all. ## What could break Graphs the ASM engines used to accept and now decline: any graph with a dimension above its field width. Those were already miscomputing, so declining them moves the work to another engine instead of returning wrong numbers. The realistic blast radius is zero, because a tensor with a dimension above 2^31 also needs strides that the existing byte-stride guard already rejects, unless the strides are degenerate (broadcast or zero). The forward sequence-length limit is `INT32_MAX` rather than the field's `UINT32_MAX`. That is the stricter of the two defensible limits for the forward pass alone. The argument for it is in the table above (the grid ceil-division), and it keeps one rule across both passes. Reviewers who want the forward to accept up to `UINT32_MAX` should say so; the alternative is a tile-size dependent bound. `byteStrideFitsU32` now declines `elementBytes <= 0`. No current caller passes one, and a tile size of 0 would mean an unresolved registry row. ## Follow-up Commit `6b5c2ab2a0`, on top of a merge of current develop. Three changes. 1. `TestSdpaFwdPlanBuilder.cpp` formed `2 * 192` in `int` before widening it to `int64_t`, which `bugprone-implicit-widening-of-multiplication-result` rejects and the Linux superbuild treats as an error. It reads `int64_t{2} * 192` now. 2. The flash2 `seq_len_kv` ceiling, from review. The engine accepted `INT32_MAX` while both kernels count KV tiles in signed `int`, so the count overflowed at the accepted boundary. `K_FLASH2_MAX_SEQ_LEN_KV` is `INT32_MAX - 64` and the builder and the test read the same constant. The acceptance case moved there, with a one-past case next to it; `TestHipFlash2DimRange` is 12 cases. 3. The forward input element size, from review. The Ts bound and the q/k/v per-field stride check assumed two bytes, so an fp8 descriptor the kernel can run was declined. `fwdInputBytesPerElement` derives the size from the resolved `dataTypeId` and `buildPlan` fills `inBytesPerElement` from the same function. `IsApplicableBoundsTheTsProductByTheResolvedTileSize` gained fp8 boundary cases beside the bf16 ones, with `ts_qo` read from the `fp8bf16` registry row. clang-tidy 20.1.2 with `projects/hipdnn/.clang-tidy` is clean over all seven translation units this PR touches, and clang-format 18.1.3 is clean over all nine files. Rebuilt and rerun on an MI300X (gfx942) and an MI355X (gfx950), same Release/Ninja configure as above. - Whole `hip_kernel_provider_tests`: gfx942, 929 ran, 922 passed, 7 skipped; gfx950, 929 ran, 884 passed, 45 skipped. No failures. - Forward filter: gfx942, 109 ran, 109 passed; gfx950, 109 ran, 107 passed and 2 skipped (the existing fp8 descale cases). - Backward filter: gfx942, 103 ran, 102 passed and 1 skipped; gfx950, 103 ran, 87 passed and 16 skipped. - flash2 filter: gfx942, 46 ran, 46 passed, including all three `seq_len_kv` cases; gfx950, 46 ran, 25 passed and 21 skipped at the existing gfx942-only device guard. - `hipdnn_integration_tests_unit_tests` filtered to `*GpuSdpaFwdPlan*`: 16 passed on both. - `hip_kernel_provider_integration_tests` filtered to `*SdpaFwd*` on gfx942: 50 ran, 29 passed, 21 skipped (the gfx950 rows). The three `fp8bf16` GPU cases pass, so the fp8 dispatch path is unchanged. The counts above are one more test than the earlier table because the single `seq_len_kv at INT32_MAX` case became two.
BradPepersAMD
added a commit
that referenced
this pull request
Oct 29, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 30, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 31, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 31, 2026
cuDNN's left_bound = L keeps the keys j > i + off - L: L keys counting the diagonal. Every hipDNN component kept j >= i + off - L, one key more, and because they all agreed no test caught it. - CPU dense and ragged references, GPU forward/ragged/backward kernels and the golden generator (1.5.0) use the cuDNN rule. - Kernels count the keys left of the diagonal (FlashAttention window_size_left), so the provider converts once: resolveDiagonalBand stores L - 1 (asm forward/backward, hip_flash2) and the gfx1151 WMMA pack passes L - 1 as window_left. - left_bound = 0 keeps no keys; like cuDNN, the frontend nodes, the reference plans and the provider reject it. - The gfx1151 pack now composes the deprecated causal flags with the bounds (#13065 changed every other consumer; this pack still let them override). Co-Authored-By: Claude Opus 5.5 <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
set_causal_mask()/set_causal_mask_bottom_right()only set deprecated booleans, which overrode the diagonal-band bounds andalignment. That diverged from cuDNN, where these setters write
right_bound = 0plus the matchingdiagonal_alignment. This PR matchescuDNN's setter semantics and sends every provider and reference consumer through one shared mask resolution.
JIRA ID : ALMIOPEN-2586
Risk Assessment
3/5. This changes how SDPA masks are classified and dispatched in ASM SDPA fwd/bwd, HipFlash2 fwd, the gfx950 kernel-ingestor
dense-attention pack, and the test_sdk CPU/GPU reference plans. A wrong mask fails silently as wrong attention output, so
golden-numerics integration tests are the backstop. There is no schema or ABI change: the deprecated FlatBuffer booleans are still
serialized, and only how they are interpreted changes.
Intentional behavior change for graphs that set the deprecated booleans directly (persisted graphs, or Python
attrs.causal_mask = True): the booleans are now merged with the stored band instead of overriding it.causal_mask+ BOTTOM_RIGHT alignmentcausal_mask+left_bound=64causal_mask+right_bound=5ASIC Coverage
Only gfx942 and gfx950 are affected: ASM SDPA (gfx942/gfx950), HipFlash2 (gfx942), and the kernel-ingestor dense-attention pack
(gfx950). PR CI alone is not enough, because the gfx950-only integration tests skip on gfx942. Verification needs a TheRock multi-arch
run covering gfx942 and gfx950. A full sweep is not needed.
Testing Summary
falseis a no-op, and later bound/alignment setters take precedence (C++ andPython bindings).
DeprecatedCausalMaskCasetable covers ASM fwd/bwd classification, HipFlash2extractParams().causal, and the gfx950maskTypeFor. The table also checks that setting both booleans, or a bound below -1, isrejected.
extractDiagonalBandParamsmerges explicit bounds in the CPU fwd/bwd reference plans.Testing Checklist
hip_kernel_provider_tests- Status: Pending re-run on HEADhipdnn_test_sdk_tests- Status: Pending re-run on HEADhipdnn_frontend_tests- Status: Pending re-run on HEADpytest frontend_wheel_package/tests/test_sdpa.py- Status: Pendinghip_kernel_provider_integration_tests --gtest_filter='*Sdpa*:*Flash2*:*Gfx950AttentionDense*'- Status: Pending re-run on HEADTechnical Changes
SdpaAttributes.hpp,SdpaBackwardAttributes.hpp):set_causal_mask*()now setsdiagonal_alignmentandright_bound = 0, matching cudnn-frontend, and no longer writes the deprecated members. Those members are documented as kept only for persistedgraphs.
SdpaPlanUtils.hpp): newresolveDiagonalBand()merges the deprecated booleans as if their setter ran first.An unset right bound or one of -1 becomes 0,
causal_mask_bottom_rightforces BOTTOM_RIGHT, andleft_boundis kept. Setting bothbooleans, or a bound below -1, throws.
getMaskType()now classifies the resolved band.causalnow comes fromgetMaskType(). It used to read the raw boolean, soset_causal_mask(true)would otherwiserun non-causal.
isApplicablenow declines when mask resolution fails.maskTypeFor: now delegates togetMaskType()instead of keeping its own divergent copy of the classifier.PlanUtils.hpp,GpuSdpaFwdPlan.hpp):extractDiagonalBandParamsapplies the same merge, and the GPU plan'sduplicated precedence logic is removed.