fix(hipdnn): follow cuDNN semantics for the SDPA left diagonal-band bound

dnn-providers/hip-kernel-provider/src/engines/asm_sdpa_engine/plans/SdpaPlanUtils.hpp

@@ -57,6 +57,12 @@ enum class MaskType : int
 
 // The diagonal band an SDPA (forward or backward) attribute set requests, with
 // the deprecated causal_mask / causal_mask_bottom_right booleans merged in.
+//
+// The bounds are in the kernels' FlashAttention convention (window_size_left /
+// window_size_right), not the graph's: leftBound is the number of keys kept
+// before the diagonal, so it is the graph's cuDNN-semantics left_bound minus 1
+// (see resolveDiagonalBand()). rightBound means the same in both conventions.
+// -1 is unbounded.
 struct DiagonalBand
 {
     int64_t leftBound;
@@ -100,7 +106,12 @@ inline DiagonalBand makeCoveringBoundsUnbound(DiagonalBand band, int64_t seqLenQ
 // alignment; causal_mask_bottom_right always resolves to BOTTOM_RIGHT, because
 // TOP_LEFT is the schema default and indistinguishable from an unset alignment.
 //
-// Setting both booleans, or a bound below -1, throws
+// The graph's left_bound follows cuDNN: left_bound = L keeps the L keys ending at
+// the diagonal, diagonal included. The returned band holds the FlashAttention
+// window_size_left the kernels take, L - 1 (see DiagonalBand).
+//
+// Setting both booleans, a bound below -1, or left_bound = 0 (which keeps no key at or
+// before the diagonal, and which L - 1 would turn into "unbounded") throws
 // HipdnnPluginException(INVALID_VALUE).
 template <typename SdpaAttrsT>
 DiagonalBand resolveDiagonalBand(const SdpaAttrsT& attrs)
@@ -128,13 +139,22 @@ DiagonalBand resolveDiagonalBand(const SdpaAttrsT& attrs)
             "SDPA: left_bound and right_bound must be >= -1 (got left_bound=" + std::to_string(left)
                 + ", right_bound=" + std::to_string(right) + ")");
     }
+    if(left == 0)
+    {
+        throw hipdnn_plugin_sdk::HipdnnPluginException(
+            HIPDNN_PLUGIN_STATUS_INVALID_VALUE,
+            "SDPA: left_bound must be -1 (unbounded) or >= 1; left_bound = L keeps the L keys "
+            "ending at the diagonal, so 0 keeps no key at or before it");
+    }
+
+    const int64_t windowLeft = left > 0 ? left - 1 : -1;
 
     if(!causalDeprecated && !bottomRightDeprecated)
     {
-        return {left, right, attrs.diagonal_alignment()};
+        return {windowLeft, right, attrs.diagonal_alignment()};
     }
 
-    return {left,
+    return {windowLeft,
             right >= 0 ? right : 0,
             bottomRightDeprecated ? DiagonalAlignment::BOTTOM_RIGHT : attrs.diagonal_alignment()};
 }
