Commit Graph

9 Commits

Author SHA1 Message Date
ViperEkura 9e63cb9ed0 feat: MMA head-packing decode kernel with scalar fallback dispatch
- Add gqa_decode_attn_mma.cuh for tensor-core decode path
- Add dispatch_decode<> selecting MMA vs scalar based on G and mask
- Add TORCH_CHECK for unsupported head_dim instead of silent scalar launch
2026-07-07 22:56:02 +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 f1cc7cedce feat: ldmatrix + smem padding for mma prefill kernel
- replace scalar fragment loads with ldmatrix.sync.x4/x2
- add smem row-stride padding (LD = HEAD_DIM + 8) to eliminate 8-way bank conflicts from HEAD_DIM being a 32-bank multiple
- switch build flag from positive to negative: -DASTRAI_NO_MMA for pre-sm_80 only; mma is the default path
- vectorize scalar path smem loads with float4 ld8
- fix pure-C test configs for ld8 alignment
2026-07-06 20:55:22 +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 cc36530c73 perf: group-split register-blocking gqa_prefill kernel
- one query row per group of G=8 lanes, each owning HEAD_DIM/G dims of qreg[]/acc[] in registers
- removes full 32-lane warp_reduce_sum; S dot reduces over only G lanes
- templated on <HEAD_DIM,G,ROWS,P_BC>, block=(G,ROWS)=(8,32)
- per-group shuffle mask so causal loop-bound divergence doesn't deadlock the shuffle
- update pure-C test to the templated launch
2026-07-06 18:33:08 +08:00
ViperEkura 11fa807cfc fix: correct prefill mask index, unify GQA kernel interface
- Fix mask indexing: batch*q_len*kv_len -> batch*kv_len
- Add csrc/kernels/gqa_common.cuh with shared GQAParams struct
- Unify decode/prefill Python API: both accept (q,k,v,mask=None,...)
- Decode now supports optional mask, is_causal, causal_offset, scale
- Rename struct fields: B->batch, Hq->q_head, Hk->kv_head, D->head_dim
- Use py::arg() for correct None/defaults handling in pybind11
- Update pure C tests and build instructions (-arch=sm_89)
2026-07-06 17:21:23 +08:00
ViperEkura bcdd93e0eb feat: split kernel defs from bindings, add prefill tiled kernel and pure C tests
- Split .cuh/.cu for gqa_decode_attn and gqa_prefill_attn
- gqa_prefill_attn: tiled shared-memory K/V, fused load, compute-opt, mask support
- Add pure C tests under csrc/tests/ for fast nvcc-only iteration
- Update .gitignore for build artifacts
2026-07-06 16:14:55 +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