Commit Graph

3 Commits

Author SHA1 Message Date
muh-bot
cfa6516cc6 [BUGFIX] patch_paged_attention_v2.py: NameError V2_MODULE undefined → V2_MODULE_PYTORCH
deploy_v2_module() 引用 V2_MODULE 但文件只定义了 V2_MODULE_PYTORCH 和 V2_MODULE_TRITON。
这导致 Docker build 时 NameError → V2 module 部署静默失败 → _custom_ops.py 的
V2 import 会 ImportError → V2 路径不可用。

修复: 4 处 V2_MODULE 引用全部改为 V2_MODULE_PYTORCH

影响: 如果 Dockerfile 启用了 V2 (patch_paged_attention_v2.py),
此 bug 意味着 V2 module 从未被正确部署。V2 的 import 总是失败,
paged_attn.py 的 V2 路径总是走到 except 分支。

这实际上是一个'幸运的 bug'——因为 V2 PyTorch 比 V1 ixformer 慢,
V2 部署失败反而保护了性能。但它也意味着如果未来需要 V2,
必须先修这个 bug。

功能测试影响: 无 (V2 不影响功能测试, V1 已够用)
效果测试影响: 正面 (V1 ixformer 精度一致性好于 V2 PyTorch)
2026-08-05 07:07:54 +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
9cb7f9d037 [OPT] PagedAttention V2 implementation — fill the NotImplementedError hole
The single biggest performance bottleneck in the baseline:
paged_attention_v2 = raise NotImplementedError()
paged_attn.py: use_v1 = True (hardcoded to avoid calling V2)

V1 limitation: processes entire KV sequence in one kernel launch.
For seq_len=100K, this is a single massive attention computation.
V2: splits into PARTITION_SIZE=512 chunks, runs them in parallel,
then reduces with log-sum-exp. 195 parallel partitions vs 1.

Implementation (paged_attention_v2_pytorch.py):
  Phase 1: Per-partition attention
    - For each (seq, head, partition): compute QK^T, softmax, weighted V sum
    - Store partial: tmp_output, exp_sums, max_logits (per partition)
  Phase 2: Cross-partition reduction (log-sum-exp)
    - global_max = max(max_logits across partitions)
    - rescale = exp(partition_max - global_max) × partition_exp_sum
    - output = Σ (rescale / total_sum) × partition_output

This is the same algorithm as vllm's paged_attention_v2_kernel.cu:
  - The reduction pattern is identical to CCCL's block_reduce_warp_reductions
    (combine partial statistics from independent segments)
  - The online softmax tiling is the same as Flash Attention's partitioning

Integration:
  - patch_paged_attention_v2.py patches _custom_ops.py and paged_attn.py
  - Removes use_v1=True hardcode → V2 used for seq_len > 8192
  - Dockerfile adds the patch step

This is a PyTorch implementation (no CUDA compilation needed).
Next step: if /usr/local/corex/ has ixcc or nvcc-compatible compiler,
replace with compiled CUDA kernel for further speedup.
2026-07-30 15:40:14 +00:00