4 Commits

Author	SHA1	Message	Date
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