@@ -163,8 +183,9 @@ inline MaskType getMaskType(int64_t leftBound,
 // Sliding-window mask coordinate transformation
 // =============================================================================
 //
-// Converts raw window sizes (left_bound, right_bound from the hipDNN graph)
-// into the precomputed mask coordinates (mask_y, mask_x) that the AITER ASM
+// Converts FlashAttention-convention window sizes (a DiagonalBand's bounds, not
+// the graph's cuDNN-semantics left_bound) into the precomputed mask coordinates
+// (mask_y, mask_x) that the AITER ASM
 // DQDKDV kernel expects in its argument struct.
 //
 // AITER reference: ck_tile_shim.h::compute_mask_coordinates()

dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/packs/Gfx1151WmmaAttentionNative.cpp

@@ -396,10 +396,15 @@ std::optional<MaskParameters> maskFor(const data_objects::SdpaAttributes& attrib
         = attributes.left_bound().has_value() ? attributes.left_bound().value() : UNBOUNDED;
     const int64_t right
         = attributes.right_bound().has_value() ? attributes.right_bound().value() : UNBOUNDED;
-    if(left < UNBOUNDED || right < UNBOUNDED || !fitsI32(left) || !fitsI32(right))
+    // left_bound = 0 keeps no key at or before the diagonal; hipDNN rejects it.
+    if(left < UNBOUNDED || left == 0 || right < UNBOUNDED || !fitsI32(left) || !fitsI32(right))
     {
         return std::nullopt;
     }
+    // The graph's left_bound follows cuDNN (L keys ending at the diagonal, diagonal
+    // included); the kernel's window_left follows FlashAttention (keys before the
+    // diagonal), so the same window is L - 1.
+    const int64_t windowLeft = left > 0 ? left - 1 : UNBOUNDED;
 
     MaskParameters result;
     // The deprecated causal flags override any bounds and alignment, exactly as
@@ -413,11 +418,11 @@ std::optional<MaskParameters> maskFor(const data_objects::SdpaAttributes& attrib
     }
     result.bottomRight = static_cast<int32_t>(attributes.diagonal_alignment()
                                               == data_objects::DiagonalAlignment::BOTTOM_RIGHT);
-    result.left = static_cast<int32_t>(left);
+    result.left = static_cast<int32_t>(windowLeft);
     result.right = static_cast<int32_t>(right);
-    if(left != UNBOUNDED || right != UNBOUNDED)
+    if(windowLeft != UNBOUNDED || right != UNBOUNDED)
     {
-        result.kind = left == UNBOUNDED && right == 0 ? MaskKind::CAUSAL : MaskKind::WINDOW;
+        result.kind = windowLeft == UNBOUNDED && right == 0 ? MaskKind::CAUSAL : MaskKind::WINDOW;
     }
     return result;
 }

dnn-providers/hip-kernel-provider/src/tests/engines/asm_sdpa_engine/DeprecatedCausalMaskCases.hpp

@@ -224,6 +224,17 @@ inline const std::vector<DeprecatedCausalMaskCase>& deprecatedCausalMaskCases()
                                  std::nullopt},
         DeprecatedCausalMaskCase{
             "LeftBoundBelowMinusOne", false, false, -2, 0, ALIGN_TOP_LEFT, std::nullopt},
