3 Commits

Include renames

Author	SHA1	Message	Date
ViperEkura	a1a1a6bf0f	perf: pack gqa q-heads per prefill block to reuse kv tiles ... - pack HB = min(G, WARPS) q heads per block; K/V tiles stream once per block instead of once per q head - G=1 keeps the old grid; paged path splits 64-row host Q tiles into HB blocks along grid.x (host maps unchanged) Benchmark: NVIDIA RTX 6000D, short-q/long-kv prefill 1.4-3.4x (G=8 B=16 q=16 kv=16k 4.22 -> 1.26 ms); full prefill/MHA/paged unchanged (compute-bound); verified vs SDPA G in {1,2,3,4,8,32}, 99 tests pass	2026-08-31 13:39:47 +08:00
ViperEkura	31ca357c61	refactor: namespace csrc kernels and extract common helpers ... - attention family -> astrai::attention; fp8 family -> astrai::fp8 - new common/reduce.cuh (warp/group reductions, atomic_max_float) - new common/cp_async.cuh (predicated cp_async_16, commit/wait group) - move MAX_SPLITS into attention/common.h; delete warp_utils.cuh - .cu bindings and pure C tests open family namespaces via using	2026-08-24 14:49:55 +08:00
ViperEkura	16a55bb474	refactor: reorganize CUDA kernels into per-family directories ... - move attention kernels to csrc/kernels/attention/ and rotary to rotary/ - add shared common/mma.cuh (mma_sync, ldmatrix) and device.cuh (sm checks) - split fp8_mm into three-layer fp8/common.h, gemm.cuh, mm.cu - fix fused FP8 GEMM ldmatrix lane indexing to fix OOB shared reads - update extension ops, loader, and kernel tests	2026-08-22 20:40:31 +08:00