4 Commits

Include renames

Author	SHA1	Message	Date
ViperEkura	8055027df7	perf: enable paged MMA kernel for page_size=1 ... - Replace per-tile page lookup with per-element lookup in load_tile - Remove page_ok gate and scalar fallback in launch_paged_decode_mma - Unified path works for any page_size (L1-cached when page_size >= BC) - HBM BW: 12% → 73%, decode throughput: 2,250 → 2,606 tok/s (B=32) - Scales to 5,232 tok/s at B=128 (2.54x vs torch native)	2026-07-30 22:06:41 +08:00
ViperEkura	2e29ed45d3	perf: shrink decode tile to BC=16 for higher occupancy ... - BC=32→16 halves smem (32KB→16KB for D=128), doubling blocks/SM (3→6) - D=256 now fits STAGES=2 double-buffer in 32KB, eliminating 176-byte spill - min_tiles_per_split=2 avoids excessive split overhead on small kv - paged decode: require page_size multiple of BC so tiles stay page-aligned Benchmark (L20 sm_89, D=128): - B=1 kv=4096: 0.0134→0.0122ms (+9% BW) - B=16 kv=2048: 0.0434→0.0352ms (+23% BW) - B=32 kv=1024: 0.0343→0.0282ms (+22% BW)	2026-07-27 22:44:02 +08:00
ViperEkura	20041d7aa9	perf: extend MMA decode to arbitrary GQA ratio, add launch bounds, vectorize combine ... - Multi-pass MMA: encode pass in grid blockIdx.x, compute q_head0/G in-kernel - Fixes crash for G>32 (previously block(32,G) exceeded 1024 threads) - Fixes alloc_split_partials using uninitialized num_splits (MAX_SPLITS=32) - __launch_bounds__ on all MMA and prefill kernels for better register allocation - 4x vectorized combine kernel (4 head_dim per thread) - uint4 vectorized K loads in scalar decode kernels - cp.async .L2::128B cache hint for K/V tile streaming - Extract warp_reduce_sum, bf16, MAX_SPLITS to attn_warp_utils.cuh	2026-07-27 00:35:34 +08:00
ViperEkura	f7a16efc9d	refactor: extract shared dispatcher header, unify MMA/scalar dispatch format ... - Merge 3 duplicated dispatch blocks into single attn_dispatchers.cuh - Merge compute_num_splits from attn_utils.cuh into dispatcher header - All dim3 grid/block declarations and <<<>>> launches are single-line - Production .cu files (35-42 loc) only handle torch wrapping + pybind11 - Test files include dispatcher header directly, removing all #ifndef ASTRAI_NO_MMA duplication	2026-07-21 22:21:39 +08:00