+        // cuDNN semantics: left_bound = L keeps L keys, so 0 keeps no key at or before the
+        // diagonal and is rejected.
+        DeprecatedCausalMaskCase{"LeftBoundZero", false, false, 0, 0, ALIGN_TOP_LEFT, std::nullopt},
+        // left_bound = 1 keeps only the diagonal: a band, not causal.
+        DeprecatedCausalMaskCase{"LeftBoundOneIsDiagonalOnly",
+                                 false,
+                                 false,
+                                 1,
+                                 0,
+                                 ALIGN_TOP_LEFT,
+                                 plan_utils::MaskType::SLIDING_WINDOW},
         DeprecatedCausalMaskCase{"BoundsOnlyTopLeftCausal",
                                  false,
                                  false,

dnn-providers/hip-kernel-provider/src/tests/engines/asm_sdpa_engine/TestSdpaBwdPlanBuilder.cpp

@@ -1194,18 +1194,20 @@ TEST(TestSdpaBwdDeprecatedCausalMaskBand, CausalWithLeftBoundResolvesToCausalWin
               .front()
               ->attributesAs<hipdnn_flatbuffers_sdk::data_objects::SdpaBackwardAttributes>();
 
+    // left_bound = 64 keeps 64 keys (cuDNN), which is FlashAttention's window_size_left 63.
     const auto band = plan_utils::resolveDiagonalBand(attrs);
-    EXPECT_EQ(band.leftBound, 64);
+    EXPECT_EQ(band.leftBound, 63);
     EXPECT_EQ(band.rightBound, 0);
     EXPECT_EQ(band.alignment, ALIGN_BOTTOM_RIGHT);
 
     constexpr int64_t SEQ_LEN_ABOVE_BOUND = 256;
     const auto windowed
         = plan_utils::makeCoveringBoundsUnbound(band, SEQ_LEN_ABOVE_BOUND, SEQ_LEN_ABOVE_BOUND);
-    EXPECT_EQ(windowed.leftBound, 64);
+    EXPECT_EQ(windowed.leftBound, 63);
     EXPECT_EQ(windowed.rightBound, 0);
 
-    constexpr int64_t SEQ_LEN_COVERED_BY_BOUND = 65;
+    // 64 keys cover every row of a 64-long sequence; at 65 the last row would drop key 0.
+    constexpr int64_t SEQ_LEN_COVERED_BY_BOUND = 64;
     const auto covered = plan_utils::makeCoveringBoundsUnbound(
         band, SEQ_LEN_COVERED_BY_BOUND, SEQ_LEN_COVERED_BY_BOUND);
     EXPECT_EQ(covered.leftBound, -1);

dnn-providers/hip-kernel-provider/src/tests/engines/asm_sdpa_engine/TestSdpaFwdPlanBuilder.cpp

@@ -26,6 +26,7 @@
 #include <hipdnn_plugin_sdk/PluginException.hpp>
 #include <hipdnn_test_sdk/utilities/MockEngineConfig.hpp>
 #include <optional>
+#include <utility>
 
 namespace asm_sdpa_engine
 {
@@ -875,6 +876,35 @@ INSTANTIATE_TEST_SUITE_P(,
                          ::testing::ValuesIn(deprecatedCausalMaskCases()),
                          deprecatedCausalMaskCaseName);
 
+// The graph's left_bound follows cuDNN (L keys ending at the diagonal, diagonal
+// included); the band the kernels take is FlashAttention's window_size_left, L - 1.
+TEST(TestSdpaFwdResolveDiagonalBand, ConvertsCudnnLeftBoundToWindowLeft)
+{
+    // Literal expectations, so the test does not restate the conversion it checks.
+    for(const auto& [leftBound, expectedWindowLeft] : {std::pair<int64_t, int64_t>{-1, -1},
+                                                       std::pair<int64_t, int64_t>{1, 0},
+                                                       std::pair<int64_t, int64_t>{2, 1},
+                                                       std::pair<int64_t, int64_t>{64, 63}})
+    {
+        auto builder = createSdpaFwdGraphWithMask(
+            /*causalMask=*/false,
+            /*causalMaskBottomRight=*/false,
+            leftBound,
+            /*rightBound=*/0,
+            hipdnn_flatbuffers_sdk::data_objects::DiagonalAlignment::TOP_LEFT);
+        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>();
+
+        const auto band = plan_utils::resolveDiagonalBand(attrs);
+        EXPECT_EQ(band.leftBound, expectedWindowLeft) << "left_bound=" << leftBound;
+        EXPECT_EQ(band.rightBound, 0);
+    }
+}
+
 // =============================================================================
 // plan_utils::makeCoveringBoundsUnbound
 // =============================================================================

dnn-providers/hip-kernel-provider/src/tests/engines/hip_flash2_engine/TestHipFlash2FwdPlanBuilder.cpp

@@ -334,8 +334,9 @@ INSTANTIATE_TEST_SUITE_P(,
                          ::testing::ValuesIn(asm_sdpa_engine::deprecatedCausalMaskCases()),
                          asm_sdpa_engine::deprecatedCausalMaskCaseName);
 
-// The mask graph is 256x256, so a top-left left bound of 255 covers the whole
-// sequence and is plain causal, while 254 is a genuine sliding window.
+// The mask graph is 256x256. left_bound follows cuDNN (L keeps the L keys ending at
+// the diagonal), so a top-left left bound of 256 covers the whole sequence and is
+// plain causal, while 255 is a genuine sliding window (the last query drops key 0).
 bool extractCausalForTopLeftWindow(int64_t leftBound)
 {
     auto builder = asm_sdpa_engine::createSdpaFwdGraphWithMask(
@@ -352,13 +353,13 @@ bool extractCausalForTopLeftWindow(int64_t leftBound)
 
 TEST(TestHipFlash2ExtractParamsWindow, CoveringLeftBoundIsCausal)
 {
-    EXPECT_TRUE(extractCausalForTopLeftWindow(255));
+    EXPECT_TRUE(extractCausalForTopLeftWindow(256));
     EXPECT_TRUE(extractCausalForTopLeftWindow(4294967298));
 }
 
 TEST(TestHipFlash2ExtractParamsWindow, GenuineWindowIsNotCausal)
 {
-    EXPECT_FALSE(extractCausalForTopLeftWindow(254));
+    EXPECT_FALSE(extractCausalForTopLeftWindow(255));
 }
 
 } // namespace

dnn-providers/hip-kernel-provider/src/tests/engines/kernel_ingestor_engine/packs/TestGfx1151WmmaAttentionMatchers.cpp

@@ -603,6 +603,32 @@ TEST(TestGfx1151WmmaAttentionMatchers, DeclinesOverrideShapeGraphs)
     EXPECT_FALSE(matchGraph(spec).has_value());
 }
 
+// The graph's left_bound follows cuDNN (L keys ending at the diagonal, diagonal
+// included); the kernel's window_left follows FlashAttention (keys before the
+// diagonal), so the same window binds as L - 1. left_bound = 0 is declined.
+TEST(TestGfx1151WmmaAttentionMatchers, BindsCudnnLeftBoundAsFlashAttentionWindowLeft)
+{
+    using hipdnn_plugin_sdk::ingestor::tryGetBoundInt;
+    GraphSpec spec;
+    spec.leftBound = 7;
+    spec.rightBound = 3;
+    const auto bound = matchGraph(spec);
+    ASSERT_TRUE(bound.has_value());
+    EXPECT_EQ(tryGetBoundInt(*bound, "gfx1151_wmma_attention.mask"), 2); // window
+    EXPECT_EQ(tryGetBoundInt(*bound, "gfx1151_wmma_attention.window_left"), 6);
+    EXPECT_EQ(tryGetBoundInt(*bound, "gfx1151_wmma_attention.window_right"), 3);
+
+    spec.leftBound = 1;
+    spec.rightBound = 0;
+    const auto diagonalOnly = matchGraph(spec);
+    ASSERT_TRUE(diagonalOnly.has_value());
+    EXPECT_EQ(tryGetBoundInt(*diagonalOnly, "gfx1151_wmma_attention.mask"), 2); // window
+    EXPECT_EQ(tryGetBoundInt(*diagonalOnly, "gfx1151_wmma_attention.window_left"), 0);
+
+    spec.leftBound = 0;
+    EXPECT_FALSE(matchGraph(spec).has_value());
+}
+
 // The deprecated causal flags override the bounds and alignment, as in the
 // hipDNN reference; otherwise a left bound would drop the causal limit.
 TEST(TestGfx1151WmmaAttentionMatchers, DeprecatedCausalFlagsOverrideBounds)
@@ -942,7 +968,7 @@ TEST(TestGfx1151WmmaAttentionMatchers, DecodeObjectServesTheDispatcherDecodeRule
     window.kvHeads = 8;
     EXPECT_TRUE(matchesKernel(window, decodeKernel(128), gfx12));
     window.bottomRight = true;
-    window.leftBound = 1023;
+    window.leftBound = 1024; // 1024 keys including the diagonal: window_left 1023
     window.rightBound = 0;
     EXPECT_FALSE(matchesKernel(window, decodeKernel(128), gfx12));
 }
@@ -1004,7 +1030,7 @@ TEST(TestGfx1151WmmaAttentionMatchers, DecodeWorkspaceMirrorsTheDispatcherPlan)
     // A bottom-right 1024-key window: 2 splits (still merged by the merge kernel).
     GraphSpec window = decodeGraph();
     window.bottomRight = true;
-    window.leftBound = 1023;
+    window.leftBound = 1024; // 1024 keys including the diagonal: window_left 1023
     window.rightBound = 0;
     EXPECT_EQ(workspaceOf(window, decodeKernel(128), gfx1201), 33024U);
     // Batch 4: 2 splits.

dnn-providers/integration-tests/gpu-ref/kernels/sdpa/GpuRefSdpaFwd.cpp

@@ -126,8 +126,9 @@ extern "C" __global__ void sdpaFwdRef(SdpaFwdArgs args)
             s += toAccum(mask[maskOffset]);
         }
 
-        // (b) Sliding-window mask (applied after the additive mask).
-        // Asymmetric: +1 on the right bound, none on the left bound.
+        // (b) Sliding-window mask (applied after the additive mask), cuDNN semantics:
+        // keeps skv in [sq + windowOffset - leftBound + 1, sq + windowOffset + rightBound],
+        // so leftBound = L keeps the L keys ending at the diagonal.
         if(args.rightBound >= 0)
         {
             long long startKv = sq + 1 + windowOffset + args.rightBound;
@@ -142,7 +143,7 @@ extern "C" __global__ void sdpaFwdRef(SdpaFwdArgs args)
         }
         if(args.leftBound >= 0)
         {
-            if(skv < sq + windowOffset - args.leftBound)
+            if(skv <= sq + windowOffset - args.leftBound)
             {
                 s = negInf;
             }

dnn-providers/integration-tests/gpu-ref/kernels/sdpa/GpuRefSdpaRaggedFwd.cpp

@@ -182,7 +182,8 @@ extern "C" __global__ void sdpaRaggedFwdRef(SdpaRaggedFwdArgs args)
         // Fold in the fp8 Q/K descale.
         COMPUTE_TYPE s = dot * descaleQK * static_cast<COMPUTE_TYPE>(args.scale);
 
-        // Sliding-window mask, aligned per batch. Only the right bound gets the +1.
+        // Sliding-window mask, aligned per batch, cuDNN semantics: keeps skv in
+        // [sq + windowOffset - leftBound + 1, sq + windowOffset + rightBound].
         // No additive bias: it is gated off on the ASM v3 path.
         if(args.rightBound >= 0)
         {
@@ -198,7 +199,7 @@ extern "C" __global__ void sdpaRaggedFwdRef(SdpaRaggedFwdArgs args)
         }
         if(args.leftBound >= 0)
         {
-            if(skv < sq + windowOffset - args.leftBound)
+            if(skv <= sq + windowOffset - args.leftBound)
             {
                 s = negInf;
             }

dnn-providers/integration-tests/reference-data-scripts/README.md

@@ -124,3 +124,18 @@ and for template-sweep bundles:
 - `golden.path` validation and per-case tensor file/NaN/size checks (input and output tensors)
 - per-case metadata field validation
 - a directory with only one of `graph.template.json`/`sweep.json`
+
+## Sliding-window bounds in the generators
+
+A left window bound has two common conventions, and each generator flag is named after the
+one it uses:
+
+| Convention | Name | `L = 3` keeps, for the query on the diagonal |
+|---|---|---|
+| hipDNN graph, cuDNN | `left_bound` (`--left-bound`) | the 3 keys ending at the diagonal, diagonal included |
+| FlashAttention, AITER | `window_size_left` (`--window-left`) | the 3 keys before the diagonal, plus the diagonal (4 keys) |
+
+`generate_sdpa_fwd_golden.py --left-bound L` writes `left_bound = L` to the graph.
+`--window-left W` writes `left_bound = W + 1`; it exists for `--validate`, which passes the
+window to AITER. The two flags are mutually exclusive. `left_bound` must be positive or -1
+(unbounded).