diff --git a/csrc/attention/flash_attention.cu b/csrc/attention/flash_attention.cu
index a891122..c9c69cd 100644
--- a/csrc/attention/flash_attention.cu
+++ b/csrc/attention/flash_attention.cu
@@ -306,16 +306,27 @@ __global__ void flash_attention_sinks_bf16_kernel(
                 row_max = fmaxf(row_max, s);
             }
 
-            float m_new = fmaxf(m_val, row_max);
-            float psum = 0.0f;
-            for (int c = 0; c < kv_tile_cols; c++) {
-                P[c] = expf(P[c] - m_new);
-                psum += P[c];
+            // A fully-masked KV tile (every key causal- or window-masked) has
+            // row_max == -INFINITY. Folding it in computes expf(-inf - (-inf))
+            // = NaN, and a later valid tile's 0*NaN correction then poisons the
+            // whole row. This happens for sliding-window layers whenever a
+            // query's window starts past an early tile (the causal `continue`
+            // above only skips fully-future tiles, not out-of-window ones).
+            // A masked tile contributes nothing to the softmax — skip it.
+            if (row_max != -INFINITY) {
+                float m_new = fmaxf(m_val, row_max);
+                float psum = 0.0f;
+                for (int c = 0; c < kv_tile_cols; c++) {
+                    P[c] = expf(P[c] - m_new);
+                    psum += P[c];
+                }
+                float correction = expf(m_val - m_new);
+                l_val = correction * l_val + psum;
+                for (int d = 0; d < head_dim; d++) O_acc[d] *= correction;
+                m_val = m_new;
+            } else {
+                for (int c = 0; c < kv_tile_cols; c++) P[c] = 0.0f;
             }
-            float correction = expf(m_val - m_new);
-            l_val = correction * l_val + psum;
-            for (int d = 0; d < head_dim; d++) O_acc[d] *= correction;
-            m_val = m_new;
         }
 
         __syncthreads();