- inline decode_use_mma() into dispatch_decode() - drop G>1 guard, MMA works correctly for G>=1 - decode is memory-bound, tensor cores + cp.async still win at G=1
- inline decode_use_mma() into dispatch_decode() - drop G>1 guard, MMA works correctly for G>=1 - decode is memory-bound, tensor cores + cp.async still win at G=1