From 07ba03bcbfd0d51cc10ff6d81cf85fd5215c2a98 Mon Sep 17 00:00:00 2001 From: root Date: Wed, 1 Apr 2026 18:00:19 +0000 Subject: [PATCH] Fix sliding window mask: use window_generic when left >= 0 mask_info::decode('b:left,right,sink') always created mask_bottom_right (IsLocal=false) which ignores the left window boundary. For sliding window attention (left >= 0), use window_generic (IsLocal=true) so the kernel respects the left boundary. Fixes: CK split-KV producing identical results with and without sliding window. Now 724/724 shapes pass correctness vs Triton. Made-with: Cursor --- example/ck_tile/01_fmha/mask.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/example/ck_tile/01_fmha/mask.hpp b/example/ck_tile/01_fmha/mask.hpp index 03e1537c5d..bc0cf1e25b 100644 --- a/example/ck_tile/01_fmha/mask.hpp +++ b/example/ck_tile/01_fmha/mask.hpp @@ -122,7 +122,7 @@ struct mask_info v1 = atoi(v.substr(found_1 + 1).c_str()); sink = 0; } - tmp.type = mask_enum::mask_bottom_right; + tmp.type = (v0 >= 0) ? mask_enum::window_generic : mask_enum::mask_bottom_right; auto r = ck_tile::make_generic_attention_mask_coordinates_from_lr_window( v0, v1, sink, y_total, x_total, false); tmp.y = r.at(ck_tile::number<0>{});