Replace shared-memory butterfly stages h=1,2,4,8,16 with
__shfl_xor_sync() in turbo3, turbo2, and turbo4 SET_ROWS kernels.
Pairs with distance < 32 are always in the same warp — no barrier
needed. Only h=32 and h=64 (cross-warp) retain __syncthreads().

Mathematically identical: __shfl_xor_sync(mask, v, h) gives thread j
the value from j^h, then (j & h) ? (other - v) : (v + other) is the
exact butterfly computation.

Zero PPL regression: turbo3 6.24, turbo4 6.23 (unchanged).

Credit to seanrasch (perf/ftz-and-wht-shuffle branch).

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Auto-asymmetric (K→q8_0) creates FA dimension mismatch on models
where head_dim_k != head_dim_v (e.g. DeepSeek: K=192, V=128). The
q8_0 K at non-standard D has no CUDA FA kernel, falling to slow
CPU FA.

Fix: when K and V head dims differ, keep symmetric turbo which pads
both K and V consistently for CUDA FA.

Note: turbo4 symmetric on DeepSeek still falls to CPU FA (21 t/s)
because padded K D=256 != V D=128. This is a pre-existing FA D
matching limitation, not a regression. turbo3 on DeepSeek works at
full CUDA speed (172 t/s) because the VEC kernel handles asymmetric
D natively.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Precompute Q[d] × centroid[c] into shared memory LUT once per decode
step, then the KQ inner loop does a single shmem read per element
instead of centroid lookup + multiply.

- turbo3: 127→146 t/s at 32K depth (+15%)
- turbo2: 148→154 t/s at 32K depth (+4%)
- Only for ncols==1 (decode path, not prefill)
- turbo4 excluded: 16 centroids × D exceeds shmem budget
- LUT stride = n_centroids+1 to avoid bank conflicts
- 8-wide processing (2 qs bytes + 1 signs byte per iteration)
- L2 prefetch hints for next K block

PPL unchanged (6.31). Zero quality impact — the LUT stores the exact
same Q×centroid values, just precomputed.

Based on Madreag/turbo3-cuda (release/cuda-optimized branch).

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Four general VEC flash attention optimizations:

1. __launch_bounds__ occupancy 1→3: allows 3 blocks per SM, better
   latency hiding. This is the biggest win (~16% on q8_0 baseline).

2. V_is_unquantized: remove turbo types from V unquantized path,
   matching the K change. Turbo V uses quantized dequant path.

3. Aggressive sparse V threshold: 5e-3 for turbo3/4 (was 1e-6),
   1e-2 for turbo2. Validated zero PPL impact per Madreag.

4. L2 prefetch: add V block prefetch alongside existing K prefetch.

Results (tg128 @ d32768, Qwen3.5 35B Q4_K_M, RTX 5090):

| Type | Before | After | vs q8_0 |
|------|--------|-------|---------|
| q8_0 | 156 | 181 | baseline |
| turbo2 | 148 | 185 | +2.5% faster |
| turbo3 | 127 | 171 | -5.5% |
| turbo4 | 101 | 117 | -35% (no LUT) |

turbo2 at 7.5x compression is now FASTER than q8_0 at 32K context.
PPL unchanged (6.31).

Based on Madreag/turbo3-cuda optimizations.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>