From 096a48b5cbdbd2fa1f3e02a7ca06c780c51b493a Mon Sep 17 00:00:00 2001 From: biondizzle Date: Sat, 23 May 2026 19:37:34 +0000 Subject: [PATCH] SMEM-P: fix scoping - define tTMEM_LOADcS_frg unconditionally --- dsv4/kernels/attention/fmha.py | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/dsv4/kernels/attention/fmha.py b/dsv4/kernels/attention/fmha.py index 531d4ccf..96a11b19 100644 --- a/dsv4/kernels/attention/fmha.py +++ b/dsv4/kernels/attention/fmha.py @@ -316,9 +316,9 @@ class FmhaKernel: frg_cnt = 4 frg_tile = cute.size(tTMEM_LOADrS) // frg_cnt tTMEM_LOADrS_frg = cute.logical_divide(tTMEM_LOADrS, cute.make_layout(frg_tile)) - # Coordinate fragments for SMEM-P mapping + # Coordinate fragments for SMEM-P mapping (needed unconditionally for scoping) + tTMEM_LOADcS_frg = cute.logical_divide(tTMEM_LOADcS, cute.make_layout(frg_tile)) if self.use_smem_p: - tTMEM_LOADcS_frg = cute.logical_divide(tTMEM_LOADcS, cute.make_layout(frg_tile)) print(f"[SMEM-P CUTLASS] Created tTMEM_LOADcS_frg shape: {cute.shape(tTMEM_LOADcS_frg)}") for j in range(frg_cnt):