submission 690605

  ⋯ 4 unchanged lines
  # leaderboard = "amd-mxfp4-mm"
   
  """
- isa_v573: Multi-attack breakthrough attempt.
- Attack A: Non-power-of-2 SplitK (SK=7 for shape 2, K=7168/7=1024 per split = exact division!)
- Attack B: Pre-quantized A cache (skip quant on repeated calls with same A data)
- Attack C: AITER deep probe (print all APIs to stderr)
-  
- KEY INSIGHT: v439 rounds SK to power-of-2 (line 290), forcing SK=7→8.
- SK=7 gives EXACT K-division: 7168/7=1024, with BK=512 → 2 iters per split.
- SK=7 × 33 N-tiles = 231 WGs — excellent CU utilization on 256 CUs.
+ v754c: Best config — preshuffle for M<=16 K<=1024 (shape 1: 6.30us) +
+ BSM=16 quant for shapes 5,6 (13.8/12.4us) + v690 fused for shapes 2-4.
+ Bench geomean: 8.86us (best ever). LB-safe (no caching).
  """
- import os, sys, subprocess
+ import os, sys
  os.environ["HIP_FORCE_DEV_KERNARG"] = "1"
   
  from task import input_t, output_t
  ⋯ 7 unchanged lines
  _FP4X2 = _dt.fp4x2
  _E8M0 = _dt.fp8_e8m0
  _cache = {}
- _a_quant_cache = {}  # Attack B: cache pre-quantized A
+ _e8m0_shuffle = None
   
   
- # ============ Attack C: AITER Deep Probe (runs at import time) ============
- def _aiter_probe():
-     P("\n=== AITER DEEP PROBE ===")
-     try:
-         # Check all top-level exports
-         all_attrs = [a for a in dir(aiter) if not a.startswith('_')]
-         gemm_attrs = [a for a in all_attrs if 'gemm' in a.lower() or 'quant' in a.lower() or 'fp4' in a.lower() or 'mxfp' in a.lower()]
-         P(f"GEMM/quant-related attrs: {gemm_attrs}")
-  
-         # Check for fused quant+gemm
-         for name in ['fused_quant_gemm', 'bf16_fp4_gemm', 'bf16_to_fp4_gemm', 'quant_gemm',
-                       'hk_gemm', 'gemm_bf16_fp4', 'mxfp4_gemm', 'fused_mxfp4_gemm',
-                       'gemm_a4w4_fused', 'gemm_fp4_fused']:
-             if hasattr(aiter, name):
-                 P(f"FOUND: aiter.{name} = {getattr(aiter, name)}")
-  
-         # Check gemm_a4w4_asm signature
-         if hasattr(aiter, 'gemm_a4w4_asm'):
-             import inspect
-             try:
-                 sig = inspect.signature(aiter.gemm_a4w4_asm)
-                 P(f"gemm_a4w4_asm signature: {sig}")
-             except: pass
-  
-         # Check for blockscale
-         if hasattr(aiter, 'gemm_a4w4_blockscale'):
-             import inspect
-             try:
-                 sig = inspect.signature(aiter.gemm_a4w4_blockscale)
-                 P(f"gemm_a4w4_blockscale signature: {sig}")
-             except: pass
-  
-         # Check aiter.ops namespace
-         if hasattr(aiter, 'ops'):
-             ops_attrs = [a for a in dir(aiter.ops) if 'gemm' in a.lower() or 'quant' in a.lower()]
-             P(f"aiter.ops gemm/quant attrs: {ops_attrs}")
-  
-         # Check for per_1x32_f4_quant (fast quant function)
-         if hasattr(aiter, 'per_1x32_f4_quant'):
-             import inspect
-             try:
-                 sig = inspect.signature(aiter.per_1x32_f4_quant)
-                 P(f"per_1x32_f4_quant signature: {sig}")
-             except: pass
-  
-         # Check new .co files
-         try:
-             result = subprocess.run(["find", "/home/runner/aiter/hsa", "-name", "*fp4*", "-o", "-name", "*quant*"],
-                                    capture_output=True, text=True, timeout=5)
-             if result.stdout.strip():
-                 P(f"FP4/quant .co files: {result.stdout.strip()[:500]}")
-         except: pass
-  
-         # Check git log for recent changes
-         try:
-             result = subprocess.run(["git", "-C", "/home/runner/aiter", "log", "--oneline", "-5"],
-                                    capture_output=True, text=True, timeout=5)
-             P(f"AITER recent commits: {result.stdout.strip()}")
-         except: pass
-  
-     except Exception as e:
-         P(f"AITER probe error: {e}")
-     P("=== END AITER PROBE ===\n")
-  
- _aiter_probe()
-  
-  
- # ============ Triton kernels (same as v439) ============
-  
  def _knl_name(tile_m, tile_n):
      base = f"f4gemm_bf16_per1x32Fp4_BpreShuffle_{tile_m}x{tile_n}"
      return f"_ZN5aiter{len(base)}{base}E"
  ⋯ 7 unchanged lines
      MBITS_FP4: tl.constexpr = 1
      EBITS_F32: tl.constexpr = 8
      EBITS_FP4: tl.constexpr = 2
+     max_normal: tl.constexpr = 6
+     min_normal: tl.constexpr = 1
      QUANT: tl.constexpr = 32
      NUM_QB: tl.constexpr = BLOCK_K // QUANT
      x = x.reshape(BLOCK_M, NUM_QB, QUANT)
  ⋯ 10 unchanged lines
      s = qx & 0x80000000
      qx = qx ^ s
      qx_fp32 = qx.to(tl.float32, bitcast=True)
-     saturate_mask = qx_fp32 >= 6
-     denormal_mask = (not saturate_mask) & (qx_fp32 < 1)
+     saturate_mask = qx_fp32 >= max_normal
+     denormal_mask = (not saturate_mask) & (qx_fp32 < min_normal)
      normal_mask = not (saturate_mask | denormal_mask)
      denorm_exp: tl.constexpr = (EXP_BIAS_FP32 - EXP_BIAS_FP4) + (MBITS_F32 - MBITS_FP4) + 1
      denorm_mask_int: tl.constexpr = denorm_exp << MBITS_F32
  ⋯ 25 unchanged lines
  @triton.jit
  def xcd_swizzle(pid, domain_size, XCD_SWIZZLE: tl.constexpr):
      pids_per_group = domain_size // XCD_SWIZZLE
-     extra = domain_size % XCD_SWIZZLE
+     extra_pid_groups = domain_size % XCD_SWIZZLE
      group = pid % XCD_SWIZZLE
      local_pid = pid // XCD_SWIZZLE
-     return group * pids_per_group + tl.minimum(group, extra) + local_pid
+     new_pid = group * pids_per_group + tl.minimum(group, extra_pid_groups) + local_pid
+     return new_pid
   
   
+ # ============ K<=1024: FUSED (same as v127) ============
  @triton.jit
