Repository navigation
fix(hipdnn): range-check every SDPA kernarg narrowed to 32 bits - #14075
Merged
Merged
Conversation
Each SDPA dimension reaches the ASM kernels through a 32-bit field: the kernarg structs store them as uint32, the launch grid takes unsigned int, and the kernel registries are keyed on int. Batch, the head counts and the sequence lengths had no magnitude check at all, and the head dims were cast to int for the registry lookup unchecked, so a dimension whose low 32 bits were a supported value resolved a kernel built for that value. Add dimFitsU32 / dimFitsI32 next to byteStrideFitsU32 and gate both ASM plan builders on them before the first narrowing cast, with the same field-by-field logging the stride guard uses. Give hip_flash2 the same gate: it narrows every dim to int and never looked at batch at all. byteStrideFitsU32 formed the product before testing it, so an element count near 2^62 overflowed into range and the log line printed the wrapped product; it now divides instead. Add the forward Ts guard the backward already had, and fold the backward tile size into the element size so that product is not formed outside the predicate either. buildPlan re-read the sequence lengths as unsigned int and classified the mask on the truncated values while isApplicable classified on the graph dims; both now read the dims. Port TestGpuSdpaFwdPlanBuilder.DeprecatedCausalMaskWithLeftBoundIsSlidingWindow from #12923: causal_mask plus left_bound is a sliding window, checked against the GPU reference.
✅ 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. |
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Flash2 still permits overflowing tile-count arithmetic, and the forward Ts check unnecessarily rejects valid FP8 strides.
2 open issues
What changed in this PR
Adds range checks to hipDNN’s SDPA plan builders to prevent silent narrowing of dimensions and strides into 32-bit kernel arguments.
Changes:
- Adds dimension guards for ASM SDPA and hip_flash2.
- Makes stride checks overflow-safe and validates tile-scaled strides.
- Adds boundary tests and a sliding-window reference regression test.
| File | Description |
|---|---|
| dnn-providers/integration-tests/tests/TestGpuSdpaFwdPlan.cpp | Tests deprecated causal flags with sliding windows. |
| dnn-providers/hip-kernel-provider/src/tests/engines/hip_flash2_engine/TestHipFlash2FwdPlanBuilder.cpp | Adds Flash2 dimension boundary tests. |
| dnn-providers/hip-kernel-provider/src/tests/engines/asm_sdpa_engine/TestSdpaFwdPlanBuilder.cpp | Tests range predicates and forward limits. |
| dnn-providers/hip-kernel-provider/src/tests/engines/asm_sdpa_engine/TestSdpaBwdPlanBuilder.cpp | Tests backward dimension and stride limits. |
| dnn-providers/hip-kernel-provider/src/engines/hip_flash2_engine/HipFlash2FwdPlanBuilder.cpp | Guards dimensions before signed narrowing. |
| dnn-providers/hip-kernel-provider/src/engines/asm_sdpa_engine/plans/SdpaPlanUtils.hpp | Adds dimension predicates and safe stride comparisons. |
| dnn-providers/hip-kernel-provider/src/engines/asm_sdpa_engine/plans/SdpaFwdPlanBuilder.cpp | Guards dimensions and Ts; classifies original graph dimensions. |
| dnn-providers/hip-kernel-provider/src/engines/asm_sdpa_engine/plans/SdpaBwdPlanBuilder.cpp | Guards dimensions and safely checks tile-scaled strides. |
Review effort: Balanced
1 task done
The forward Ts bound and the q/k/v byte-stride checks assumed two bytes per input element. The kernarg builder uses one for fp8bf16, so an fp8 descriptor the kernel can run was declined. fwdInputBytesPerElement derives the size from the resolved dataTypeId and both checks read it; buildPlan sets params.inBytesPerElement from the same function so the two sites cannot drift. O stays two bytes. IsApplicableBoundsTheTsProductByTheResolvedTileSize gained fp8 boundary cases next to the bf16 ones. Flash2 isApplicable accepted seq_len_kv up to INT32_MAX, but both kernels count KV tiles in signed int ((seq_len_kv + 63) / 64 in HipFlash2FwdPlan.hip and (seq_len_kv + F2_BK - 1) / F2_BK with F2_BK at most 64 in HipFlash2FwdPlanVariant.hip), so the count overflowed at that boundary. The ceiling is now INT32_MAX - 64, and the acceptance case moved there with a new one-past case beside it. TestSdpaFwdPlanBuilder.cpp formed 2 * 192 in int before widening to int64_t, which bugprone-implicit-widening-of-multiplication-result rejects and the Linux superbuild treats as an error.
Keeps the attention-sink declines from #14246 and the 32-bit kernarg range checks from this PR. The only textual conflict was in TestHipFlash2FwdPlanBuilder.cpp, where both sides appended a test at the same spot.
brockhargreaves-amd
approved these changes
Oct 27, 2026
brockhargreaves-amd
left a comment
brockhargreaves-amd
left a comment
Contributor
There was a problem hiding this comment.
🤖 AI-generated review (Claude Code), posted by @brockhargreaves-amd. Checked against head 5892e0620b7 and disassembly of the shipped .co kernels.
The dimension gate, the divide-first byteStrideFitsU32, the new forward Ts check, and reading the dims in both isApplicable and buildPlan look correct, and the D_qk = 2^32 + 128 truncation fix closes a real bug. Recommendation: approve. No changes requested.
For the record: some "accepted at the limit" test cases (the dimension at-limit cases in both ASM builders, and the gfx942 fp8 and gfx950 bf16 Ts cases) assert acceptance of layouts the kernels address with wrapped 32-bit offsets. This predates the PR (develop accepts the same graphs) and is tracked in ALMIOPEN-3044, whose fix will relabel those cases and the related comments after this PR lands. Suggested merge order: this PR first; the other ASM_SDPA_ENGINE fixes under ALMIOPEN-3035 will rebase onto it.
jerehartAMD
approved these changes
Oct 27, 2026
Add IsApplicableBoundsEachBf16VByteStrideField and IsApplicableBoundsEachFp8VByteStrideField to TestSdpaFwdPlanBuilder. Each one moves a single V stride (batch_stride_v, nhead_stride_v, stride_v) to its uint32 byte bound while Q, K and O stay contiguous. bf16 accepts an element stride of UINT32_MAX / 2 and declines one more. fp8 (gfx942 only) scales by one byte, so it accepts UINT32_MAX and declines UINT32_MAX + 1, and it must still accept the UINT32_MAX / 2 + 1 stride that bf16 declines. createSdpaFwdFp8Graph gains a vStrides parameter for the fp8 cases.
Silences clang-tidy modernize-use-auto in HipDNN Linux Superbuild.
BradPepersAMD
added a commit
that referenced
this pull request
Oct 29, 2026
…e options API DeprecatedCausalMaskWithLeftBoundIsSlidingWindow called the positional GpuFpReferenceSdpa::fprop that this PR replaces with GpuSdpaFwdOptions. Its comment also predates #14483: left_bound=2 now keeps 2 keys per row, as cuDNN does. Regenerate SUPPORT_MATRIX_OVERVIEW.md after the rebase. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
1 task
brockhargreaves-amd
added a commit
that referenced
this pull request
Oct 29, 2026
…youts SdpaBwdPlanBuilder::isApplicable accepted tensor layouts that the gfx942 backward kernels address wrongly. On gfx942 (MI300X) each one gave silent wrong gradients instead of a decline. The evidence comes from two sources. One is GPU runs of the plugin from a CI build of #14312, which accepts all of these layouts. The other is disassembly of the dispatched .co files. Q/K/V/O/dO/dQ/dK/dV strides. The ODO, DQDKDV and DQ_CONVERT argument blocks carry only batch, head and sequence byte strides. isApplicable now declines: - A head-dim (innermost) stride other than 1. The kernels read it as 1: a stride-2 Q gave Inf dK/dV, and a stride-2 dQ left elements unwritten. - A sequence stride below D, also on a length-1 sequence axis. Each buffer ends S * seq stride bytes into its (batch, head) slice, so the end of the last row reads as 0, and writes there are dropped. A seq stride of 1 on a length-1 axis gave wrong gradients for each of the eight tensors. Seq strides of D/2 and 2 did the same for Q, K, V and dO. - An odd stride on an axis whose offset DQDKDV puts in its 4-byte record index (byte offset >> 2). That is the sequence axis of Q, K, V, dO, dK and dV, and the head axis of dK and dV. Rows that start 2 bytes into a dword are read or written one element early. A seq stride of D+1 failed on each of the six. A dO row pitch of D+1 gave 47-97 % wrong dQ, dK and dV. Odd strides on the other axes are addressed correctly and stay accepted: - batch strides, which are added to the 64-bit buffer base; - head strides of Q, K, V, O, dO and dQ; - seq strides of O and dQ, which ODO and DQ_CONVERT address by byte offset. The causal kernels also add the dK/dV head offset to the base, but the check does not depend on the mask. Stats (LSE). All 28 dispatched batch-mode DQDKDV kernels compute the LSE and D offset as ((b * H_q + h) * S_q + s) * 4 and never load Hs_lsed. A stats head stride padded to S_q + 64 gave wrong gradients that changed between runs with identical inputs. A padded batch stride, a seq stride of 2 and [B, S_q, H] storage also gave wrong gradients. isApplicable now requires stats dims [B, H_q, S_q] or [B, H_q, S_q, 1] matching Q, with packed strides on every axis longer than 1. buildPlan sets the LSE/D strides from Q's dims, so the stride of a size-1 stats axis no longer reaches a kernel argument or the 32-bit stride check. These layouts stay accepted, all GPU-verified on gfx942: BSHD on all tensors, dO strides that differ from O, seq strides D+2 and D+8, B and H above 1, hd192, fp16, causal and sliding-window masks, BR causal with Sq above Skv, Sq in {1, 16, 17}, and tails. New tests in TestSdpaBwdPlanBuilder.cpp (gfx942 only, like the other isApplicable tests). Each decline arm has a case that only it declines, and each exemption has an accept case: - IsApplicableAcceptsAddressableLayouts - IsApplicableRejectsNonUnitHeadDimStride - IsApplicableRejectsMisalignedOrShortRows - IsApplicableRejectsNonPackedStats Checked on gfx942 (MI300X). This commit's plugin replaced the one in a CI build of #14312, and results were compared with the CPU reference: - Each declined layout above failed with the old plugin, which accepts it, and now declines. Each layout named above as addressed correctly passed with the old plugin and still passes. The probes cover hd128 and hd192, bf16 and fp16, and the no-mask, causal and sliding-window kernels. - The shipped bf16 tolerance hides some short-row errors, so those were also compared at atol=1e-9 with packed controls. Their error was 7-34x the controls'; for Q, the dQ error sits in the cut-off last row. - The develop quick/standard SdpaBackward bundles and the CI Sdpa bundles give the same results as before. - The unit tests pass. Each of 29 mutants that removes or widens one arm of these checks makes one of the new tests fail. Rebased onto develop 98b504d, after #14075, whose range gate now runs first in isApplicable. The GPU results above are from develop 61c2675. On the new base, the four ALMIOPEN-3035 ASM SDPA fix branches (backward layouts, attention scale, a16 accumulator, forward layouts) were built and run together: - hip_kernel_provider_tests: 912 pass and 7 skip on gfx942 (MI325X), 886 pass and 33 skip on gfx950 (MI355X); none fail. - hip_kernel_provider_integration_tests: none fail on either. - develop's SDPA bundles on ASM_SDPA_ENGINE, run like the CI lane (claims enforced): 62 pass and 56 skip on gfx942, 15 pass and 103 skip on gfx950, and no support claim fails. A host-side run of the real isApplicable over these bundles finds no outcome change from develop. - #14622's SdpaBackward bundles (138218c) with its temporary [[test_skips]] removed: 275 pass and 225 skip on gfx942, none fail, and no support claim fails. This branch's checks decline the 20 cases behind the DEF-01, DEF-02, DEF-04, DEF-05/H61, prbu01, prbu03 and prbu13 skips, and no other case (host-checked for this branch alone). - clang-tidy 20 and clang-format 18 are clean on the changed files. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
brockhargreaves-amd
added a commit
that referenced
this pull request
Oct 29, 2026
SdpaBwdPlanBuilder::isApplicable accepted dQ, dK and dV with a zero batch or head stride on an axis longer than 1. Such a graph has no well-defined result in general, because the slices that share memory hold different gradients. No layer rejects it: the frontend and backend check only that dims are positive and that strides are set, and isApplicable's stride checks pass a zero stride on these axes, which fits in 32 bits and is even. Each kernel launches its own workgroups for each (batch, head) slice (grid y and z are head and batch). DQDKDV writes dK and dV, and DQ_CONVERT (A32 path) writes dQ, with plain stores at batch * batch stride + head * head stride. With a zero stride several slices share memory, so each element ends up holding whichever slice was stored last. On gfx942 (MI300X), B=2, H=2 (H=3 where noted), S=256, bf16, three runs in each of two GPU jobs: - dK and dV with batch stride 0: each memory location held one batch's gradient, so half of the dK and dV elements differed from the packed run, and 17k-54k of them changed between runs with identical inputs. - dQ with batch stride 0: the same, with 33k-65k changing between runs. - dK and dV with head stride 0, H=3: two thirds differed from the packed run, and 23k-147k changed between runs. - dQ with head stride 0, H=3: two thirds differed from the packed run, and 64k-148k changed between runs. There was no NaN or Inf and no write into the guard regions around each buffer, and dQ was unaffected when only dK and dV aliased. isApplicable now declines dQ, dK and dV with a zero batch or head stride on an axis longer than 1. The sequence and head-dim axes are already covered: their strides must be at least D and exactly 1. These stay accepted: - Zero strides on the inputs Q, K, V, O and dO (broadcasts), which the kernels just re-read. On gfx942 these gave gradients that matched the CPU reference and were bitwise equal to the packed runs, in each of three runs: K and V together, or dO alone, with a zero batch stride, head stride or both; Q, O or V alone with a zero batch or head stride; Q and O together with a zero head stride; and all five inputs with a zero batch stride. The K and V batch broadcast also passed with fp16, with a top-left causal mask, and at S=64. - A zero stride on a length-1 axis, on outputs too, since it only ever multiplies index 0. On gfx942, dQ, dK and dV with batch stride 0 at B=1, or head stride 0 at H=1, passed the same way. Nonzero strides that make an output overlap itself are not checked. On gfx942, dK and dV with equal batch and head strides (B=H=2) are still accepted. They differed from the packed run in a quarter of their elements, 6% of dK and 17% of dV were outside the CPU reference's bf16 tolerance, and the result was not always the same from run to run. The zero stride is checked on its own because it is how a broadcast is written, so an output gets one when it is given the strides of a broadcast input; other overlap has to be built by hand. PyTorch's check for tensors that an operation writes to also rejects only a zero stride on an axis longer than 1. A general rule (sort the axes by stride and require each stride to cover the span of the axes inside it) is left for a follow-up, in this plugin or in hipDNN graph validation, where every engine would get it. It would also need the range tests that #14075 added to switch to non-overlapping strides: they give dQ, dK and dV overlapping strides, such as equal batch and head strides with H=8, and expect them to be accepted. New test IsApplicableRejectsZeroStrideOutputs has one case per output tensor and axis, and covers the four layouts that failed on gfx942. IsApplicableAcceptsAddressableLayouts gains the broadcast inputs and the zero strides on length-1 output axes. Checked on gfx942 (MI300X) with a standalone hipDNN probe that drives ASM_SDPA_ENGINE on its own buffers, because the bundle harness's test tensors reject zero strides. It compares dQ, dK and dV with a float64 CPU reference and with the same data stored packed, three runs per case. - The develop plugin accepts every layout above. With this commit's plugin the four failing layouts decline, every broadcast and length-1 case above still passes bitwise, and the forward cases give the same results as with the develop plugin. - The unit and integration tests pass. Each of 14 mutants that removes or widens one arm of the check (dropping a length-1 exemption, or treating an input as an output) makes one of the two tests fail. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
brockhargreaves-amd
added a commit
that referenced
this pull request
Oct 29, 2026
IsApplicableBoundsTheTsProductByTheResolvedTileSize (added by #14075) probes K sequence strides at and one past the Ts limit. With the gfx942 bf16 hd128 tile (ts = 192) the past-limit stride is odd, and isApplicable now also declines an odd stride on K's sequence axis and dK's head axis, because DQDKDV indexes them in 4-byte records. The test still passed with the Ts guard deleted, so it no longer tested the guard. The at-limit stride would also fail the parity check for a tile size that makes the limit odd. Both probes are now the nearest even strides on each side of the limit, so the Ts guard is the only check that separates them. Deleting the guard makes the test fail again. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
brockhargreaves-amd
added a commit
that referenced
this pull request
Oct 29, 2026
…ales SdpaFwdPlanBuilder::isApplicable and SdpaBwdPlanBuilder::isApplicable accepted any attention scale, and execute passed it to the kernels unchecked. cuDNN and hipDNN accept any scale value. On gfx942 (MI300X) a scale of 0, -0.0, a negative value or a subnormal gave NaN or Inf outputs and gradients instead of a decline. The mechanism comes from the .co disassembly. The forward and backward kernels write -inf into every masked score (causal and sliding-window masks, and the padding past S_kv in the last K tile), and only then multiply the scores by scale * log2(e). A scale of 0 turns those scores into NaN, and a negative scale turns them into +inf. The bf16 forward kernels also start the running row max at -FLT_MAX / (scale * log2(e)), which is positive for a negative scale or -0.0, so every exp2 overflows even with no mask. 32-bit denormals are flushed, so a subnormal scale acts as 0. An infinite scale gives NaN everywhere. The gfx942 runs used the plugin from a CI build of #14312, which accepts all of these scales, and compared results with the CPU reference: - Backward: 0, -0.0, 1e-40 and -1.2 gave NaN or Inf dQ, dK and dV with top-left or bottom-right causal masks and sliding windows, in bf16 and fp16, hd128 and hd192. With no mask and S_kv = 128 (the hd128 K tile is 192), all of dQ was NaN. - Forward: -1.2, -0.088 and -0.0 gave NaN O and Inf LSE even with no mask. 0 gave NaN O with a bottom-right causal mask or S_kv = 200. - A runtime pass-by-value scale tensor gave the same results as attn_scale_value. Both isApplicable functions now decline a scale that is not a finite, positive, normal float (FLT_MIN <= scale <= FLT_MAX; NaN fails). The check covers attn_scale_value, and the value of a compile-time constant or runtime-with-default pass-by-value scale tensor, narrowed to the f32 kernarg as execute narrows it. They also decline a scale tensor whose value the plugin SDK cannot read, on which buildPlan would throw. The value of a pure runtime pass-by-value scale tensor arrives only at execute, so both execute functions now throw BAD_PARAM for such a scale before the first launch, as the gfx1151 WMMA attention pack does for its runtime scale. The check is shared in SdpaPlanUtils.hpp. An unset scale (1.0) and every finite, positive, normal scale stay accepted; 1e-37 and 1/sqrt(128) were correct on gfx942. The rule also declines some cases that happened to compute correctly: no mask with S_kv a multiple of the K tile, for any of these scales in the backward and for +0.0 or a subnormal scale in the forward. A rule keyed on the mask and the tile would break once group mode makes the K tail per-sequence, so the decline does not depend on them. Scales near FLT_MAX (1e37, 1e38) still give NaN in the forward; that is a separate issue. New tests: - TestSdpaPlanUtils (host only): IsSupportedAttnScaleRequiresFinitePositiveNormalFloat and AttnScaleBeforeExecute. - TestSdpaFwdPlanBuilder and TestSdpaBwdPlanBuilder (gfx942, and gfx950 for the forward, like the other isApplicable tests): IsApplicableRejectsUnsupportedAttnScale declines 0, -0.0, -1.2, -1/sqrt(128), -FLT_MIN, three subnormals, NaN and +-inf, each from attn_scale_value, a constant scale tensor and a runtime scale tensor default, plus an FP8_E4M3 scale tensor. IsApplicableAcceptsSupportedAttnScale accepts 1e-37, 1/sqrt(128) and 1 from each source, the unset default and a pure runtime scale tensor. - TestSdpaFwdPlanExecute and TestSdpaBwdPlanExecute: execute returns BAD_PARAM for each unsupported runtime scale (host only, on a plan with no loaded kernels), and passes supported ones on to the launch. Checked on gfx942 (MI300X). This commit's plugin replaced the one in a CI build of #14312, and results were compared with the CPU reference on the 82 scale cases behind the results above: - The 61 cases whose unsupported scale comes from attn_scale_value or a constant scale tensor are now declined. With the old plugin 48 of them failed, 12 passed (the no-mask cases named above) and one was already declined (fp16 forward). - The 4 cases with an unsupported runtime scale tensor (0 and -1.2, forward and backward) now fail at execute with BAD_PARAM, and the message gives the scale. The old plugin gave NaN or Inf in 3 of them. - The 15 cases with a supported scale pass with both plugins: 1/sqrt(128) with each mask, dtype and scale source, and 1e-37. The forward with 1e37 or 1e38 fails with both, as noted above. - On the develop CI build, develop's quick and standard SdpaFwd, SdpaFwdRuntimeScale and SdpaBackward bundles give the same result for each case as with develop's plugin: 42 pass and 39 are declined, for the same reasons, none of them the scale. - The unit tests pass. Each of 14 mutants that removes, widens or narrows one arm of these checks makes one of the new tests fail. Rebased onto develop 98b504d, after #14075, whose range gate now runs first in isApplicable. The GPU results above are from develop 61c2675. On the new base, the four ALMIOPEN-3035 ASM SDPA fix branches (backward layouts, attention scale, a16 accumulator, forward layouts) were built and run together: - hip_kernel_provider_tests: 912 pass and 7 skip on gfx942 (MI325X), 886 pass and 33 skip on gfx950 (MI355X); none fail. - hip_kernel_provider_integration_tests: none fail on either. - develop's SDPA bundles on ASM_SDPA_ENGINE, run like the CI lane (claims enforced): 62 pass and 56 skip on gfx942, 15 pass and 103 skip on gfx950, and no support claim fails. A host-side run of the real isApplicable over these bundles finds no outcome change from develop. - #14622's SdpaBackward bundles (138218c) with its temporary [[test_skips]] removed: 275 pass and 225 skip on gfx942, none fail, and no support claim fails. This change declines def06a, def06b and def06c (scale 0 and -1.2), the cases behind the DEF-06 skip, and no other case (host-checked for this branch alone). The execute-path scale tests (PassesSupportedRuntimeAttnScaleToLaunch) pass on both GPUs. - clang-tidy 20 and clang-format 18 are clean on the changed files. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
brockhargreaves-amd
added a commit
that referenced
this pull request
Oct 29, 2026
SdpaBwdPlanBuilder advertised the sdpa.bwd.accumulator_type knob with valid values a32 and a16, and initializeExecutionSettings accepted a16. isApplicable checks only the A32 kernels. With an A16 setting, buildPlan asks for a DQDKDV row with atomic32=0, pssk=1, pddv=1, and the gfx942 registry has none: its 24 a16 rows all have pssk=0, hd64 or hd128, and mask none or top-left causal. So buildPlan threw INTERNAL_ERROR "failed to resolve dqdkdv kernel ... (isApplicable should have rejected)" for every graph isApplicable accepts. On gfx942 (MI300X), develop's builder built all 42 accepted graphs of a sweep with a32 and failed all 42 with a16: bf16 and fp16; hd128 with no mask, top-left and bottom-right causal and sliding window; hd192 with no mask and both causal masks. hipDNN users cannot reach this today. AsmSdpaEngine publishes no builder knobs and never calls initializeExecutionSettings, so it always builds A32 plans, and the frontend refuses or drops knobs that an engine does not advertise. This change leaves the engine alone: with only a32 valid there is nothing to forward, and forwarding the knobs before this fix would have exposed the failure. Selecting the a16 rows that do exist would not fix it either. The gfx942 a16 kernels add dQ atomically, in the 16-bit output type, into the output dQ, address dQ and dO with Q's strides, and never load seqlen_k, the window bounds, or any V, dO, dK or dV stride other than dK's seq stride. In a bf16 diagnostic run with those rows patched in, cases inside that envelope matched a32 to 1 bf16 ulp. A dQ holding stale values, a dQ or dO layout that differs from Q's, Sq != Skv, and S not a multiple of 64 each gave wrong gradients, and a padded Q/dO row pitch and S=300 also wrote past the end of dQ. Now: - getCustomKnobs lists only a32. - initializeExecutionSettings rejects a16 with INVALID_VALUE and a message of its own. Other unknown strings still get INVALID_VALUE and non-string values BAD_PARAM. - buildPlan refuses an A16 setting with an INTERNAL_ERROR that says so, before it queries the device or resolves a kernel. It is the only gate: getMaxWorkspaceSize still sizes an A16 setting without dq_acc, but no A16 plan is ever built to use that size, and the A16 code in buildPlan and SdpaBwdPlan is kept for a future enablement. - Comments in the builder, the plan, and the params, kernel-args and workspace-size headers that said A16 can be selected, or that its kernels write dQ directly in BF16 and need no dQ zeroing, now say that the a16 kernels add dQ atomically into the output dQ, which must be zeroed first, and that A16 is refused. The comment on the rejection lists the graphs that getCustomKnobs and initializeExecutionSettings would have to limit A16 to before it could be enabled. isApplicable cannot do that: it sees no knob, and it also validates A32, which runs graphs A16 cannot. An unset knob and a32 behave as before, and isApplicable and the A32 plans are unchanged. Tests in TestSdpaBwdPlanBuilder.cpp: - GetCustomKnobsReturnsAccumulatorKnob now expects only a32 (gfx942 only, like the other isApplicable tests). - InitializeExecutionSettingsAcceptsA32 covers a32, and an unset knob from an invalid config, a valid config with no knobs, and a valid config with only another knob. - InitializeExecutionSettingsRejectsA16 - InitializeExecutionSettingsRejectsUnknownAccumulator - BuildPlanRejectsA16Setting, which needs no GPU. - Gfx942HasNoA16DqdkdvForDispatchedLayout pins the registry fact that the rejection comment cites, with the pssk=0 a16 rows as positive controls. Checked on gfx942 (MI300X). All 228 asm_sdpa unit tests pass except IsApplicableRejectsGfx950, which skips on gfx942 as it does on develop (223 tests there). The host-only tests also pass on a machine without a gfx942 device. Each of 9 mutants makes one of the tests above fail: mapping a16 back to A16, dropping the a16 arm, letting a32 fall through, listing a16 again, dropping the buildPlan check (buildPlan then fails to resolve a DQDKDV kernel, as on develop), accepting unknown strings, dropping the has-knob check (a valid config without the knob then throws std::out_of_range), pointing the registry check at gfx950, which has such rows, and making every lookup in it miss. develop's quick and standard SdpaBackward bundles (8 cases) pass against the CPU reference with this plugin, as with develop's, and still build A32 plans. clang-format 18 and clang-tidy 20 are clean on the changed files. Rebased onto develop 98b504d, after #14075, whose range gate now runs first in isApplicable. The GPU results above are from develop 61c2675. On the new base, the four ALMIOPEN-3035 ASM SDPA fix branches (backward layouts, attention scale, a16 accumulator, forward layouts) were built and run together: - hip_kernel_provider_tests: 912 pass and 7 skip on gfx942 (MI325X), 886 pass and 33 skip on gfx950 (MI355X); none fail. - hip_kernel_provider_integration_tests: none fail on either. - develop's SDPA bundles on ASM_SDPA_ENGINE, run like the CI lane (claims enforced): 62 pass and 56 skip on gfx942, 15 pass and 103 skip on gfx950, and no support claim fails. A host-side run of the real isApplicable over these bundles finds no outcome change from develop. - No #14622 SdpaBackward case changes outcome: no bundle sets a knob (host-checked for this branch alone). - clang-tidy 20 and clang-format 18 are clean on the changed files. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
brockhargreaves-amd
added a commit
that referenced
this pull request
Oct 29, 2026
…outs SdpaFwdPlanBuilder::isApplicable accepted Q/K/V/O and stats (LSE) layouts that the forward kernels address wrongly. Its only stride check was that byte strides fit in 32 bits. On gfx942 (MI300X) each kind of layout below gave a silently wrong O or wrong stats instead of a decline, in the cases that were run (the exceptions are noted). The evidence comes from two sources. One is GPU runs of the develop CI plugin, which accepts all of these layouts, compared with the CPU reference. The other is disassembly of the 14 forward .co files. Q/K/V/O strides. The forward argument block carries only batch, head and sequence byte strides, and the kernels read or write each row as D contiguous elements. isApplicable now declines: - A head-dim (innermost) stride other than 1. The kernels read it as 1. Stride-2 Q, K or V gave a wrong O that changed between runs, because the kernels read the gaps between elements. A stride-2 O left half of O unwritten. Q stored as [B, H, D, S] and hd192 Q with stride 2 also failed. - A sequence stride below the tensor's own head dim, also 0 and also on a length-1 sequence axis. Each buffer ends S * seq stride bytes into its (batch, head) slice, so the end of the last row reads as 0 and writes there are dropped. A seq stride of D/2 gave a wrong O for Q and V, and for K at S_kv = 2. At S_kv = 256 the K case stayed within the bf16 tolerance, but its error was about 6x the packed run's. With S_q = 1 and an O seq stride of 1, O was left NaN. With S_kv = 1 and a V seq stride of 1, O was 0. K and V with seq stride 0 gave O and stats off by up to 1.0 and 1.3, Q with seq stride 0 gave a wrong O and stats, and O with seq stride 0 was left unwritten. With S_kv = 1 a K seq stride of 1 still gave the right O, since the one key takes all the softmax weight; by the address formula its score, and so the stats, are wrong. Each tensor is compared with its own head dim. That relies on K having Q's head dim and O being [B, H_q, S_q, D_v], because the kernels size K rows by Q's head dim and O's buffer by S_q. The forward isApplicable does not check those dims yet; a forward counterpart of the backward's mismatched-dims decline is a separate change. - A zero O batch or head stride on an axis longer than 1. Each (batch, head) slice of O is written by its own workgroups, so such slices overwrite each other: with an O batch or head stride of 0, half or two thirds of O differed from the packed run, and O changed between some runs with identical inputs. Nonzero strides that make O overlap itself are not checked. On gfx942, O with equal batch and head strides (B = H = 2), or with a head stride of half a slice, differed from the packed run in a quarter of its elements, with both plugins. As in the backward, only a zero stride is checked: a zero stride is how a broadcast is written, so O gets one when it is given a broadcast input's strides, while overlapping nonzero strides have to be built by hand. PyTorch's check for tensors that an operation writes to also rejects only a zero stride on an axis longer than 1. Stats do get an overlap check (below), because their address formula leaves a single comparison to make; for O it would take a footprint check over three axes in any order, which would fit better in hipDNN's graph validation, where every engine gets it. The zero O batch stride check assumes batch mode, the only mode with forward kernels. In group (ragged) mode the kernels locate each batch by its sequence offsets and get no O batch stride, so when group-mode forward kernels are added, a ragged O must be exempt from the batch-axis part of this check. Stats. Every forward kernel writes the LSE of (b, h, s) at ((b * H_q + h) * lse_Hs + s) * 4 bytes, with lse_Hs the stats head stride, and takes no stats batch or sequence stride. isApplicable now declines stats whose seq stride is not 1 when S_q > 1, or whose batch stride is not H_q times the head stride when B > 1. It also declines a head stride below S_q when B * H_q > 1, with which (batch, head) slices overlap. On gfx942 each of these gave wrong or NaN stats while O was correct: - a batch stride padded to H_q * S_q + 128, with hd128, hd192 and GQA; - a seq stride of 2; - [B, S_q, H] storage, and H_q = 1 with strides [S_q, 1, 1, 1]; - [H, B, S_q] storage, which by the kernels' address formula also writes 2 KiB past the end of the stats buffer. This is looser than the backward's packed-only stats rule on purpose. The backward kernels ignore the stats strides, but the forward kernels use lse_Hs: a head stride padded to S_q + 64, with a batch stride of H_q times that, gave correct stats on gfx942 and stays accepted. The stride of a size-1 stats batch or sequence axis is not checked. A size-1 head axis is, on purpose: buildPlan passes the head stride as lse_Hs, so with H_q = 1 the kernel steps between batches by the head stride. Passing the batch stride instead when H_q = 1 would also accept packed [B, 1, S_q] stats with any head stride, such as the H_q = 1 case above; they are rare and decline. The stats decline reasons name the stats dims and strides and Q's dims. These stay accepted, all GPU-verified with bf16 on gfx942: - BSHD on all tensors, and O BSHD with BHSD inputs; - seq strides D+2 and D+8; - odd strides: a seq stride of D+1 and odd head and batch strides on each of Q, K, V and O, with hd128 and hd192, gave O bitwise equal to the packed runs (the backward declines some of these); - hd192, GQA, bottom-right causal, tails, and S_q = 1 or S_kv = 1 packed; - a zero batch or head stride on Q, K or V (a broadcast): K and V with a zero batch or head stride, and Q with a zero batch stride, gave O bitwise equal to the packed runs and stats that matched the CPU reference. A zero stride on a length-1 O axis is accepted too, since it only ever multiplies index 0. The gfx950 kernels have the same argument block and the same LSE address formula (from their disassembly), but load rows with 16-byte LDS-DMA where the gfx942 kernels use 4-byte. On gfx950 (MI355X), this commit's plugin gave the BSHD, D+2, D+8 and odd-stride layouts above, with hd128 and hd192, the same mismatch counts and max and mean error against the CPU reference at atol=1e-9 as the packed layouts. The fp8 kernels also have the same argument block and LSE formula, so the same rules apply to them. They were not run; in particular odd strides, which put fp8 rows 1 or 3 bytes past a 4-byte boundary, are unverified with fp8. New tests in TestSdpaFwdPlanBuilder.cpp (gfx942 and gfx950, like the other forward isApplicable tests). Each decline arm has a case that only it declines, and each exemption has an accept case: - IsApplicableAcceptsAddressableLayouts - IsApplicableRejectsNonUnitHeadDimStride - IsApplicableRejectsShortRows, including hd192 Q and K rows long enough for V's and O's head dim but not their own - IsApplicableRejectsZeroStrideOutput - IsApplicableRejectsUnaddressableStats The zero-stride and overlapping layouts were run with a standalone hipDNN probe that drives ASM_SDPA_ENGINE on its own buffers, because the bundle harness's test tensors reject zero strides. Checked on gfx942 (MI300X). This commit's plugin replaced the one in a develop CI build whose ASM SDPA engine source matches this commit's parent, and results were compared with the CPU reference: - 84 forward layout probes. The 25 that this commit declines now skip, each with a reason that names the tensor, its dims and its strides. With the develop plugin, 21 of them gave a wrong O or wrong stats. Two could not be checked there: the harness rejects a zero stride (the standalone probe covers it), and the CPU reference rejects rank-3 stats. The last two passed but are wrong in ways the check missed: the K seq stride D/2 case stays within the bf16 tolerance, and with S_kv = 1 a K seq stride of 1 only gives the right O because one key takes all the softmax weight. - The other 59 give the same results as with the develop plugin, down to the mismatch counts and the max and mean error at atol=1e-9. 47 pass; they cover every layout named above as staying accepted except the broadcasts, which the next item covers. 3 were already declined. 9 still fail: non-positive or denormal scales, which are left to a separate change, and rank-3 stats, which the CPU reference cannot check. - On the standalone probe, K/V and Q broadcasts still match the CPU reference and are bitwise equal to their packed twins in each of 3 runs. An O batch or head stride of 0 and a Q, K/V or O seq stride of 0 now decline; with the develop plugin each gave a wrong O. The two overlapping O layouts are accepted by both plugins, with the same mismatch counts. - 19 probes with tensors of up to 4 GiB, which exercise the kernels' 32-bit address math, give the same results as before (10 of them still fail with both plugins; the 32-bit address math is a separate change). - The develop SdpaFwd and SdpaFwdRuntimeScale bundles (102 tests, quick to full) give the same results as before: 46 pass and 56 skip, with the same decline reasons. - The provider's own forward GPU integration tests (hip_kernel_provider_integration_tests, *IntegrationGpuSdpaFwd*: Smoke, SmokeFp8, SmokeStats, ShapeSweep, Gqa and AsymmetricSeq) pass against this commit's plugin: 29 pass, and 21 skip as gfx950-only or causal with stats. - The unit tests pass. Each of 33 mutants that removes, narrows or widens one arm of these checks, or compares a row stride with another tensor's head dim or a fixed 128, makes one of the new tests fail. On gfx950 the 84 forward layout probes ran with this commit's plugin only, since the develop CI build has no gfx950 kernels. The same 28 skip, and every probe that passes on gfx942 and has a gfx950 kernel passes. The forward unit tests pass there too. Rebased onto develop 98b504d, after #14075, whose range gate now runs first in isApplicable. The GPU results above are from develop 61c2675. On the new base, the four ALMIOPEN-3035 ASM SDPA fix branches (backward layouts, attention scale, a16 accumulator, forward layouts) were built and run together: - hip_kernel_provider_tests: 912 pass and 7 skip on gfx942 (MI325X), 886 pass and 33 skip on gfx950 (MI355X); none fail. - hip_kernel_provider_integration_tests: none fail on either. - develop's SDPA bundles on ASM_SDPA_ENGINE, run like the CI lane (claims enforced): 62 pass and 56 skip on gfx942, 15 pass and 103 skip on gfx950, and no support claim fails. A host-side run of the real isApplicable over these bundles finds no outcome change from develop. - No develop or #14622 bundle case changes outcome (host-checked for this branch alone). The new forward tests pass on gfx942 and gfx950. - clang-tidy 20 and clang-format 18 are clean on the changed files. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
This was referenced Oct 29, 2026
adickin-amd
added a commit
that referenced
this pull request
Oct 30, 2026
Brings in the SDPA backward O/LSE derivation (#14312), the cuDNN left-bound semantics (#14483), the SDPA kernarg range checks (#14075) and --benchmark (#13915). Conflicts: the outcome publishing moved into runBody() on this branch, so the benchmark report now sits there; TestFillInputs and the docs keep both sides.
Contributor
TheRock Submodule Bump ActivityNewest first
|
brockhargreaves-amd
added a commit
that referenced
this pull request
Oct 30, 2026
…ales SdpaFwdPlanBuilder::isApplicable and SdpaBwdPlanBuilder::isApplicable accepted any attention scale, and execute passed it to the kernels unchecked. cuDNN and hipDNN accept any scale value. On gfx942 (MI300X) a scale of 0, -0.0, a negative value or a subnormal gave NaN or Inf outputs and gradients instead of a decline. The mechanism comes from the .co disassembly. The forward and backward kernels write -inf into every masked score (causal and sliding-window masks, and the padding past S_kv in the last K tile), and only then multiply the scores by scale * log2(e). A scale of 0 turns those scores into NaN, and a negative scale turns them into +inf. The bf16 forward kernels also start the running row max at -FLT_MAX / (scale * log2(e)), which is positive for a negative scale or -0.0, so every exp2 overflows even with no mask. 32-bit denormals are flushed, so a subnormal scale acts as 0. An infinite scale gives NaN everywhere. The gfx942 runs used the plugin from a CI build of #14312, which accepts all of these scales, and compared results with the CPU reference: - Backward: 0, -0.0, 1e-40 and -1.2 gave NaN or Inf dQ, dK and dV with top-left or bottom-right causal masks and sliding windows, in bf16 and fp16, hd128 and hd192. With no mask and S_kv = 128 (the hd128 K tile is 192), all of dQ was NaN. - Forward: -1.2, -0.088 and -0.0 gave NaN O and Inf LSE even with no mask. 0 gave NaN O with a bottom-right causal mask or S_kv = 200. - A runtime pass-by-value scale tensor gave the same results as attn_scale_value. Both isApplicable functions now decline a scale that is not a finite, positive, normal float (FLT_MIN <= scale <= FLT_MAX; NaN fails). The check covers attn_scale_value, and the value of a compile-time constant or runtime-with-default pass-by-value scale tensor, narrowed to the f32 kernarg as execute narrows it. They also decline a scale tensor whose value the plugin SDK cannot read, on which buildPlan would throw. The value of a pure runtime pass-by-value scale tensor arrives only at execute, so both execute functions now throw BAD_PARAM for such a scale before the first launch, as the gfx1151 WMMA attention pack does for its runtime scale. The check is shared in SdpaPlanUtils.hpp. An unset scale (1.0) and every finite, positive, normal scale stay accepted; 1e-37 and 1/sqrt(128) were correct on gfx942. The rule also declines some cases that happened to compute correctly: no mask with S_kv a multiple of the K tile, for any of these scales in the backward and for +0.0 or a subnormal scale in the forward. A rule keyed on the mask and the tile would break once group mode makes the K tail per-sequence, so the decline does not depend on them. Scales near FLT_MAX (1e37, 1e38) still give NaN in the forward; that is a separate issue. New tests: - TestSdpaPlanUtils (host only): IsSupportedAttnScaleRequiresFinitePositiveNormalFloat and AttnScaleBeforeExecute. - TestSdpaFwdPlanBuilder and TestSdpaBwdPlanBuilder (gfx942, and gfx950 for the forward, like the other isApplicable tests): IsApplicableRejectsUnsupportedAttnScale declines 0, -0.0, -1.2, -1/sqrt(128), -FLT_MIN, three subnormals, NaN and +-inf, each from attn_scale_value, a constant scale tensor and a runtime scale tensor default, plus an FP8_E4M3 scale tensor. IsApplicableAcceptsSupportedAttnScale accepts 1e-37, 1/sqrt(128) and 1 from each source, the unset default and a pure runtime scale tensor. - TestSdpaFwdPlanExecute and TestSdpaBwdPlanExecute: execute returns BAD_PARAM for each unsupported runtime scale (host only, on a plan with no loaded kernels), and passes supported ones on to the launch. Checked on gfx942 (MI300X). This commit's plugin replaced the one in a CI build of #14312, and results were compared with the CPU reference on the 82 scale cases behind the results above: - The 61 cases whose unsupported scale comes from attn_scale_value or a constant scale tensor are now declined. With the old plugin 48 of them failed, 12 passed (the no-mask cases named above) and one was already declined (fp16 forward). - The 4 cases with an unsupported runtime scale tensor (0 and -1.2, forward and backward) now fail at execute with BAD_PARAM, and the message gives the scale. The old plugin gave NaN or Inf in 3 of them. - The 15 cases with a supported scale pass with both plugins: 1/sqrt(128) with each mask, dtype and scale source, and 1e-37. The forward with 1e37 or 1e38 fails with both, as noted above. - On the develop CI build, develop's quick and standard SdpaFwd, SdpaFwdRuntimeScale and SdpaBackward bundles give the same result for each case as with develop's plugin: 42 pass and 39 are declined, for the same reasons, none of them the scale. - The unit tests pass. Each of 14 mutants that removes, widens or narrows one arm of these checks makes one of the new tests fail. Rebased onto develop 98b504d, after #14075, whose range gate now runs first in isApplicable. The GPU results above are from develop 61c2675. On the new base, the four ALMIOPEN-3035 ASM SDPA fix branches (backward layouts, attention scale, a16 accumulator, forward layouts) were built and run together: - hip_kernel_provider_tests: 912 pass and 7 skip on gfx942 (MI325X), 886 pass and 33 skip on gfx950 (MI355X); none fail. - hip_kernel_provider_integration_tests: none fail on either. - develop's SDPA bundles on ASM_SDPA_ENGINE, run like the CI lane (claims enforced): 62 pass and 56 skip on gfx942, 15 pass and 103 skip on gfx950, and no support claim fails. A host-side run of the real isApplicable over these bundles finds no outcome change from develop. - #14622's SdpaBackward bundles (138218c) with its temporary [[test_skips]] removed: 275 pass and 225 skip on gfx942, none fail, and no support claim fails. This change declines def06a, def06b and def06c (scale 0 and -1.2), the cases behind the DEF-06 skip, and no other case (host-checked for this branch alone). The execute-path scale tests (PassesSupportedRuntimeAttnScaleToLaunch) pass on both GPUs. - clang-tidy 20 and clang-format 18 are clean on the changed files. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 30, 2026
…e options API DeprecatedCausalMaskWithLeftBoundIsSlidingWindow called the positional GpuFpReferenceSdpa::fprop that this PR replaces with GpuSdpaFwdOptions. Its comment also predates #14483: left_bound=2 now keeps 2 keys per row, as cuDNN does. Regenerate SUPPORT_MATRIX_OVERVIEW.md after the rebase. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 31, 2026
…e options API DeprecatedCausalMaskWithLeftBoundIsSlidingWindow called the positional GpuFpReferenceSdpa::fprop that this PR replaces with GpuSdpaFwdOptions. Its comment also predates #14483: left_bound=2 now keeps 2 keys per row, as cuDNN does. Regenerate SUPPORT_MATRIX_OVERVIEW.md after the rebase. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
BradPepersAMD
added a commit
that referenced
this pull request
Oct 31, 2026
…e options API DeprecatedCausalMaskWithLeftBoundIsSlidingWindow called the positional GpuFpReferenceSdpa::fprop that this PR replaces with GpuSdpaFwdOptions. Its comment also predates #14483: left_bound=2 now keeps 2 keys per row, as cuDNN does. Regenerate SUPPORT_MATRIX_OVERVIEW.md after the rebase. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
brockhargreaves-amd
added a commit
that referenced
this pull request
Nov 3, 2026
…ales SdpaFwdPlanBuilder::isApplicable and SdpaBwdPlanBuilder::isApplicable accepted any attention scale, and execute passed it to the kernels unchecked. cuDNN and hipDNN accept any scale value. On gfx942 (MI300X) a scale of 0, -0.0, a negative value or a subnormal gave NaN or Inf outputs and gradients instead of a decline. The mechanism comes from the .co disassembly. The forward and backward kernels write -inf into every masked score (causal and sliding-window masks, and the padding past S_kv in the last K tile), and only then multiply the scores by scale * log2(e). A scale of 0 turns those scores into NaN, and a negative scale turns them into +inf. The bf16 forward kernels also start the running row max at -FLT_MAX / (scale * log2(e)), which is positive for a negative scale or -0.0, so every exp2 overflows even with no mask. 32-bit denormals are flushed, so a subnormal scale acts as 0. An infinite scale gives NaN everywhere. The gfx942 runs used the plugin from a CI build of #14312, which accepts all of these scales, and compared results with the CPU reference: - Backward: 0, -0.0, 1e-40 and -1.2 gave NaN or Inf dQ, dK and dV with top-left or bottom-right causal masks and sliding windows, in bf16 and fp16, hd128 and hd192. With no mask and S_kv = 128 (the hd128 K tile is 192), all of dQ was NaN. - Forward: -1.2, -0.088 and -0.0 gave NaN O and Inf LSE even with no mask. 0 gave NaN O with a bottom-right causal mask or S_kv = 200. - A runtime pass-by-value scale tensor gave the same results as attn_scale_value. Both isApplicable functions now decline a scale that is not a finite, positive, normal float (FLT_MIN <= scale <= FLT_MAX; NaN fails). The check covers attn_scale_value, and the value of a compile-time constant or runtime-with-default pass-by-value scale tensor, narrowed to the f32 kernarg as execute narrows it. They also decline a scale tensor whose value the plugin SDK cannot read, on which buildPlan would throw. The value of a pure runtime pass-by-value scale tensor arrives only at execute, so both execute functions now throw BAD_PARAM for such a scale before the first launch, as the gfx1151 WMMA attention pack does for its runtime scale. The check is shared in SdpaPlanUtils.hpp. An unset scale (1.0) and every finite, positive, normal scale stay accepted; 1e-37 and 1/sqrt(128) were correct on gfx942. The rule also declines some cases that happened to compute correctly: no mask with S_kv a multiple of the K tile, for any of these scales in the backward and for +0.0 or a subnormal scale in the forward. A rule keyed on the mask and the tile would break once group mode makes the K tail per-sequence, so the decline does not depend on them. Scales near FLT_MAX (1e37, 1e38) still give NaN in the forward; that is a separate issue. New tests: - TestSdpaPlanUtils (host only): IsSupportedAttnScaleRequiresFinitePositiveNormalFloat and AttnScaleBeforeExecute. - TestSdpaFwdPlanBuilder and TestSdpaBwdPlanBuilder (gfx942, and gfx950 for the forward, like the other isApplicable tests): IsApplicableRejectsUnsupportedAttnScale declines 0, -0.0, -1.2, -1/sqrt(128), -FLT_MIN, three subnormals, NaN and +-inf, each from attn_scale_value, a constant scale tensor and a runtime scale tensor default, plus an FP8_E4M3 scale tensor. IsApplicableAcceptsSupportedAttnScale accepts 1e-37, 1/sqrt(128) and 1 from each source, the unset default and a pure runtime scale tensor. - TestSdpaFwdPlanExecute and TestSdpaBwdPlanExecute: execute returns BAD_PARAM for each unsupported runtime scale (host only, on a plan with no loaded kernels), and passes supported ones on to the launch. Checked on gfx942 (MI300X). This commit's plugin replaced the one in a CI build of #14312, and results were compared with the CPU reference on the 82 scale cases behind the results above: - The 61 cases whose unsupported scale comes from attn_scale_value or a constant scale tensor are now declined. With the old plugin 48 of them failed, 12 passed (the no-mask cases named above) and one was already declined (fp16 forward). - The 4 cases with an unsupported runtime scale tensor (0 and -1.2, forward and backward) now fail at execute with BAD_PARAM, and the message gives the scale. The old plugin gave NaN or Inf in 3 of them. - The 15 cases with a supported scale pass with both plugins: 1/sqrt(128) with each mask, dtype and scale source, and 1e-37. The forward with 1e37 or 1e38 fails with both, as noted above. - On the develop CI build, develop's quick and standard SdpaFwd, SdpaFwdRuntimeScale and SdpaBackward bundles give the same result for each case as with develop's plugin: 42 pass and 39 are declined, for the same reasons, none of them the scale. - The unit tests pass. Each of 14 mutants that removes, widens or narrows one arm of these checks makes one of the new tests fail. Rebased onto develop 98b504d, after #14075, whose range gate now runs first in isApplicable. The GPU results above are from develop 61c2675. On the new base, the four ALMIOPEN-3035 ASM SDPA fix branches (backward layouts, attention scale, a16 accumulator, forward layouts) were built and run together: - hip_kernel_provider_tests: 912 pass and 7 skip on gfx942 (MI325X), 886 pass and 33 skip on gfx950 (MI355X); none fail. - hip_kernel_provider_integration_tests: none fail on either. - develop's SDPA bundles on ASM_SDPA_ENGINE, run like the CI lane (claims enforced): 62 pass and 56 skip on gfx942, 15 pass and 103 skip on gfx950, and no support claim fails. A host-side run of the real isApplicable over these bundles finds no outcome change from develop. - #14622's SdpaBackward bundles (138218c) with its temporary [[test_skips]] removed: 275 pass and 225 skip on gfx942, none fail, and no support claim fails. This change declines def06a, def06b and def06c (scale 0 and -1.2), the cases behind the DEF-06 skip, and no other case (host-checked for this branch alone). The execute-path scale tests (PassesSupportedRuntimeAttnScaleToLaunch) pass on both GPUs. - clang-tidy 20 and clang-format 18 are clean on the changed files. 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.
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 onint. The plan builders checked the byte strides and nothing else, so a batchsize, 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::dimFitsU32andplan_utils::dimFitsI32sit next to the existingbyteStrideFitsU32and follow its shape: log the offending field, returnfalse, let the caller report every violation of a tensor in one pass. Both ASM
plan builders call a per-pass wrapper from
isApplicableimmediately after therank checks, which is ahead of the first narrowing cast in each builder.
UINT32_MAXINT32_MAXcomputeMaskCoordinates, which is int32 throughout, and both passes compute(extent + tile - 1) / tileinunsigned int, which wraps nearUINT32_MAXINT32_MAXintfor the kernel-registry lookup before they reach the uint32 fieldINT32_MAXon the sequence lengths is also what makes the window narrowing inSdpaBwdPlanBuilderexact.makeCoveringBoundsUnboundreplaces any bound thatreaches its span with -1, so a surviving bound is strictly below a sequence
length, and every
1 + bound + offsetthatcomputeMaskCoordinatesforms thenstays inside int32.
The head dims are the case with teeth.
D_qk = 2^32 + 128truncated 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 * elementBytesbefore testing the product, so an elementcount 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 sizeinstead of dividing by it, and prints the largest element count that would fit.
Ts
Tsis the one kernarg that is a product of a stride and a tile size. Thebackward checked it but built
tsKv * stride_kin int64 before calling thepredicate, and the forward had no check at all even though
s_Ts = tileSizeQo * qStrideRow * inBytesPerElementis computed inunsigned intinSdpaFwdArgsBuilder. Both now fold the tile size into theelement size argument, so the product is never formed outside the predicate,
and the forward resolves
ts_qofrom the row it already looked up. Theforward element size is the resolved input size, one byte for fp8 and two for
bf16, read from the same
dataTypeIdbuildPlanuses; the per-field q/k/vstride check takes it too, and O stays two bytes.
isApplicable and buildPlan now classify on the same values
buildPlanre-read the sequence lengths asunsigned intand classified themask on the truncated values while
isApplicableclassified on the graph dims.A sequence length above
UINT32_MAXcould therefore be approved as one masktype 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
Flash2FwdParamsnarrows every dimension toint, andbatchhad no checkanywhere:
isApplicablenever read it, so a batch of2^32 + 1was dispatchedas 1. The same dimension gate now runs in that engine's
isApplicable, aheadof its casts, with
INT32_MAXfor all six dims because they are allint.seq_len_kvstops one KV tile lower, atINT32_MAX - 64: both kernels countKV tiles in signed
int, and 64 is the largest KV tile the engine dispatches.Its stride handling is untouched:
checkedStridealready rejects out-of-rangestrides, 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 keepsits 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
e0b0ad5f6dand branchfeef30d804on gfx942 andgfx950 with the same Release/Ninja configure. The gfx942 run used an MI325X
node in partition
miopen. The gfx950 run used an MI355X node in partitionmiopen; the rocke MI355X nodes had no free GPU at that time. Both jobs usedaccount
miopen.On gfx942:
hip_kernel_provider_tests: develop, 898 ran, 891 passed, 7 skipped;branch, 928 ran, 921 passed, 7 skipped. No failures.
ran, 102 passed and 1 skipped. The skip is the gfx950-only case.
hipdnn_integration_tests_unit_tests --gtest_filter=TestGpuSdpaFwdPlanBuilder.*:develop, 15 passed; branch, 16 passed.
On gfx950:
hip_kernel_provider_tests: develop, 898 ran, 869 passed, 29 skipped;branch, 928 ran, 884 passed, 44 skipped. No failures.
ran, 107 passed and 2 skipped. Both skips are the existing fp8 descale cases.
ran, 87 passed and 16 skipped. The four additional skips are the new
backward boundary test methods, each guarded to run only on gfx942.
25 passed and 20 skipped. All 11 new
TestHipFlash2DimRangecases skip atthe 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
TestSdpaKernargRangestests, and the ported sliding-window GPU-reference testran 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 exactlimit and one past it, the negative-dimension case, the tile-factor case, and
the 2^62 case that the old multiply-first form accepted.
TestSdpaFwdPlanBuilderandTestSdpaBwdPlanBuilder: each dimension at thelimit 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 fix(hipdnn): keep left_bound under causal_mask in SDPA engines #12923:
causal_mask=truewithleft_bound=2runs asleftBound=2, rightBound=0against the GPU reference, with a control that thewindow actually changes the output. It holds under fix(hipdnn): Set Sdpa diagonal alignment and left/right_bounds correctly #13065's contract because
makeCoveringBoundsUnboundleaves a bound of 2 alone at Sq = Skv = 8.Two backward cases cannot be isolated to the new guard:
nhead_qandseqlen_qpast 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 areisolated, and the predicate suite covers the boundary exactly.
Negative control, same node and configure line, with the three plan builders
restored to
f86ff449andbyteStrideFitsU32restored to its multiply-firstform (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
byteStrideFitsU32cases the old body accepted. The cases thatassert 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 guardalready rejected that stride before this change; reordering the multiply only
matters when
tsKv * stride_kitself overflows int64, which the per-fieldstride 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_MAXrather than the field'sUINT32_MAX. That is the stricter of the two defensible limits for the forwardpass 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_MAXshould say so; the alternative is a tile-sizedependent bound.
byteStrideFitsU32now declineselementBytes <= 0. No current caller passesone, 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.TestSdpaFwdPlanBuilder.cppformed2 * 192inintbefore widening it toint64_t, whichbugprone-implicit-widening-of-multiplication-resultrejects and the Linux superbuild treats as an error. It reads
int64_t{2} * 192now.seq_len_kvceiling, from review. The engine acceptedINT32_MAXwhile both kernels count KV tiles in signedint, so the countoverflowed at the accepted boundary.
K_FLASH2_MAX_SEQ_LEN_KVisINT32_MAX - 64and the builder and the test read the same constant. Theacceptance case moved there, with a one-past case next to it;
TestHipFlash2DimRangeis 12 cases.per-field stride check assumed two bytes, so an fp8 descriptor the kernel
can run was declined.
fwdInputBytesPerElementderives the size from theresolved
dataTypeIdandbuildPlanfillsinBytesPerElementfrom thesame function.
IsApplicableBoundsTheTsProductByTheResolvedTileSizegainedfp8 boundary cases beside the bf16 ones, with
ts_qoread from thefp8bf16registry row.clang-tidy 20.1.2 with
projects/hipdnn/.clang-tidyis clean over all seventranslation 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.
hip_kernel_provider_tests: gfx942, 929 ran, 922 passed, 7 skipped;gfx950, 929 ran, 884 passed, 45 skipped. No failures.
2 skipped (the existing fp8 descale cases).
87 passed and 16 skipped.
seq_len_kvcases; gfx950, 46 ran, 25 passed and 21 skipped at the existing gfx942-only
device guard.
hipdnn_integration_tests_unit_testsfiltered to*GpuSdpaFwdPlan*: 16passed on both.
hip_kernel_provider_integration_testsfiltered to*SdpaFwd*on gfx942: 50ran, 29 passed, 21 skipped (the gfx950 rows). The three
fp8bf16GPU casespass, 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_MAXcase became two.