From fda82ee232389b11dad61133520cf16395e7ec8d Mon Sep 17 00:00:00 2001 From: ViperEkura <3081035982@qq.com> Date: Sat, 1 Aug 2026 16:17:34 +0800 Subject: [PATCH] perf: drop redundant smem zero-init in paged decode kernel - Removes per-step STAGES*BC*LD smem clear loop (2 buffers x 24 layers) - cp.async predicated load + softmax mask already exclude padding slots, matching the paged prefill kernel which never zero-inits - Standalone and extension tests pass; decode step time unchanged --- csrc/kernels/attn_paged_decode_split_kv_mma.cuh | 7 ------- 1 file changed, 7 deletions(-) diff --git a/csrc/kernels/attn_paged_decode_split_kv_mma.cuh b/csrc/kernels/attn_paged_decode_split_kv_mma.cuh index 32c0fdf..a67bf72 100644 --- a/csrc/kernels/attn_paged_decode_split_kv_mma.cuh +++ b/csrc/kernels/attn_paged_decode_split_kv_mma.cuh @@ -39,13 +39,6 @@ __global__ void paged_attn_decode_split_kv_mma_kernel(PagedAttentionParams __shared__ __align__(16) bf16 sK[Traits::STAGES * Traits::BC * Traits::LD]; __shared__ __align__(16) bf16 sV[Traits::STAGES * Traits::BC * Traits::LD]; - #pragma unroll - for (int i = lane; i < Traits::STAGES * Traits::BC * Traits::LD; i += 32) { - sK[i] = __float2bfloat16(0.0f); - sV[i] = __float2bfloat16(0.0f); - } - __syncwarp(); - const int q_base = batch * p.q_stride_l + q_head0 * p.q_stride_h; const int qra = gid; const int qrb = gid + 8;