diff --git a/csrc/attention/flash_attention.cu b/csrc/attention/flash_attention.cu
index b6c3c51..8086b38 100644
--- a/csrc/attention/flash_attention.cu
+++ b/csrc/attention/flash_attention.cu
@@ -216,7 +216,9 @@ __global__ void decode_attention_bf16_kernel(
     __nv_bfloat16* __restrict__ O,
     int num_q_heads, int num_kv_heads,
     int kv_len, int head_dim,
-    float scale
+    float scale,
+    const float* __restrict__ sinks,  // [num_q_heads] or nullptr (gpt-oss)
+    int window                        // >0: only attend last `window` keys; 0: full
 ) {
     int bh = blockIdx.x;
     int batch_idx = bh / num_q_heads;
@@ -228,6 +230,10 @@ __global__ void decode_attention_bf16_kernel(
 
     int tid = threadIdx.x;
 
+    // Sliding window: the query is at position kv_len-1 (decode appends it), so it
+    // attends to keys [kv_start, kv_len). kv_start=0 for full attention.
+    int kv_start = (window > 0 && kv_len > window) ? (kv_len - window) : 0;
+
     // Pointers to this batch/head's data
     // Q: [batch, num_q_heads, 1, head_dim]
     const __nv_bfloat16* Q_ptr = Q + ((long long)batch_idx * num_q_heads + q_head) * head_dim;
@@ -251,7 +257,7 @@ __global__ void decode_attention_bf16_kernel(
         local_O[d] = 0.0f;
     }
 
-    for (int pos = tid; pos < kv_len; pos += DECODE_THREADS) {
+    for (int pos = kv_start + tid; pos < kv_len; pos += DECODE_THREADS) {
         // Compute dot(Q, K[pos]) * scale
         const __nv_bfloat16* K_pos = K_base + pos * head_dim;
         float dot = 0.0f;
@@ -336,8 +342,20 @@ __global__ void decode_attention_bf16_kernel(
     __syncthreads();
     global_sum = smem_sum[0];
 
+    // gpt-oss attention sink: a per-head learned logit that joins the softmax
+    // denominator but contributes no value. Rebase max to include it, then add
+    // exp(sink - m_new) to the denominator. O was accumulated relative to
+    // global_max, so it picks up a factor exp(global_max - m_new).
+    float o_scale = 1.0f;
+    if (sinks != nullptr) {
+        float sink = sinks[q_head];
+        float m_new = fmaxf(global_max, sink);
+        o_scale = expf(global_max - m_new);
+        global_sum = global_sum * o_scale + expf(sink - m_new);
+    }
+
     // Step 4: Reduce O across block (dimension by dimension using shared mem)
-    float inv_sum = (global_sum > 0.0f) ? (1.0f / global_sum) : 0.0f;
+    float inv_sum = (global_sum > 0.0f) ? (o_scale / global_sum) : 0.0f;
 
     // Process head_dim in chunks: each iteration reduces one dimension
     // Use shared memory accumulator: each warp contributes via warp reduction + atomic