Commit Graph

4 Commits

Author SHA1 Message Date
Claude
39e32343eb [ARCH] CCCL-derived paged attention kernel architecture + Triton rewrite
Architecture document: docs/paged_attention_kernel_architecture.md
Defines every module from CCCL algorithm patterns before code.

Three-level decomposition from CCCL:
  Level 1 (warp_reduce_shfl): shfl.down butterfly for per-thread QK scores
  Level 2 (block_reduce_warp_reductions): warp partials → SMEM → block aggregate
  Level 3 (agent_scan decoupled lookback): cross-partition combine

Compound type (from summary_statistics.cu):
  attention_partial = (max_score, exp_sum, weighted_v[256])
  combine(a, b) = online softmax rescaling (same math as Flash Attention)

Key design change: Grid on num_kv_heads, not num_heads.
  Before: grid = (1, 24, 200) = 4800 blocks, KV loaded 6x redundantly
  After:  grid = (1, 4, 200) = 800 blocks, KV loaded once per kv_head
  Each block computes GQA_RATIO=6 query heads with shared KV loads.
  Reduces KV cache bandwidth by 6x (the GQA ratio).

SMEM budget verified:
  K tile [32, 256] fp16 = 16KB
  V tile [32, 256] fp16 = 16KB
  Total = 32KB ≤ 48KB ✓

Phase 1 kernel: _partition_attn_kernel
  Processes query heads sequentially within the GQA group
  to minimize register pressure (6 × 256 = 1536 registers
  too many if all loaded simultaneously).

Phase 2 kernel: _reduce_partitions_kernel
  Also gridded on kv_heads, reduces all partitions for
  GQA_RATIO heads per block.

This replaces the previous Triton V2 which was gridded on num_heads
and had no GQA awareness at the kernel level.
2026-07-31 04:13:07 +00:00
Claude
cd0d9e1a91 [OPT] Fix online softmax bug in Triton V2 Phase 1
Bug: if l_i > 0 branch in Triton is invalid (compiled as constexpr).
Also: p = exp(scores - m_i_new) computed after m_i_new update was
using the wrong reference max (should subtract m_ij first, then rescale).

Fix: Adapted exactly from prefix_prefill.py's proven-correct pattern:
  p = exp(scores - m_ij)          # probs relative to chunk max
  l_ij = sum(p)                   # chunk sum
  m_i_new = max(m_i, m_ij)       # new running max
  alpha = exp(m_i - m_i_new)     # old accumulator rescale
  beta = exp(m_ij - m_i_new)     # new chunk rescale
  l_i_new = alpha*l_i + beta*l_ij
  acc = acc*(alpha*l_i/l_i_new) + (p*beta/l_i_new) @ V

This is the Flash Attention online softmax tiling algorithm.
Same math as CCCL's parallel_reduce with compound accumulators.
2026-07-30 16:16:56 +00:00
Claude
33f6ead1b8 [OPT] Complete Triton V2 Phase 1 — paged K/V gather from prefix_prefill.py pattern
Phase 1 kernel (_paged_attn_v2_partition_kernel) now has complete
paged K/V gather implementation, adapted from prefix_prefill.py:

  K gather:
    bn = tl.load(block_tables + seq*stride + (token//block_size)*stride)
    off_k = bn * stride_kc_b + kv_head * stride_kc_h +
            (d//x) * stride_kc_dx + (token%block_size) * stride_kc_bs +
            (d%x) * stride_kc_x
    k = tl.load(key_cache + off_k, mask=valid)

  V gather (simpler layout):
    off_v = bn * stride_vc_b + kv_head * stride_vc_h +
            d * stride_vc_d + (token%block_size) * stride_vc_bs

  Online softmax (Flash Attention pattern):
    m_i_new = max(m_i, max(scores))
    alpha = exp(m_i - m_i_new)
    acc = acc * alpha * l_i / l_i_new + (p/l_i_new * beta) @ V

Key difference from prefix_prefill.py:
  - BLOCK_M=1 (decode: 1 query token) vs BLOCK_M>1 (prefill)
  - q @ k is dot product [D]•[D,N] → [N], not matrix [M,D]@[D,N] → [M,N]
  - head_dim=256 support: BLOCK_N=32 (vs 64 for head_dim=128)
    32×256×2×2 = 32KB ≤ 48KB SMEM ✓

Integration: Triton V2 tried first, PyTorch V2 as fallback.
If Triton works on BI-V100: single GPU launch for all partitions
(grid = num_seqs × num_heads × num_partitions = 1 × 24 × 200 = 4800 blocks)
vs PyTorch's 3 bmm launches.
2026-07-30 16:07:15 +00:00
Claude
a53d1a28b0 [OPT] Triton paged_attention_v2 kernel skeleton — Phase 2 reduction complete
Two-kernel design following vllm's paged_attention_v2_kernel.cu:

Phase 1: _paged_attn_v2_partition_kernel
  grid = (num_seqs, num_heads, num_partitions)
  Each instance: Q[head] @ K[partition]^T → softmax → @ V[partition]
  Status: SKELETON — paged K/V gather from indirect block_tables
  is complex in Triton (requires scatter/gather through block_tables).
  Currently falls back to PyTorch partition loop.

Phase 2: _paged_attn_v2_reduce_kernel
  grid = (num_seqs, num_heads)
  Each instance: log-sum-exp reduction across partitions
  Status: COMPLETE — replaces Python einsum with single Triton launch.
  Algorithm: global_max → rescale → weighted sum (same pattern as
  CCCL summary_statistics binary_op for combining partial statistics).

SMEM: Phase 1 needs BLOCK_N=64 × head_dim=128 × 2B × 2 = 32KB ≤ 48KB.
Phase 2 needs no SMEM (partitions fit in registers).

The Phase 1 paged gather is the hard part. The key_cache layout
[blocks, kv_heads, head_dim/x, block_size, x] requires:
  1. block_tables[seq, token // block_size] → physical_block_id
  2. key_cache[physical_block_id, kv_head, :, token % block_size, :]
This is indirect indexed access — possible in Triton via tl.load with
computed offsets, but needs careful stride arithmetic.
2026-07-30 15:59:15 +00:00