Repository navigation
fix(hipdnn): keep left_bound under causal_mask in SDPA engines - #12923
Closed
yassinsolim wants to merge 7 commits into
Closed
yassinsolim wants to merge 7 commits into
yassinsolim wants to merge 7 commits into
Conversation
…ndow in SDPA engines A graph with causal_mask (or causal_mask_bottom_right) and a left_bound is a causal sliding window, the way cuDNN's set_causal_mask(true) plus set_sliding_window_length(n) spells one. asm_sdpa_engine's getMaskType let the deprecated booleans win and ignored left_bound, so such a graph was classified as plain causal and could be served with the full causal triangle. hip_flash2_engine gates on the same helper. plan_utils::resolveMask now returns the mask kind plus the band. A deprecated boolean fixes the diagonal (right bound 0) and its alignment and keeps a real left_bound, matching the CPU and GPU SDPA references. getMaskType returns the kind. Effects: - ASM forward: no forward kernel serves a window, so these graphs are declined instead of run as causal. - ASM backward: the SWA kernel gets the resolved bounds (left, 0, the boolean's alignment) instead of the raw right_bound and diagonal_alignment. - hip_flash2: declines the window through its existing mask gate. Its kernel causal flag now comes from the same classification (requestsCausal), so a top-left causal mask spelled as right_bound 0 is masked. It used to read only causal_mask and ran that graph unmasked. Tests: two tests that pinned "boolean beats bounds" for left=64 now assert the window, in the forward and backward suites. New backward tests check the resolved bounds. New TestHipFlash2RequestsCausal cases cover the causal flag. (cherry picked from commit 0641c119c99e75acade1e6ba6eaa9302b030af70)
The CPU reference (extractDiagonalBandParams) and the GPU reference (GpuSdpaFwdPlanBuilder) mapped the deprecated causal_mask and causal_mask_bottom_right flags to left_bound=-1, right_bound=0. That discarded a left_bound set on the same node. A graph built with set_causal_mask(true).set_sliding_window_length(128), the cuDNN spelling of a causal sliding window, was checked against plain causal attention. An engine that computed the window correctly would fail verification and one that ignored it would pass. The flags now fix only the diagonal (right_bound=0 and the alignment) and keep a set left_bound. This matches the rule the gfx950 dense attention pack and the UKD worked example already state. None of the checked-in bundles set causal_mask together with a left_bound, so existing results do not change. New tests run causal_mask=true with left_bound=2 on Sq=Skv=8 through each reference. They compare against the explicit (2, 0, top-left) band, with plain causal as a control. Both fail on the old mapping and pass now. (cherry picked from commit 00bbb9ee3c9e1e9a7f185642d937095ce4670f8c)
With left_bound no longer reset by the deprecated causal flags, clang-tidy's misc-const-correctness (warnings as errors in the hipDNN build) rejects the mutable declaration in extractDiagonalBandParams and GpuSdpaFwdPlanBuilder. (cherry picked from commit 6c951b1c608b4536dda48c44ff96f1a596b763e3)
1 task done
This was referenced Oct 2, 2026
⏳ Waiting on Required Checks⏳ Required checks have not all concluded yet. This bot run is exiting without failing so those checks can report independently.
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
⏳ Waiting for required checks to start or finish:
|
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Backward dispatch can silently narrow 64-bit window bounds into incorrect 32-bit kernel parameters.
Review effort: Balanced
Findings: 1
What changed in this PR
Fixes SDPA causal-mask handling so left_bound remains a sliding-window constraint across engines and references.
Changes:
- Centralizes mask resolution and preserves causal window bounds.
- Correctly sets hip_flash2’s causal kernel flag.
- Adds CPU, GPU, forward, and backward regression tests.
| File | Description |
|---|---|
projects/hipdnn/test_sdk/tests/utilities/cpu_graph_executor/TestSdpaFwdPlan.cpp |
Tests CPU causal sliding windows. |
projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/cpu_graph_executor/detail/PlanUtils.hpp |
Preserves CPU reference left bounds. |
dnn-providers/integration-tests/tests/TestGpuSdpaFwdPlan.cpp |
Tests GPU reference behavior. |
dnn-providers/integration-tests/src/harness/gpu-graph-executor/detail/GpuSdpaFwdPlan.hpp |
Preserves GPU reference left bounds. |
dnn-providers/hip-kernel-provider/src/tests/engines/hip_flash2_engine/TestHipFlash2FwdPlanBuilder.cpp |
Tests causal-flag resolution. |
dnn-providers/hip-kernel-provider/src/tests/engines/asm_sdpa_engine/TestSdpaFwdPlanBuilder.cpp |
Updates forward mask classification tests. |
dnn-providers/hip-kernel-provider/src/tests/engines/asm_sdpa_engine/TestSdpaBwdPlanBuilder.cpp |
Tests resolved backward bounds. |
dnn-providers/hip-kernel-provider/src/engines/hip_flash2_engine/HipFlash2FwdPlanBuilder.cpp |
Derives the kernel causal flag from mask type. |
dnn-providers/hip-kernel-provider/src/engines/hip_flash2_engine/HipFlash2FwdPlanBuilder_v2.hpp |
Declares causal-mask resolution helper. |
dnn-providers/hip-kernel-provider/src/engines/asm_sdpa_engine/plans/SdpaPlanUtils.hpp |
Adds canonical resolved-mask representation. |
dnn-providers/hip-kernel-provider/src/engines/asm_sdpa_engine/plans/SdpaBwdPlanBuilder.cpp |
Passes resolved window parameters to kernels. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #12923 +/- ##
===========================================
- Coverage 71.01% 70.68% -0.34%
===========================================
Files 2944 2940 -4
Lines 485420 484568 -852
Branches 70076 69880 -196
===========================================
- Hits 344710 342480 -2230
- Misses 116771 118193 +1422
+ Partials 23939 23895 -44
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
…ernel left_bound and right_bound are int64 attributes, but the ASM SDPA backward kernel takes int32 window fields. SdpaBwdPlanBuilder cast them directly, so a causal window with left_bound = 4294967298 (2^32 + 2) became a window of 2 instead of the full causal triangle the references compute. plan_utils::kernelWindowBounds now narrows the resolved bounds against the sequence lengths. A bound that reaches the widest offset the band can span (seqLen - 1 on the matching axis for the alignment) masks nothing, so it becomes -1, which computeMaskCoordinates turns back into exactly that offset. Every bound that survives is below a sequence length, so the int32 cast is exact. resolveMask also rejects bounds below -1 with INVALID_VALUE, as the CPU and GPU references already do. Before, -4294967294 would have wrapped to a window of 2 as well. isApplicable in the ASM forward and backward builders already declines on that exception; hip_flash2's isApplicable now catches it too instead of letting it escape. The forward ASM builder and hip_flash2 do not pass bounds to a kernel, so they have no narrowing of their own.
…causal Review follow-ups on the causal sliding window change. Wording. Comments called causal_mask plus left_bound = n the same window as cuDNN's set_sliding_window_length(n). It is not: hipDNN's left_bound counts like flash-attn's window_size_left, so L keeps L + 1 keys per row with the diagonal, while cuDNN keeps L. The engines, both references and the ingestor tests all use the L + 1 count, so nothing here changes behavior. The comments now say so, and the reference tests that say "rows 3..7 lose keys" now say that is the hipDNN count. The frontend forwarder is tracked in #12982. Causal flag with a positive right_bound. resolveMask replaced it with 0 and served a narrower band than the graph asked for. It now throws INVALID_VALUE, which every engine using it turns into a decline, as cuDNN's Python binding and the gfx950 dense pack do. -1 and 0 still agree with the flag. Causal windows that cover the whole sequence. With left_bound honored, causal_mask plus left_bound = 128 at Sq = Skv <= 129 became SLIDING_WINDOW, and the ASM forward and hip_flash2 declined it although it is plain causal and was served before. New plan_utils::resolveMaskFor(attrs, seqLenQ, seqLenKv) narrows the bounds with kernelWindowBounds and classifies the narrowed band. The ASM forward and backward builders and hip_flash2 use it in isApplicable and when building the plan, so both agree on the kernel.
Member
Author
|
@BrianHarrisonAMD thanks for the review. Changes per point, all in
Tests, host only (gfx950 container, CPU-only build, no GPU): |
BrianHarrisonAMD
approved these changes
Oct 4, 2026
BrianHarrisonAMD
left a comment
BrianHarrisonAMD
left a comment
Contributor
There was a problem hiding this comment.
LGTM thanks for the fixes.
Comment on lines
912
to
921
| // Resolve the SDPA attributes from a graph buffer and classify the mask. | ||
| plan_utils::MaskType classifyMask(const flatbuffers::FlatBufferBuilder& builder) | ||
| { | ||
| const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper graphWrapper( | ||
| builder.GetBufferPointer(), builder.GetSize()); | ||
| const auto& attrs = graphWrapper.nodeWrappers() | ||
| .front() | ||
| ->attributesAs<hipdnn_flatbuffers_sdk::data_objects::SdpaAttributes>(); | ||
| return plan_utils::getMaskType(attrs); | ||
| } |
Contributor
There was a problem hiding this comment.
Flagged by AI:
getMaskType is only called here now, and could be removed + RFC17/RFC18 still mention it by name.
We could delete this, and instead swap to resolveMask(...).type + update RFCs.
Member
Author
There was a problem hiding this comment.
Done in 2942b70. getMaskType is gone from SdpaPlanUtils.hpp. classifyMask in this test now returns plan_utils::resolveMask(attrs).type, and the mask-policy comments in TestSdpaFwdPlanBuilder.cpp and TestSdpaBwdPlanBuilder.cpp name resolveMask. RFC 0017 (worked example, section 3) now names resolveMask (the type it returns) as the canonical classifier. RFC 0018's mapping table names resolveMaskFor, which is what SdpaFwdPlanBuilder.cpp calls, with the line reference moved to :451. No other caller existed on this branch, and #12926 and #12927 add none.
Checked on a CPU node with clang-tidy-20 ON (in-tree config, ASM SDPA and hip_flash2 engines on): the 44 TUs that reach the changed files are tidy clean. TestSdpa*:TestHipFlash2* ran 179 (140 passed, 39 skipped for no GPU), and the full hip_kernel_provider_tests passed 1034 with 91 device tests skipped.
Member
Author
There was a problem hiding this comment.
At head 2942b70a, getMaskType is gone from the plan builders, tests, and RFC references. The test helpers use resolveMask(attrs) and resolveMaskFor(attrs, ...); RFC 0017 and RFC 0018 name those entry points. Please resolve this thread if the rename addresses your request.
There is a separate merge-contract conflict with #13065, @jerehartAMD: that branch restores getMaskType, keeps BOTTOM_RIGHT when causal_mask=true has an explicit bottom-right alignment, and accepts a positive right_bound under causal_mask. This branch forces top-left alignment and rejects the positive right bound. Both preserve the left bound. Which contract and branch should own mask resolution? I will keep the competing changes separate until the behavior is agreed.
getMaskType was a one-line wrapper over resolveMask(attrs).type and its only caller left was the mask-policy helper in TestSdpaFwdPlanBuilder. The helper now calls resolveMask directly, the wrapper is gone, and the test comments plus RFC 0017 (worked example) and RFC 0018 (mapping table) name resolveMask / resolveMaskFor instead.
yassinsolim
marked this pull request as draft
October 5, 2026 15:33
Member
Author
|
Holding this (moved to draft) at Jeremy Hart's request while we line it up with #13065, which changes how the frontend and the engines resolve the causal flags. We'll decide together which parts land where before either merges. |
yassinsolim
marked this pull request as ready for review
October 6, 2026 02:45
Member
Author
|
Status: Brian approved this on Oct 4 and all required checks are green. Jeremy and I are confirming the landing order with #13065 before this merges, since both PRs change the SDPA mask handling and touch 11 of the same files. |
Member
Author
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.
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.
Closes #12908
Motivation
causal_mask=trueplus aleft_bounddescribes a causal sliding window (cuDNN also keeps the window next toset_causal_mask(true)). hipDNN'sleft_bound = nkeeps one more key than cuDNN'sset_sliding_window_length(n); that convention gap predates this PR and is tracked in #12982. The ASM SDPA engine and hip_flash2 treated it as plain causal, and both SDPA references did the same, so tests agreed with the wrong answer. hip_flash2 also ran a top-left causal mask unmasked when it was spelled asright_bound=0without the boolean. Details in #12908.Technical Details
Three commits:
fix(hipdnn): treat causal_mask plus left_bound as a causal sliding window in SDPA enginesasm_sdpa_engine/plans/SdpaPlanUtils.hpp: newplan_utils::resolveMask(attrs)returns{type, left, right, topLeft}. A deprecated boolean fixes the diagonal (right=0) and its alignment and keeps a realleft_bound.getMaskTypereturnsresolveMask(attrs).type.ASM forward: a causal window is now SLIDING_WINDOW. No forward CSV row has that mask, so
isApplicabledeclines instead of running plain causal.ASM backward (
SdpaBwdPlanBuilder.cpp): SLIDING_WINDOW params come from the resolved bounds, socausal_mask=true, left=Ldispatches the SWA kernel with(L, 0, top-left), the same path the explicit trio already takes.hip_flash2: the window is declined by the existing NO_MASK / TOP_LEFT_CAUSAL gate. New
HipFlash2FwdPlanBuilder::requestsCausal(attrs)derives the kernel causal flag from the same classification.extractParamsuses it, so a bounds-spelled top-left causal mask is masked.fix(hipdnn): keep left_bound when SDPA references see causal_mask: the CPU reference (extractDiagonalBandParams) and the GPU reference (GpuSdpaFwdPlanBuilder) now set onlyrightBound = 0and the alignment for a causal flag and keep a setleft_bound. Unsetleft_boundis still -1, so plain causal is unchanged.fix(hipdnn): declare the SDPA reference left_bound const: onceleftBoundis no longer reassigned, clang-tidymisc-const-correctness(warnings as errors) fails every test_sdk TU that includesPlanUtils.hpp. The GPU reference has the same pattern.Review follow-ups.
1051f6cd93: int64 window bounds are narrowed to the backward kernel's int32 fields without wrapping, and bounds below -1 are declined.27f740ee1d: a causal flag next to a positiveright_boundis declined (it used to be forced to 0); newplan_utils::resolveMaskFor(attrs, Sq, Skv)classifies the mask at the graph's sequence lengths, so a causal window that covers the whole sequence runs as plain causal in the ASM forward, ASM backward and hip_flash2 instead of being declined; comments no longer equateleft_bound = nwith cuDNN'sset_sliding_window_length(n).Tests:
IsApplicablePrefersCausalMaskOverWindowBoundsandIsApplicablePrefersBottomRightCausalOverWindowBounds(fwd and bwd) asserted thatleft=64with a boolean is plain causal, which is the bug. They now assert the window. New tests: backward resolved bounds,TestHipFlash2RequestsCausal, andDeprecatedCausalMaskWithLeftBoundIsSlidingWindowin the CPU and GPU reference suites (window must equal the explicit(2, 0, TOP_LEFT)band, with plain causal as a control).What could change: only graphs that set a causal boolean together with a
left_bound, or hip_flash2 graphs that spell top-left causal as bounds, and graphs that set a causal boolean with a positiveright_bound, which are now declined. Every ASM integration case that sets a causal boolean usesleft_bound=-1, and no integration-test bundle in this repo sets a causal flag with aleft_bound(455 checked-in graph JSON files, 11 with a causal flag), so existing bundle results do not change. The dnn-benchmarking workload corpus does have that shape: the fivegpt_oss__..._causal_swa128graphs that RFC 0017's worked example mentions setcausal_mask=true,left_bound=128and noright_bound. At their lengths (Sq = Skv = 2048 to 32768) that is a real window, so these engines now classify them as a window and decline them, where before they classified them as plain causal.The
causal_mask+diagonal_alignmentquestion (#12909) is not part of this PR. The agreed direction is for the frontend's causal setters to mirror cuDNN (alignment plusright_bound = 0, no deprecated booleans), tracked in ALMIOPEN-2586 (Jeremy Hart). #12924 is closed. Theleft_boundoff-by-one against cuDNN is #12982.Test Plan
Engine commit, host unit tests (gfx950 build with
-DENABLE_HIP_FLASH2_ENGINE=ON): new tests against the old classification (before) and the fix (after), then everyTestSdpa*:*Flash2*:*AsmSdpa*host test.Reference commit: the new CPU test built against the develop test_sdk headers and the fixed ones. The new GPU reference test on an MI355X, pre-fix and fixed.
GPU checks on gfx942 (MI325X) with throwaway test changes (not committed): the flash2 integration cases spelled as
right_bound=0, TOP_LEFT, with the branch and with the oldrequestsCausal. ASM backward window cases spelled with the boolean plusleft_bound=64, top-left at Sq=Skv=256 and bottom-right at Sq=256/Skv=512, against the CPU backward reference.Full suites on gfx950 (MI355X) and gfx942 (MI325X).
Update 2026-10-03: this branch on its own (head 3c85f83, develop merged in) was built and run on a gfx942 MI325X with ASM SDPA and hip_flash2 enabled; unit, reference and bundle-tier verdicts match develop test by test, the new GPU reference test passes (it fails against develop), and end-to-end causal-window graphs that develop served wrongly now match the CPU and GPU references or are declined. The full-stack numbers that follow were taken earlier with the now-closed alignment commit on top and are kept for reference.
Test Result
Engine commit, before: 5 new tests FAILED (
CausalMaskWithLeftBoundIsSlidingWindow,BottomRightCausalMaskWithLeftBoundIsSlidingWindowgot causal; the two backward window tests got the wrong type, right bound 64 and the wrong alignment;BoundsSpelledTopLeftCausalSetsCausalgot false). After: all pass. AllTestSdpa*:*Flash2*:*AsmSdpa*host tests: 160 passed, 0 failed (155 passed, 5 failed before).CPU reference test: fails on both assertions with the develop headers, passes with the fix (13/13 in
TestSdpaFwdPlanBuilder). GPU reference test on MI355X: fails pre-fix on both assertions, passes with the fix.gfx942 flash2, bounds spelling: 6/6 pass with the branch. With the old
requestsCausalthe 5 causal cases fail against the reference and the non-causal case passes.gfx942 ASM backward, boolean spelling: 20/20 pass (bf16 and fp16), same as the explicit trio spelling (20/20). An engine that ignores
left_boundunder a boolean fails the bottom-right window case.Full stack on gfx950 (MI355X): engine filters 235 passed. CPU reference
*Sdpa*124 passed. GPU reference (all ofhipdnn_integration_tests_unit_tests) 889 passed. Installed provider integration*Sdpa*:*Flash2*:*Gfx950AttentionDense*: 31 passed, 5 failed. The 5 are the gfx950 Dqk192 launch failures on develop, fixed separately in the hqs-asm PR (hip-kernel-provider: gfx950 ASM SDPA Dqk192/Dv128 bf16 launch fails #12910).Full stack on gfx942 (MI325X, not MI300X): engine filters 257 passed, 1 skipped. CPU reference 124 passed. GPU reference 889 passed. Installed provider integration: 38 passed, 36 skipped (gfx950-only cases).
A merged build of this stack with other runbook fixes, clang-tidy ON, on MI355X: build and install pass, and hipDNN test_sdk, hip-kernel-provider unit and integration-harness unit tests pass. It did not include the ASM backward bundle or RMSNorm claim changes.
Review follow-ups, host only (gfx950 container, CPU-only build with the ASM SDPA and hip_flash2 engines on, no GPU):
TestSdpa*,TestHipFlash2*andTestGfx950AttentionDense*287 passed, 39 skipped (they need a device); all ofhip_kernel_provider_tests1034 passed, 91 skipped. With the right-bound check and the sequence-aware reclassification removed, the 4 new tests fail. No GPU run of the follow-ups.clang-format 18 clean on every touched file.
Submission Checklist