- def _fused_quant_gemm_kernel(A_ptr, Bq_ptr, Bscale_sh_ptr, C_ptr, M, N, K: tl.constexpr, stride_a_m, stride_a_k, stride_bq_n, stride_bq_k, SN_DIV8_MUL256: tl.constexpr, stride_c_m, stride_c_n, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr):
-     pid_m = tl.program_id(0); pid_n = tl.program_id(1)
-     offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M); offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
-     acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32); QUANT: tl.constexpr = 32; NSK: tl.constexpr = BLOCK_K // QUANT
+ def _fused_quant_gemm_kernel(
+     A_ptr, Bq_ptr, Bscale_sh_ptr, C_ptr,
+     M, N, K: tl.constexpr,
+     stride_a_m, stride_a_k, stride_bq_n, stride_bq_k,
+     SN_DIV8_MUL256: tl.constexpr,
+     stride_c_m, stride_c_n,
+     BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,
+ ):
+     pid_m = tl.program_id(0)
+     pid_n = tl.program_id(1)
+     offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
+     offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
+     acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
+     QUANT: tl.constexpr = 32
+     NSK: tl.constexpr = BLOCK_K // QUANT
+  
      for ki in tl.range(0, K, BLOCK_K):
-         a_offs_k = ki + tl.arange(0, BLOCK_K); a_ptrs = A_ptr + offs_m[:, None] * stride_a_m + a_offs_k[None, :] * stride_a_k
+         a_offs_k = ki + tl.arange(0, BLOCK_K)
+         a_ptrs = A_ptr + offs_m[:, None] * stride_a_m + a_offs_k[None, :] * stride_a_k
          a_mask = (offs_m < M)[:, None] & (a_offs_k < K)[None, :]
          a_tile = tl.load(a_ptrs, mask=a_mask, other=0.0).to(tl.float32)
          a_fp4, a_scale = _mxfp4_quant_op(a_tile, BLOCK_K, BLOCK_M)
  ⋯ 1 unchanged lines
          b_ptrs = Bq_ptr + offs_n[None, :] * stride_bq_n + b_offs_k[:, None] * stride_bq_k
          b_mask = (offs_n < N)[None, :] & (b_offs_k < K // 2)[:, None]
          b_tile = tl.load(b_ptrs, mask=b_mask, other=0)
-         bs_row = offs_n; bs_col_base = ki // QUANT; bs_col_offs = tl.arange(0, NSK)
-         row = bs_row[:, None]; col = (bs_col_base + bs_col_offs)[None, :]
+         bs_row = offs_n
+         bs_col_base = ki // QUANT
+         bs_col_offs = tl.arange(0, NSK)
+         row = bs_row[:, None]
+         col = (bs_col_base + bs_col_offs)[None, :]
          shuf_idx = (row // 32) * SN_DIV8_MUL256 + (col // 8) * 256 + (col % 4) * 64 + (row % 16) * 4 + ((col % 8) // 4) * 2 + ((row % 32) // 16)
          bs_mask = (offs_n[:, None] < N) & (bs_col_offs[None, :] < (K // QUANT - bs_col_base))
          b_scale = tl.load(Bscale_sh_ptr + shuf_idx, mask=bs_mask, other=0)
          acc = tl.dot_scaled(a_fp4, a_scale, "e2m1", b_tile, b_scale, "e2m1", acc, fast_math=True)
      c_ptrs = C_ptr + offs_m[:, None] * stride_c_m + offs_n[None, :] * stride_c_n
-     c_mask = (offs_m < M)[:, None] & (offs_n < N)[None, :]; tl.store(c_ptrs, acc.to(tl.bfloat16), mask=c_mask)
+     c_mask = (offs_m < M)[:, None] & (offs_n < N)[None, :]
+     tl.store(c_ptrs, acc.to(tl.bfloat16), mask=c_mask)
   
   
+ # ============ FUSED SPLITK FOR K>1024, M<=32 ============
  @triton.jit
- def _fused_splitk_gemm(A_ptr, Bq_ptr, Bscale_sh_ptr, Y_ptr, M, N, K: tl.constexpr, stride_a_m, stride_a_k, stride_bq_n, stride_bq_k, SN_DIV8_MUL256: tl.constexpr, stride_y_k, stride_y_m, stride_y_n, grid_m, grid_n, BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr, SPLIT_K: tl.constexpr, XCD_SWIZZLE: tl.constexpr):
+ def _fused_splitk_gemm(
+     A_ptr, Bq_ptr, Bscale_sh_ptr, Y_ptr,
+     M, N, K: tl.constexpr,
+     stride_a_m, stride_a_k, stride_bq_n, stride_bq_k,
+     SN_DIV8_MUL256: tl.constexpr,
+     stride_y_k, stride_y_m, stride_y_n,
+     grid_m, grid_n,
+     BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,
+     SPLIT_K: tl.constexpr, XCD_SWIZZLE: tl.constexpr,
+     matrix_instr_nonkdim: tl.constexpr = 32,  # v752: configurable MFMA
+ ):
      pid = tl.program_id(0)
-     if XCD_SWIZZLE > 1: pid = xcd_swizzle(pid, grid_m * grid_n * SPLIT_K, XCD_SWIZZLE)
-     pid_k = pid % SPLIT_K; pid_mn = pid // SPLIT_K; pid_m = pid_mn // grid_n; pid_n = pid_mn % grid_n
-     offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M); offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
-     acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32); QUANT: tl.constexpr = 32; NSK: tl.constexpr = BLOCK_K // QUANT
-     k_per_split = ((K + SPLIT_K - 1) // SPLIT_K + BLOCK_K - 1) // BLOCK_K * BLOCK_K
-     k_start = pid_k * k_per_split; k_end = min(k_start + k_per_split, K)
+     total_tiles = grid_m * grid_n * SPLIT_K
+     if XCD_SWIZZLE > 1:
+         pid = xcd_swizzle(pid, total_tiles, XCD_SWIZZLE)
+     pid_k = pid % SPLIT_K
+     pid_mn = pid // SPLIT_K
+     pid_m = pid_mn // grid_n
+     pid_n = pid_mn % grid_n
+  
+     offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
+     offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
+     acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
+     QUANT: tl.constexpr = 32
+     NSK: tl.constexpr = BLOCK_K // QUANT
+  
+     k_per_split = (K + SPLIT_K - 1) // SPLIT_K
+     k_per_split = ((k_per_split + BLOCK_K - 1) // BLOCK_K) * BLOCK_K
+     k_start = pid_k * k_per_split
+     k_end = min(k_start + k_per_split, K)
+  
      for ki in tl.range(k_start, k_end, BLOCK_K):
-         a_offs_k = ki + tl.arange(0, BLOCK_K); a_ptrs = A_ptr + offs_m[:, None] * stride_a_m + a_offs_k[None, :] * stride_a_k
+         a_offs_k = ki + tl.arange(0, BLOCK_K)
+         a_ptrs = A_ptr + offs_m[:, None] * stride_a_m + a_offs_k[None, :] * stride_a_k
          a_mask = (offs_m < M)[:, None] & (a_offs_k < K)[None, :]
          a_tile = tl.load(a_ptrs, mask=a_mask, other=0.0).to(tl.float32)
          a_fp4, a_scale = _mxfp4_quant_op(a_tile, BLOCK_K, BLOCK_M)
  ⋯ 1 unchanged lines
          b_ptrs = Bq_ptr + offs_n[None, :] * stride_bq_n + b_offs_k[:, None] * stride_bq_k
          b_mask = (offs_n < N)[None, :] & (b_offs_k < K // 2)[:, None]
          b_tile = tl.load(b_ptrs, mask=b_mask, other=0)
-         bs_row = offs_n; bs_col_base = ki // QUANT; bs_col_offs = tl.arange(0, NSK)
-         row = bs_row[:, None]; col = (bs_col_base + bs_col_offs)[None, :]
+         bs_row = offs_n
+         bs_col_base = ki // QUANT
+         bs_col_offs = tl.arange(0, NSK)
+         row = bs_row[:, None]
+         col = (bs_col_base + bs_col_offs)[None, :]
          shuf_idx = (row // 32) * SN_DIV8_MUL256 + (col // 8) * 256 + (col % 4) * 64 + (row % 16) * 4 + ((col % 8) // 4) * 2 + ((row % 32) // 16)
          bs_mask = (offs_n[:, None] < N) & (bs_col_offs[None, :] < (K // QUANT - bs_col_base))
          b_scale = tl.load(Bscale_sh_ptr + shuf_idx, mask=bs_mask, other=0)
          acc = tl.dot_scaled(a_fp4, a_scale, "e2m1", b_tile, b_scale, "e2m1", acc, fast_math=True)
+  
      y_ptrs = Y_ptr + pid_k * stride_y_k + offs_m[:, None] * stride_y_m + offs_n[None, :] * stride_y_n
      y_mask = (offs_m < M)[:, None] & (offs_n < N)[None, :]
-     if SPLIT_K > 1: tl.store(y_ptrs, acc, mask=y_mask)
-     else: tl.store(y_ptrs, acc.to(tl.bfloat16), mask=y_mask)
+     if SPLIT_K > 1:
+         tl.store(y_ptrs, acc, mask=y_mask)  # FP32 partials
+     else:
+         tl.store(y_ptrs, acc.to(tl.bfloat16), mask=y_mask)
   
   
  @triton.jit
- def _reduce_splitk(Y_ptr, Out_ptr, M, N, stride_y_k, stride_y_m, stride_y_n, stride_o_m, stride_o_n, SPLIT_K: tl.constexpr, BLOCK_N: tl.constexpr):
-     pid_m = tl.program_id(0); pid_n = tl.program_id(1)
-     offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N); n_mask = offs_n < N
+ def _reduce_splitk(
+     Y_ptr, Out_ptr, M, N,
+     stride_y_k, stride_y_m, stride_y_n,
+     stride_o_m, stride_o_n,
+     SPLIT_K: tl.constexpr, BLOCK_N: tl.constexpr,
+ ):
+     pid_m = tl.program_id(0)
+     pid_n = tl.program_id(1)
+     offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
+     n_mask = offs_n < N
      acc = tl.zeros([BLOCK_N], dtype=tl.float32)
      for k in tl.range(0, SPLIT_K):
-         acc += tl.load(Y_ptr + k * stride_y_k + pid_m * stride_y_m + offs_n * stride_y_n, mask=n_mask, other=0.0).to(tl.float32)
-     tl.store(Out_ptr + pid_m * stride_o_m + offs_n * stride_o_n, acc.to(tl.bfloat16), mask=n_mask)
+         vals = tl.load(Y_ptr + k * stride_y_k + pid_m * stride_y_m + offs_n * stride_y_n,
+                        mask=n_mask, other=0.0)
+         acc += vals.to(tl.float32)
+     out_ptrs = Out_ptr + pid_m * stride_o_m + offs_n * stride_o_n
+     tl.store(out_ptrs, acc.to(tl.bfloat16), mask=n_mask)
   
   
+ # ============ v758: PER-ELEMENT PARALLEL QUANT (GodZmk-style) ============
  @triton.jit
- def _fused_quant_shuffle_kernel(x_ptr, x_fp4_ptr, bs_shuf_ptr, stride_x_m, stride_x_n, stride_fp4_m, stride_fp4_n, M, N, SN_DIV8_MUL256, SCALE_COLS, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, NUM_ITER: tl.constexpr, NUM_STAGES: tl.constexpr):
-     pid_m = tl.program_id(0); start_n = tl.program_id(1) * NUM_ITER
-     QUANT: tl.constexpr = 32; NUM_QB: tl.constexpr = BLOCK_SIZE_N // QUANT
+ def _per_element_quant_shuffle(
+     x_ptr, out_ptr, scale_ptr,
+     stride_xm, M, K: tl.constexpr, SN_DIV8_MUL256,
+     GROUP_SIZE: tl.constexpr = 32,
+ ):
+     row = tl.program_id(0)
+     group_id = tl.program_id(1)
+     k_start = group_id * GROUP_SIZE
+     half = tl.arange(0, GROUP_SIZE // 2)
+     k_even = k_start + half * 2
+     k_odd = k_start + half * 2 + 1
+     x_even = tl.load(x_ptr + row * stride_xm + k_even, mask=k_even < K, other=0.0).to(tl.float32)
+     x_odd = tl.load(x_ptr + row * stride_xm + k_odd, mask=k_odd < K, other=0.0).to(tl.float32)
+     # E8M0 scale: amax of 32 elements
+     abs_max = tl.maximum(tl.max(tl.abs(x_even), axis=0), tl.max(tl.abs(x_odd), axis=0))
+     abs_max = tl.maximum(abs_max, 1e-38).to(tl.float32)
+     abs_max_int = abs_max.to(tl.int32, bitcast=True)
+     abs_max_rounded = ((abs_max_int + 0x200000).to(tl.uint32, bitcast=True) & 0xFF800000).to(tl.float32, bitcast=True)
+     scale_unb = tl.floor(tl.math.log2(abs_max_rounded)).to(tl.int32) - 2
+     scale_unb = tl.minimum(tl.maximum(scale_unb, -127), 127)
+     e8m0_exp = (scale_unb + 127).to(tl.uint8)
+     quant_scale = tl.math.exp2(-scale_unb.to(tl.float32))
+     # Quantize even elements → lo nibbles
+     xs_e = x_even * quant_scale
+     xs_e_uint = xs_e.to(tl.int32, bitcast=True).to(tl.uint32)
+     s_e = xs_e_uint & 0x80000000
+     xs_e_pos_uint = xs_e_uint ^ s_e
+     xs_e_pos = xs_e_pos_uint.to(tl.float32, bitcast=True)
+     sat_e = xs_e_pos >= 6.0
+     den_e = xs_e_pos < 1.0
+     mant_odd_e = (xs_e_pos_uint >> 22) & 1
+     norm_e = ((xs_e_pos_uint.to(tl.int32) + (-1054867457)) + mant_odd_e.to(tl.int32)) >> 22
+     norm_e = norm_e.to(tl.uint8)
+     den_val_e = (xs_e_pos + 4194304.0).to(tl.int32, bitcast=True) - 0x4A800000
+     den_val_e = den_val_e.to(tl.uint8)
+     q_e = tl.full(xs_e.shape, 7, dtype=tl.uint8)
+     q_e = tl.where(~sat_e, norm_e, q_e)
+     q_e = tl.where(den_e, den_val_e, q_e)
+     sign_e = (s_e >> 28).to(tl.uint8)
+     lo = (q_e | sign_e) & 0xF
+     # Quantize odd elements → hi nibbles
+     xs_o = x_odd * quant_scale
+     xs_o_uint = xs_o.to(tl.int32, bitcast=True).to(tl.uint32)
+     s_o = xs_o_uint & 0x80000000
+     xs_o_pos_uint = xs_o_uint ^ s_o
+     xs_o_pos = xs_o_pos_uint.to(tl.float32, bitcast=True)
+     sat_o = xs_o_pos >= 6.0
+     den_o = xs_o_pos < 1.0
+     mant_odd_o = (xs_o_pos_uint >> 22) & 1
+     norm_o = ((xs_o_pos_uint.to(tl.int32) + (-1054867457)) + mant_odd_o.to(tl.int32)) >> 22
+     norm_o = norm_o.to(tl.uint8)
+     den_val_o = (xs_o_pos + 4194304.0).to(tl.int32, bitcast=True) - 0x4A800000
+     den_val_o = den_val_o.to(tl.uint8)
+     q_o = tl.full(xs_o.shape, 7, dtype=tl.uint8)
+     q_o = tl.where(~sat_o, norm_o, q_o)
+     q_o = tl.where(den_o, den_val_o, q_o)
+     sign_o = (s_o >> 28).to(tl.uint8)
+     hi = ((q_o | sign_o) & 0xF) << 4
+     # Pack and store FP4
+     packed = lo | hi
+     tl.store(out_ptr + row * (K // 2) + k_start // 2 + half, packed.to(tl.uint8), mask=half < (K // 2 - k_start // 2))
+     # Store CK-shuffled scale
+     sc = group_id
+     shuf_idx = (row // 32) * SN_DIV8_MUL256 + (sc // 8) * 256 + (sc % 4) * 64 + (row % 16) * 4 + ((sc % 8) // 4) * 2 + ((row % 32) // 16)
+     tl.store(scale_ptr + shuf_idx, e8m0_exp)
+  
+  
+ # ============ QUANT KERNEL FOR CK ASM PATH ============
+ @triton.jit
+ def _fused_quant_shuffle_kernel(
+     x_ptr, x_fp4_ptr, bs_shuf_ptr,
+     stride_x_m, stride_x_n, stride_fp4_m, stride_fp4_n,
+     M, N, SN_DIV8_MUL256, SCALE_COLS,
+     BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr,
+     NUM_ITER: tl.constexpr, NUM_STAGES: tl.constexpr,
+ ):
+     pid_m = tl.program_id(0)
+     start_n = tl.program_id(1) * NUM_ITER
+     QUANT: tl.constexpr = 32
+     NUM_QB: tl.constexpr = BLOCK_SIZE_N // QUANT
      for pid_n in tl.range(start_n, min(start_n + NUM_ITER, N), num_stages=NUM_STAGES):
-         x_offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M); x_offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+         x_offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
+         x_offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
+         x_offs = x_offs_m[:, None] * stride_x_m + x_offs_n[None, :] * stride_x_n
          x_mask = (x_offs_m < M)[:, None] & (x_offs_n < N)[None, :]
-         x = tl.load(x_ptr + x_offs_m[:, None] * stride_x_m + x_offs_n[None, :] * stride_x_n, mask=x_mask, other=0.0, cache_modifier=".cg").to(tl.float32)
+         x = tl.load(x_ptr + x_offs, mask=x_mask, other=0.0, cache_modifier=".cg").to(tl.float32)
          out_tensor, bs_e8m0 = _mxfp4_quant_op(x, BLOCK_SIZE_N, BLOCK_SIZE_M)
+         out_offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
          out_offs_n = pid_n * BLOCK_SIZE_N // 2 + tl.arange(0, BLOCK_SIZE_N // 2)
-         out_mask = (x_offs_m < M)[:, None] & (out_offs_n < (N // 2))[None, :]
-         tl.store(x_fp4_ptr + x_offs_m[:, None] * stride_fp4_m + out_offs_n[None, :] * stride_fp4_n, out_tensor, mask=out_mask)
-         bs_col = pid_n * NUM_QB + tl.arange(0, NUM_QB); row = x_offs_m[:, None]; col = bs_col[None, :]
+         out_offs = out_offs_m[:, None] * stride_fp4_m + out_offs_n[None, :] * stride_fp4_n
+         out_mask = (out_offs_m < M)[:, None] & (out_offs_n < (N // 2))[None, :]
+         tl.store(x_fp4_ptr + out_offs, out_tensor, mask=out_mask)
+         bs_row = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)
+         bs_col = pid_n * NUM_QB + tl.arange(0, NUM_QB)
+         row = bs_row[:, None]
+         col = bs_col[None, :]
          shuf_idx = (row // 32) * SN_DIV8_MUL256 + (col // 8) * 256 + (col % 4) * 64 + (row % 16) * 4 + ((col % 8) // 4) * 2 + ((row % 32) // 16)
-         tl.store(bs_shuf_ptr + shuf_idx, bs_e8m0, mask=(x_offs_m[:, None] < M) & (bs_col[None, :] < SCALE_COLS))
+         bs_mask = (bs_row[:, None] < M) & (bs_col[None, :] < SCALE_COLS)
+         tl.store(bs_shuf_ptr + shuf_idx, bs_e8m0, mask=bs_mask)
   
   
+ # ============ v728: INLINED AITER KERNEL (with OUR quant) for M<=4 ============
+ @triton.jit
+ def _our_quant_4arg(x, BLOCK_K: tl.constexpr, BLOCK_M: tl.constexpr, SGS: tl.constexpr):
+     """Our quant with 4-arg interface matching AITER's."""
+     return _mxfp4_quant_op(x, BLOCK_K, BLOCK_M)
+  
+ @triton.heuristics({
+     "EVEN_K": lambda args: (args["K"] % (args["BLOCK_SIZE_K"] // 2) == 0) and (args["SPLITK_BLOCK_SIZE"] % args["BLOCK_SIZE_K"] == 0) and (args["K"] % (args["SPLITK_BLOCK_SIZE"] // 2) == 0),
+     "GRID_MN": lambda args: triton.cdiv(args["M"], args["BLOCK_SIZE_M"]) * triton.cdiv(args["N"], args["BLOCK_SIZE_N"]),
+ })
+ @triton.jit
+ def _inlined_kernel(
+     a_ptr, b_ptr, c_ptr, b_scales_ptr, M, N, K,
+     stride_am, stride_ak, stride_bk, stride_bn,
+     stride_ck, stride_cm, stride_cn, stride_bsn, stride_bsk,
+     BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,
+     GROUP_SIZE_M: tl.constexpr, NUM_KSPLIT: tl.constexpr, SPLITK_BLOCK_SIZE: tl.constexpr,
+     EVEN_K: tl.constexpr, matrix_instr_nonkdim: tl.constexpr, GRID_MN: tl.constexpr,
+     ATOMIC_ADD: tl.constexpr, cache_modifier: tl.constexpr,
+ ):
+     tl.assume(stride_am > 0); tl.assume(stride_ak > 0); tl.assume(stride_bk > 0); tl.assume(stride_bn > 0)
+     tl.assume(stride_cm > 0); tl.assume(stride_cn > 0); tl.assume(stride_bsk > 0); tl.assume(stride_bsn > 0)
+     pid_unified = tl.program_id(axis=0); pid_k = pid_unified % NUM_KSPLIT; pid = pid_unified // NUM_KSPLIT
+     num_pid_m = tl.cdiv(M, BLOCK_SIZE_M); num_pid_n = tl.cdiv(N, BLOCK_SIZE_N)
+     if NUM_KSPLIT == 1:
+         num_pid_in_group = GROUP_SIZE_M * num_pid_n; group_id = pid // num_pid_in_group
+         first_pid_m = group_id * GROUP_SIZE_M; group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
+         pid_m = first_pid_m + ((pid % num_pid_in_group) % group_size_m); pid_n = (pid % num_pid_in_group) // group_size_m
+     else:
+         pid_m = pid // num_pid_n; pid_n = pid % num_pid_n
+     SCALE_GROUP_SIZE: tl.constexpr = 32
+     if (pid_k * SPLITK_BLOCK_SIZE // 2) < K:
+         num_k_iter = tl.cdiv(SPLITK_BLOCK_SIZE // 2, BLOCK_SIZE_K // 2)
+         offs_k_bf16 = tl.arange(0, BLOCK_SIZE_K); offs_k_split_bf16 = pid_k * SPLITK_BLOCK_SIZE + offs_k_bf16
+         offs_am = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % M
+         a_ptrs = a_ptr + offs_am[:, None] * stride_am + offs_k_split_bf16[None, :] * stride_ak
+         offs_k = tl.arange(0, BLOCK_SIZE_K // 2); offs_k_split = pid_k * (SPLITK_BLOCK_SIZE // 2) + offs_k
+         offs_bn = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % N
+         b_ptrs = b_ptr + offs_k_split[:, None] * stride_bk + offs_bn[None, :] * stride_bn
+         offs_ks = (pid_k * (SPLITK_BLOCK_SIZE // SCALE_GROUP_SIZE)) + tl.arange(0, BLOCK_SIZE_K // SCALE_GROUP_SIZE)
+         b_scale_ptrs = b_scales_ptr + offs_bn[:, None] * stride_bsn + offs_ks[None, :] * stride_bsk
+         accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
+         for k in range(pid_k * num_k_iter, (pid_k + 1) * num_k_iter):
+             b_scales = tl.load(b_scale_ptrs)
+             if EVEN_K:
+                 a_bf16 = tl.load(a_ptrs); b = tl.load(b_ptrs, cache_modifier=cache_modifier)
+             else:
+                 a_bf16 = tl.load(a_ptrs, mask=offs_k_bf16[None, :] < 2 * K - k * BLOCK_SIZE_K, other=0)
+                 b = tl.load(b_ptrs, mask=offs_k[:, None] < K - k * (BLOCK_SIZE_K // 2), other=0, cache_modifier=cache_modifier)
+             a, a_scales = _our_quant_4arg(a_bf16, BLOCK_SIZE_K, BLOCK_SIZE_M, 32)
+             accumulator += tl.dot_scaled(a, a_scales, "e2m1", b, b_scales, "e2m1")
+             a_ptrs += BLOCK_SIZE_K * stride_ak; b_ptrs += (BLOCK_SIZE_K // 2) * stride_bk
+             b_scale_ptrs += (BLOCK_SIZE_K // SCALE_GROUP_SIZE) * stride_bsk
+         c = accumulator.to(c_ptr.type.element_ty)
+         offs_cm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M).to(tl.int64)
+         offs_cn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N).to(tl.int64)
+         c_ptrs = c_ptr + stride_cm * offs_cm[:, None] + stride_cn * offs_cn[None, :] + pid_k * stride_ck
+         c_mask = (offs_cm[:, None] < M) & (offs_cn[None, :] < N)
+         tl.store(c_ptrs, c, mask=c_mask)
+  
+ @triton.jit
+ def _reduce_kernel(y_pp, y, M, N, syk, sym, syn, som, son, BM: tl.constexpr, BN: tl.constexpr, ACTUAL_SK: tl.constexpr, MAX_SK: tl.constexpr):
+     pm = tl.program_id(0); pn = tl.program_id(1)
+     offs_m = pm * BM + tl.arange(0, BM); offs_n = pn * BN + tl.arange(0, BN)
+     acc = tl.zeros((BM, BN), dtype=tl.float32)
+     for k in range(MAX_SK):
+         if k < ACTUAL_SK:
+             vals = tl.load(y_pp + k * syk + offs_m[:, None] * sym + offs_n[None, :] * syn,
+                           mask=(offs_m[:, None] < M) & (offs_n[None, :] < N), other=0.0)
+             acc += vals.to(tl.float32)
+     mask = (offs_m[:, None] < M) & (offs_n[None, :] < N)
+     tl.store(y + offs_m[:, None] * som + offs_n[None, :] * son, acc.to(tl.bfloat16), mask=mask)
+  
+ _inline_scale_cache = {}
+ _inline_bad = set()
+  
+ def _unshuffle_scales(B_scale_sh, n, k):
+     # Cache index tensor only (deterministic). Recompute gather every call (LB-safe).
+     idx_key = (n, k, 'idx')
+     if idx_key not in _inline_scale_cache:
+         sc = (k + 31) // 32; sn = ((sc + 7) // 8) * 8; s = (sn // 8) * 256
+         r = torch.arange(n, device=B_scale_sh.device, dtype=torch.int64).unsqueeze(1)
+         c = torch.arange(sc, device=B_scale_sh.device, dtype=torch.int64).unsqueeze(0)
+         idx = (r // 32) * s + (c // 8) * 256 + (c % 4) * 64 + (r % 16) * 4 + ((c % 8) // 4) * 2 + ((r % 32) // 16)
+         _inline_scale_cache[idx_key] = idx
+     idx = _inline_scale_cache[idx_key]
+     f = B_scale_sh.view(torch.uint8).flatten()
+     if idx.max() >= f.shape[0]: return None
+     return f[idx]  # Recompute every call (GPU-side indexed gather, ~0.1μs)
+  
+  
  def _init(m, k, n, device):
      QUANT = 32
      scale_cols = (k + QUANT - 1) // QUANT
      sn = ((scale_cols + 7) // 8) * 8
      sn_div8_mul256 = (sn // 8) * 256
   
+     CU_COUNT = 256  # MI355X has 256 CUs
+  
      if k <= 1024:
-         BLOCK_K = max(128, triton.next_power_of_2(k))
+         # Path A: Fused kernel for small K
+         BLOCK_K = max(128, triton.next_power_of_2(k))  # min 128 for dot_scaled
          BLOCK_M = 16 if m <= 32 else max(16, min(32, triton.next_power_of_2(m)))
          BLOCK_N = 64
          NW = 4
          grid = (triton.cdiv(m, BLOCK_M), triton.cdiv(n, BLOCK_N))
+         total_wgs = grid[0] * grid[1]
+         wpe = 2 if total_wgs > CU_COUNT else 1
          out = torch.empty(m, n, dtype=torch.bfloat16, device=device)
          return {
              'mode': 'fused', 'out': out,
              'grid': grid, 'BM': BLOCK_M, 'BN': BLOCK_N, 'BK': BLOCK_K,
-             'sn_div8_mul256': sn_div8_mul256, 'NW': NW,
+             'sn_div8_mul256': sn_div8_mul256, 'NW': NW, 'wpe': wpe,
          }
      elif m <= 32:
-         # ===== ATTACK A: Non-power-of-2 SplitK =====
-         # K=7168: try SK=7 (7168/7=1024 per split, exact division!)
+         # Path B: Fused SplitK for small-M K>1024 (shape 2)
          BLOCK_K = 512
          BLOCK_N = 64
          BLOCK_M = 16 if m <= 16 else 32
+  
          m_tiles = triton.cdiv(m, BLOCK_M)
          n_tiles = triton.cdiv(n, BLOCK_N)
          total_mn = m_tiles * n_tiles
   
-         # Smart SK selection: prefer exact K-division
-         k_iters_512 = triton.cdiv(k, 512)
-         # Try non-power-of-2 SK values that divide K evenly
-         best_sk = 8  # default
-         if k == 7168:
-             # K=7168 = 7×1024 = 14×512
-             # SK=7: 7168/7=1024 per split, 1024/512=2 iters per split
-             # SK=14: 7168/14=512 per split, 512/512=1 iter per split
-             best_sk = 14  # 1 iter per split = minimal quant overhead!
-         elif k % 7 == 0:
-             best_sk = 7
-         elif k % 14 == 0:
-             best_sk = 14
-         else:
-             # Fallback: power-of-2 logic from v439
-             best_sk = max(1, min(16, 256 // max(1, total_mn)))
-             while best_sk > 1 and k_iters_512 < best_sk * 2:
-                 best_sk //= 2
-             best_sk = 1 << (best_sk - 1).bit_length() if best_sk > 1 else 1
+         # Exact-division SplitK: use k_iters for 1 iter per split (zero waste)
+         # v573 bug: power-of-2 rounding gave SK=8 for K=7168 → 12.9μs
+         # isa_v573 proved SK=14 → 11.1μs (exact: 7168/512=14, 462 WGs, 1.8/CU)
+         k_iters = k // BLOCK_K if k % BLOCK_K == 0 else triton.cdiv(k, BLOCK_K)
+         # Use k_iters directly — oversubscription (1-2 WGs/CU) is fine
+         SPLIT_K = min(k_iters, 16)  # cap at 16 to limit reduce overhead
   
-         SPLIT_K = best_sk
          total_wgs = m_tiles * n_tiles * SPLIT_K
          XCD_SWIZZLE = 8 if total_wgs >= 16 else 1
+         wpe = 2 if total_wgs > CU_COUNT else 1
          out = torch.empty(m, n, dtype=torch.bfloat16, device=device)
   
-         P(f"Shape ({m},{k},{n}): SK={SPLIT_K}, BK={BLOCK_K}, total_wgs={total_wgs}, "
-           f"k_per_split={k//SPLIT_K}, iters_per_split={triton.cdiv(k//SPLIT_K, BLOCK_K)}")
-  
          if SPLIT_K > 1:
              scratch = torch.empty(SPLIT_K, m, n, dtype=torch.float32, device=device)
              reduce_grid = (m, triton.cdiv(n, 128))
  ⋯ 4 unchanged lines
          return {
              'mode': 'splitk', 'out': out, 'scratch': scratch,
              'BM': BLOCK_M, 'BN': BLOCK_N, 'BK': BLOCK_K,
-             'SPLIT_K': SPLIT_K, 'XCD_SWIZZLE': XCD_SWIZZLE,
+             'SPLIT_K': SPLIT_K, 'XCD_SWIZZLE': XCD_SWIZZLE, 'wpe': wpe,
              'grid_m': m_tiles, 'grid_n': n_tiles, 'total_wgs': total_wgs,
              'sn_div8_mul256': sn_div8_mul256, 'reduce_grid': reduce_grid,
          }
      else:
+         # Path C: CK ASM for large-M K>1024 (shapes 5, 6)
          sm = ((m + 255) // 256) * 256
          x_fp4 = torch.empty((m, k // 2), dtype=torch.uint8, device=device)
          bs_shuffled = torch.empty(sm, sn, dtype=torch.uint8, device=device)
          out = torch.empty(m, n, dtype=torch.bfloat16, device=device)
+  
+         # v754c: BSM=16 NW=4 for shapes 5,6 quant (proven optimal)
          if m <= 64:
-             BSM = triton.next_power_of_2(m)
+             BSM = 16
              NUM_ITER, BSN, NW, NS = 1, 128, 4, 1
          else:
-             NUM_ITER, BSM, BSN, NW, NS = 2, 32, 64, 4, 2
+             NUM_ITER, BSM, BSN, NW, NS = 2, 16, 64, 4, 2
+  
          grid = (triton.cdiv(m, BSM), triton.cdiv(k, BSN * NUM_ITER))
          knl = _knl_name(32, 128)
-         l2ks = None
          gemm_wgs = triton.cdiv(m, 32) * triton.cdiv(n, 128)
          if gemm_wgs < 32:
              l2ks = 3
          elif gemm_wgs < 64:
              l2ks = 2
+         elif gemm_wgs < CU_COUNT:
+             l2ks = 1
+         else:
+             l2ks = None
+  
+         # Dynamic waves_per_eu for quant kernel
+         quant_wgs = grid[0] * grid[1]
+         quant_wpe = 2 if quant_wgs > CU_COUNT else 0
+  
          return {
              'mode': 'asm',
              'x_fp4': x_fp4, 'bs_shuffled': bs_shuffled, 'out': out,
              'sc': scale_cols, 'sn_div8_mul256': sn_div8_mul256,
              'grid': grid, 'BSM': BSM, 'BSN': BSN,
              'NW': NW, 'NS': NS, 'NI': NUM_ITER,
-             'knl': knl, 'l2ks': l2ks,
+             'knl': knl, 'l2ks': l2ks, 'quant_wpe': quant_wpe,
          }
   
   
+ _preshuffle_fn = None
+ _preshuffle_bad = set()
+  
  def custom_kernel(data: input_t) -> output_t:
      A, B, B_q, B_shuffle, B_scale_sh = data
      m, k = A.shape
      n = B_q.shape[0]
+  
+     # v754c: Preshuffle ONLY for M<=16 K<=1024 (proven: shape 1 at 6.30us)
+     # Shapes 5,6: preshuffle is 2-3x slower than CK ASM (v755 confirmed 32/34us vs 14/12us)
+     global _preshuffle_fn
+     # v754c: Preshuffle for M<=16 K<=1024 only (shape 1 proven at 6.30us)
+     if m <= 16 and k <= 1024 and (m, k, n) not in _preshuffle_bad:
+         try:
+             if _preshuffle_fn is None:
+                 from aiter.ops.triton.gemm.basic.gemm_a16wfp4 import gemm_a16wfp4_preshuffle
+                 _preshuffle_fn = gemm_a16wfp4_preshuffle
+  
+             sc = (k + 31) // 32
+             sn = ((sc + 7) // 8) * 8
+             padN = B_scale_sh.view(torch.uint8).shape[0]
+             bs_reshaped = B_scale_sh.view(torch.uint8).reshape(padN // 32, sn * 32)
+             K_half = k // 2
+             b_shuf_reshaped = B_shuffle.view(torch.uint8).reshape(n // 16, K_half * 16)
+  
+             preshuffle_config = {
+                 'BLOCK_SIZE_M': 8,
+                 'BLOCK_SIZE_N': 128,
+                 'BLOCK_SIZE_K': 256,
+                 'GROUP_SIZE_M': 1,
+                 'NUM_KSPLIT': 1,
+                 'SPLITK_BLOCK_SIZE': k,
+                 'matrix_instr_nonkdim': 16,
+                 'num_warps': 8,
+                 'num_stages': 2,
+                 'waves_per_eu': 2,
+                 'cache_modifier': '.cg',
+             }
+             out = _preshuffle_fn(
+                 A, b_shuf_reshaped, bs_reshaped,
+                 prequant=True, dtype=torch.bfloat16,
+                 config=preshuffle_config,
+             )
+             return out
+         except Exception as e:
+             P(f"PRESHUFFLE FAIL ({m},{n},{k}): {e}")
+             _preshuffle_bad.add((m, k, n))
+  
+     # v751: Inlined kernel DISABLED — both transpose and strided access are slow
+     # The inlined kernel REQUIRES (K/2,N) contiguous layout for coalesced loads
+     # Cannot avoid the transpose, and transpose is 70-80μs with cold L2
+     if False and m <= 16 and (m, k, n) not in _inline_bad:
+         try:
+             scales = _unshuffle_scales(B_scale_sh, n, k)
+             if scales is not None:
+                 K_half = k // 2
+                 BK = 512
+                 NS = 1
+                 # SplitK for large K
+                 k_iters = K_half // (BK // 2) if K_half % (BK // 2) == 0 else triton.cdiv(K_half, BK // 2)
+                 SK = min(k_iters, 16) if k > 1024 else 1
+                 SPK_BS = triton.cdiv(K_half, SK) * 2 if SK > 1 else 2 * K_half
+                 # Align SPK_BS to BK
+                 if SK > 1:
+                     SPK_BS = triton.cdiv(SPK_BS // 2, BK // 2) * (BK // 2) * 2
+                 # Per-M config
+                 if m <= 4:
+                     BM = 4; BN = 128; NW = 4; NKDIM = 16
+                 elif m <= 8:
+                     BM = 8; BN = 128; NW = 8; NKDIM = 16
+                 elif m <= 16:
+                     BM = 16; BN = 128; NW = 4; NKDIM = 16
+                 else:
+                     # M=32: BM=32 BN=64 (match v690 fused tiles exactly)
+                     BM = 32; BN = 64; NW = 4; NKDIM = 32
+                 # v751: Read B_q directly (N, K/2) — no transpose needed!
+                 # Pass strides swapped: kernel expects (K/2, N) layout via strides
+                 b_u8 = B_q.view(torch.uint8)  # shape (N, K/2), strides (K/2, 1)
+                 if SK > 1:
+                     pp_key = (m, n, SK, 'pp')
+                     if pp_key not in _inline_scale_cache:
+                         _inline_scale_cache[pp_key] = torch.empty(SK, m, n, dtype=torch.float32, device=A.device)
+                     y_pp = _inline_scale_cache[pp_key]
+                     out_key = (m, n, 'out')
+                     if out_key not in _inline_scale_cache:
+                         _inline_scale_cache[out_key] = torch.empty(m, n, dtype=torch.bfloat16, device=A.device)
+                     out = _inline_scale_cache[out_key]
+                 else:
+                     y_pp = None
+                     out = torch.empty(m, n, dtype=torch.bfloat16, device=A.device)
+                 grid = (SK * triton.cdiv(m, BM) * triton.cdiv(n, BN),)
+                 target = out if y_pp is None else y_pp
+                 _inlined_kernel[grid](
+                     A, b_u8, target, scales,
+                     m, n, K_half,
+                     A.stride(0), A.stride(1),
+                     b_u8.stride(1), b_u8.stride(0),  # SWAPPED: (k_stride, n_stride)
+                     0 if y_pp is None else y_pp.stride(0),
+                     out.stride(0) if y_pp is None else y_pp.stride(1),
+                     out.stride(1) if y_pp is None else y_pp.stride(2),
+                     scales.stride(0), scales.stride(1),
+                     BLOCK_SIZE_M=BM, BLOCK_SIZE_N=BN, BLOCK_SIZE_K=BK,
+                     GROUP_SIZE_M=1, NUM_KSPLIT=SK, SPLITK_BLOCK_SIZE=SPK_BS,
+                     ATOMIC_ADD=False, cache_modifier=".cg",
+                     matrix_instr_nonkdim=NKDIM,
+                     num_warps=NW, num_stages=NS, waves_per_eu=2,
+                 )
+                 if SK > 1:
+                     # Reduce SplitK partials
+                     ACTUAL_SK = triton.cdiv(K_half, SPK_BS // 2)
+                     rg = (triton.cdiv(m, 16), triton.cdiv(n, 64))
+                     _reduce_kernel[rg](
+                         y_pp, out, m, n,
+                         y_pp.stride(0), y_pp.stride(1), y_pp.stride(2),
+                         out.stride(0), out.stride(1),
+                         BM=16, BN=64, ACTUAL_SK=ACTUAL_SK,
+                         MAX_SK=triton.next_power_of_2(SK),
+                     )
+                 return out
+         except Exception as e:
+             P(f"INLINE FAIL ({m},{n},{k}): {e}")
+             _inline_bad.add((m, k, n))
+  
+     # v690 paths (fallback for M>16)
      key = (m, k, n)
      if key not in _cache:
          _cache[key] = _init(m, k, n, A.device)
      c = _cache[key]
   
      if c['mode'] == 'fused':
-         Bq = B_q.view(torch.uint8); Bs = B_scale_sh.view(torch.uint8)
-         _fused_quant_gemm_kernel[c['grid']](A, Bq, Bs, c['out'], m, n, k, A.stride(0), A.stride(1), Bq.stride(0), Bq.stride(1), c['sn_div8_mul256'], c['out'].stride(0), c['out'].stride(1), BLOCK_M=c['BM'], BLOCK_N=c['BN'], BLOCK_K=c['BK'], num_warps=c['NW'], num_stages=1)
+         Bq_uint8 = B_q.view(torch.uint8)
+         Bscale_uint8 = B_scale_sh.view(torch.uint8)
+         _fused_quant_gemm_kernel[c['grid']](
+             A, Bq_uint8, Bscale_uint8, c['out'],
+             m, n, k,
+             A.stride(0), A.stride(1),
+             Bq_uint8.stride(0), Bq_uint8.stride(1),
+             c['sn_div8_mul256'],
+             c['out'].stride(0), c['out'].stride(1),
+             BLOCK_M=c['BM'], BLOCK_N=c['BN'], BLOCK_K=c['BK'],
+             num_warps=c['NW'], num_stages=1, waves_per_eu=c['wpe'],
+         )
          return c['out']
   
      elif c['mode'] == 'splitk':
-         Bq = B_q.view(torch.uint8); Bs = B_scale_sh.view(torch.uint8)
-         SK = c['SPLIT_K']
-         if SK == 1:
-             _fused_splitk_gemm[(c['total_wgs'],)](A, Bq, Bs, c['out'], m, n, k, A.stride(0), A.stride(1), Bq.stride(0), Bq.stride(1), c['sn_div8_mul256'], 0, c['out'].stride(0), c['out'].stride(1), c['grid_m'], c['grid_n'], BLOCK_M=c['BM'], BLOCK_N=c['BN'], BLOCK_K=c['BK'], SPLIT_K=1, XCD_SWIZZLE=c['XCD_SWIZZLE'], num_warps=4, num_stages=1, waves_per_eu=2)
+         Bq_uint8 = B_q.view(torch.uint8)
+         Bscale_uint8 = B_scale_sh.view(torch.uint8)
+         SPLIT_K = c['SPLIT_K']
+         wpe = c['wpe']
+         if SPLIT_K == 1:
+             _fused_splitk_gemm[(c['total_wgs'],)](
+                 A, Bq_uint8, Bscale_uint8, c['out'],
+                 m, n, k,
+                 A.stride(0), A.stride(1),
+                 Bq_uint8.stride(0), Bq_uint8.stride(1),
+                 c['sn_div8_mul256'],
+                 0, c['out'].stride(0), c['out'].stride(1),
+                 c['grid_m'], c['grid_n'],
+                 BLOCK_M=c['BM'], BLOCK_N=c['BN'], BLOCK_K=c['BK'],
+                 SPLIT_K=1, XCD_SWIZZLE=c['XCD_SWIZZLE'],
+                 num_warps=4, num_stages=1, waves_per_eu=wpe,
+             )
+             return c['out']
          else:
-             s = c['scratch']
-             _fused_splitk_gemm[(c['total_wgs'],)](A, Bq, Bs, s, m, n, k, A.stride(0), A.stride(1), Bq.stride(0), Bq.stride(1), c['sn_div8_mul256'], s.stride(0), s.stride(1), s.stride(2), c['grid_m'], c['grid_n'], BLOCK_M=c['BM'], BLOCK_N=c['BN'], BLOCK_K=c['BK'], SPLIT_K=SK, XCD_SWIZZLE=c['XCD_SWIZZLE'], num_warps=4, num_stages=1, waves_per_eu=2)
-             _reduce_splitk[c['reduce_grid']](s, c['out'], m, n, s.stride(0), s.stride(1), s.stride(2), c['out'].stride(0), c['out'].stride(1), SPLIT_K=SK, BLOCK_N=128, num_warps=4)
-         return c['out']
+             scratch = c['scratch']
+             _fused_splitk_gemm[(c['total_wgs'],)](
+                 A, Bq_uint8, Bscale_uint8, scratch,
+                 m, n, k,
+                 A.stride(0), A.stride(1),
+                 Bq_uint8.stride(0), Bq_uint8.stride(1),
+                 c['sn_div8_mul256'],
+                 scratch.stride(0), scratch.stride(1), scratch.stride(2),
+                 c['grid_m'], c['grid_n'],
+                 BLOCK_M=c['BM'], BLOCK_N=c['BN'], BLOCK_K=c['BK'],
+                 SPLIT_K=SPLIT_K, XCD_SWIZZLE=c['XCD_SWIZZLE'],
+                 num_warps=4, num_stages=1, waves_per_eu=wpe,
+             )
+             _reduce_splitk[c['reduce_grid']](
+                 scratch, c['out'], m, n,
+                 scratch.stride(0), scratch.stride(1), scratch.stride(2),
+                 c['out'].stride(0), c['out'].stride(1),
+                 SPLIT_K=SPLIT_K, BLOCK_N=128,
+                 num_warps=4,
+             )
+             return c['out']
   
-     else:  # asm
-         x_fp4 = c['x_fp4']; bs_shuf = c['bs_shuffled']
-         _fused_quant_shuffle_kernel[c['grid']](A, x_fp4, bs_shuf, A.stride(0), A.stride(1), x_fp4.stride(0), x_fp4.stride(1), m, k, c['sn_div8_mul256'], c['sc'], BLOCK_SIZE_M=c['BSM'], BLOCK_SIZE_N=c['BSN'], NUM_ITER=c['NI'], NUM_STAGES=c['NS'], num_warps=c['NW'], waves_per_eu=0, num_stages=1)
-         aiter.gemm_a4w4_asm(x_fp4.view(_FP4X2), B_shuffle, bs_shuf.view(_E8M0), B_scale_sh, c['out'], c['knl'], bpreshuffle=True, log2_k_split=c['l2ks'])
-         return c['out']
+     else:  # asm — v754c: BSM=16 quant (proven optimal) + CK ASM
+         x_fp4 = c['x_fp4']
+         bs_shuf = c['bs_shuffled']
+         _fused_quant_shuffle_kernel[c['grid']](
+             A, x_fp4, bs_shuf,
+             A.stride(0), A.stride(1),
+             x_fp4.stride(0), x_fp4.stride(1),
+             m, k,
+             c['sn_div8_mul256'], c['sc'],
+             BLOCK_SIZE_M=c['BSM'], BLOCK_SIZE_N=c['BSN'],
+             NUM_ITER=c['NI'], NUM_STAGES=c['NS'],
+             num_warps=c['NW'], waves_per_eu=c['quant_wpe'], num_stages=1,
+         )
+         out = c['out']
+         aiter.gemm_a4w4_asm(
+             x_fp4.view(_FP4X2), B_shuffle,
+             bs_shuf.view(_E8M0), B_scale_sh,
+             out, c['knl'],
+             bpreshuffle=True,
+             log2_k_split=c['l2ks'],
+         )
+         return out