Round 9 Track C Audit: load_tab<T_IQ3_S> Dead load_signs Removal + Barrier-Free TAB == 1 in moe_gu / moe_down + Decode moe.metal NE == 1 dn_body Specialization #

1. Summary of Changes (metal/kernels/quant.metal & metal/kernels/moe.metal) #

• Patch file: /tmp/argon_r9_quant/argon_r9_quant.patch (verified clean via git -C vendor/strata apply --check).
• Files touched: Only metal/kernels/quant.metal and metal/kernels/moe.metal in scratch copy /tmp/argon_r9_quant/strata (tracked vendor/strata/ untouched).

Change 1: Remove Dead load_signs from load_tab<T_IQ3_S> & Shrink fmt<T_IQ3_S>::TAB (quant.metal lines 41 & 100–103) #

• Every T_IQ3_S decoder across the Metal codebase (deq<T_IQ3_S> in quant.metal, DW<T_IQ3_S> in mma.metal, W4<T_IQ3_S> and W8<T_IQ3_S> in moe_mm.metal) reads raw 8-bit sign bytes directly from the block (b[74 + 4 * ib + l]) and never accesses ksigns_iq2xs (tab + 512).
• Changed fmt<T_IQ3_S>::TAB from 512 + 32 (2176 B) to 512 (2048 B) and removed load_signs(tab, 512, tid, nth); from load_tab<T_IQ3_S>.
• Saves 128 B of threadgroup memory and 128 byte loads/stores per threadgroup across gemv, moe_gu, mm_gu*, dq_f16, and mtp_chain.

Change 2: Eliminate Redundant threadgroup_barrier When fmt<F>::TAB == 1 in moe_gu & moe_down (moe.metal lines 331 & 421) #

• Guarded load_tab<F> + threadgroup_barrier(mem_flags::mem_threadgroup) with if (fmt<F>::TAB > 1 && moe_probe != 1) in moe_gu and if (fmt<F>::TAB > 1) in moe_down (matching dq_f16 in dense_mm.metal).
• Since fmt<F>::TAB is a compile-time enum constant, the Metal compiler dead-code eliminates the barrier in all TAB == 1 specializations (T_Q2_0, T_Q2_0R, T_Q3_K, T_Q4_K, T_Q5_K, T_Q6_K, T_Q8_0, T_Q5_1, T_Q4_0, T_Q5_0, T_BF16, T_F32).
• Eliminates 48 moe_down (T_Q2_0 / T_Q2_0R) and 12 moe_gu (T_Q6_K) unconditional threadgroup barriers per decode / MTP verify step.

Change 3: Specialize NE == 1 in dn_body (moe.metal lines 373–382) #

• Mirrored gu_body's if (NE == 1) branch into dn_body: for single-entry expert groups (NE == 1, which is 100% of groups at ntok = 1 and the majority of groups in ntok = 2..4 verify windows), X<F>::load(h + c * int(a.n_ff), ...) is hoisted outside the for (int r = 0; r < NR; ++r) loop instead of reloading xs[16] inside the r loop (for Q2_0 / Q2_0R with NR = 2) or allocating float q[NR][UNIT] across rows before loading x.
• Ruled out: Manual unrolling of for (int l2 = 0; l2 < 2; ++l2) in deq<T_IQ2_XS> (quant.metal lines 182–193) was benchmarked and rejected because it increased live register pressure across NR rows in gu_body on Apple M5, regressing moe_gu IQ2_XS (QUANT_CTAB=1) from 100.7 us (94.0 GB/s) to 105.2 us (90.0 GB/s) at ntok = 1 and from 283.1 us (98.2 GB/s) to 293.4 us (94.8 GB/s) at ntok = 3. Keeping the loop form in deq<T_IQ2_XS> preserves optimal register allocation.

2. Bitwise Parity Verification (100% Bit-for-Bit Identical) #

1. expert_parity (48 configurations):
   • Tested all combinations of gu in {IQ3_S, IQ2_XS, Q6_K, Q2_0}, dn in {Q2_0, Q8_0, IQ4_NL}, and ntok in {1, 2, 3, 4}.
   • Verified 64-bit FNV-1a hashes of GPU output buffers h (4 * ntok * TOPK * N_FF bytes) and y (4 * ntok * N_EMBD bytes) between baseline and patched kernels: 48 / 48 configurations matched with 0 bit differences (mismatches: 0).
2. moe_mm_parity --v3 --tile 16,64,32,2 --gu IQ3_S --dn Q2_0 --dump:
   • Verified 6,512,640-byte dump of HS and y: 100% byte-for-byte identical (b1 == b2).
3. gemv_parity:
   • All 13 formats (F32, BF16, Q8_0, Q4_0, Q5_0, IQ4_NL, IQ4_XS, Q2_0, IQ3_S, Q3_K, Q4_K, Q5_K, Q6_K), n_in in {2560, 640, 10240}, ntok in {1, 3, 8} across 20 tile configs: ALL PASS (bitwise window == one token).
4. dense_mm_parity:
   • All formats including IQ3_S (2560 -> 100, 2560 -> 48, 6144 -> 100, 6144 -> 48): ALL PASS (memcmp(wh)=0, memcmp(xh)=0).

3. Decode expert_parity --bench Results on Apple M5 (ntok = 1, 2, 3, 4) #

3-of-3 median GPU timings (--layers 4 --reps 48 --share 0.3, production tile configs: gu = 32,1,4 for ntok = 1,2 and 16,1,4 for ntok = 3,4; dn = 8,2,4):