perf: cp.async K/V loads, shared sQ staging, causal skip, XOR swizzle
- cp.async global→shared for K/V full-tile loads, eliminates 99.6% of shared-store bank conflicts (612K→2.7K per ncu)
- add cp_async_16/commit/wait_all/wait_group<N> helpers in mma utils
- shared sQ staging (single area, serialized per-warp load), cuts smem from (2*BC + WARPS*BR)*LD to (2*BC + BR)*LD bf16
- pre-scale Q by attention scale during Q load, removes per-tile scale multiply in softmax loop
- causal tile skipping: block-level early break + warp-level skip
- scalar fallback only for last partial tile
- XOR swizzle (swiz_col) at 8-bf16 chunk granularity, eliminates ldmatrix bank conflicts without LD padding, LD=HEAD_DIM (zero smem waste), saves 1280 bytes/block vs HEAD_DIM+8 padding
ViperEkura
9ebaea840f