ViperEkura/AstrAI
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

perf: drop redundant smem zero-init in paged decode kernel

Browse Source 
- 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
This commit is contained in:
ViperEkura
2026-08-01 16:17:34 +08:00
parent 4b25664c79
commit fda82ee232
1 changed files with 0 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files
csrc/kernels/attn_paged_decode_split_kv_mma.cuh
-7
View File
@@ -39,13 +39,6 @@ __global__ void paged_attn_decode_split_kv_mma_kernel(PagedAttentionParams<bf16>
__shared__ __align__(16) bf16 sK[Traits::STAGES * Traits::BC * Traits::LD]; __shared__ __align__(16) bf16 sK[Traits::STAGES * Traits::BC * Traits::LD];
__shared__ __align__(16) bf16 sV[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 q_base = batch * p.q_stride_l + q_head0 * p.q_stride_h;
const int qra = gid; const int qra = gid;
const int qrb = gid + 8; const int qrb = gid + 8;