Commit Graph

8 Commits

Author SHA1 Message Date
ViperEkura 8f89c82d55 chore: bump version to 1.3.9 2026-07-12 21:04:20 +08:00
ViperEkura 89ece26c25 feat: paged decode attention with split-KV (scalar + MMA)
- PagedAttentionParams merged into attn_common.h
- Scalar variant: warp-per-query-head split-KV, resolves page table per-position for the K/V shared-memory tile load
- MMA variant (sm_80+): tensor-core head-packing with cp.async, single page-table lookup per tile (BC=32 fits within page_size>=32)
- Standalone test: 14 cases across head_dim 32/64/128/256, GQA, multi-batch, both paths verified against CPU reference
2026-07-11 18:14:51 +08:00
ViperEkura a4ae7d17fb perf: increase decode split-K parallelism for short sequences
- Remove tiles_total/8 min-work cap that limited splits for small workloads
- Simplify decode_num_splits to only use base_blocks and tiles_total
- Short sequences now generate more blocks, improving SM utilization
2026-07-11 11:26:42 +08:00
ViperEkura d923ebe38d refactor: rename gqa_* to attn_*, split-KV for all decode paths
- Rename all csrc/kernels/gqa_*.cuh/cu to attn_*, with _split_q / _split_kv
  strategy suffix and optional _mma compute suffix
- Remove non-split MMA decode kernel, keep only split-KV path
- Convert scalar decode fallback to split-KV (o_part/ml_part + combine)
- Move combine kernel to attn_decode_split_kv.cuh (shared by both paths)
- Rename GQAParams to AttentionParams
- Update all C++ #include, PYBIND11, and Python extension references
2026-07-10 23:35:14 +08:00
ViperEkura 4225518cf3 perf: add fast-math and vectorization nvcc/cxx build flags
- centralize CXX_FLAGS/NVCC_FLAGS in csrc/build.py as single source
- add --use_fast_math, --ptxas-options=-O3,-v, --extra-device-vectorization
- add -march=native -funroll-loops host flags
- setup.py reads shared cxx_flags/nvcc_flags from registry
- sync pure-C test build commands with new flags
2026-07-07 22:28:32 +08:00
ViperEkura ddc4bd1cf6 feat: tensor-core mma prefill with build-time dispatch
- add register-resident flash-attention kernel using mma.sync.m16n8k16
- dispatch mma vs scalar at build time: pre-sm_80 defines
  -DASTRAI_NO_MMA, else defaults to mma
- scalar path vectorized with float4 smem loads (ld8)
2026-07-06 20:33:24 +08:00
ViperEkura 579b8c3129 fix: correct gqa_decode_attn reduction + add gqa_prefill_attn
- gqa_decode_attn: rewrite to per-KV-head, K in smem
- gqa_prefill_attn: new kernel for Q_len > 1 with GQA
2026-07-06 13:45:18 +08:00
ViperEkura e8e228d035 feat: add optional CUDA kernel system (csrc/) + fused GQA decode attention
Structure:
  csrc/               -- .cu sources + build.py registry
  astrai/extension/   -- compiled .so + __init__.py (import dispatcher)
  setup.py            -- CUDAExtension from csrc/build.py REGISTRY

Control: CSRC_KERNELS=true|false env var at install time.
Fallback: astrai.extension.available dict for runtime detection.
2026-07-06 12:09:58 +08:00