#616 RTX A3000 / 0.1.38: AVX2 i-quant gate/up codebook gather — bit-exact ~1.3x isolated, +1.0% e2e, kept as opt-in

Descrição

## Context

Follow-up to #494 (the CPU-pool decode profile on this host). P4 sized i-quant gate/up at 28.2 ms/round = 68% of the CPU expert pool = 36% of a fresh round on 0.1.35, and the kernel documents itself as codebook-lookup bound (`src/kernels/cpu/native_expert.cpp:81-82`). This tests that on the AVX2 path (this host has no AVX-512) and lands the cheap fix: replace the scalar `_mm256_set_epi32` grid assembly with an AVX2 `_mm256_i32gather_epi32`.

**Result: bit-exact, ~1.3x isolated, +1.0% serve end to end, kept as opt-in `STRATA_IQ256_GATHER`.** The codebook lookup is real single-core; the 15-worker pool is not decode-bound enough for it to move the whole token, so the *large* lever remains fewer bytes/pack format, but the gather is a free compounding win where it lands.

Host: RTX A3000 12 GB (sm_86) + i7-12850HX (8 P + 8 E, AVX2 + AVX-VNNI, no AVX-512), 128 GB. Engine **0.1.38** (local source build), Swift 1.5 IQ3_XXS native pack, `--adapt-every 1 --pool-workers 15`, fixed prompt (`--stats`: `hit_rate 0.5509`, 1.78 tok/round). Serve unit stopped for the measurements.

## The split (0.1.38, this config)

| Term | ms/round |
| --- | ---: |
| round | 48.2 (27.09 ms/token × 1.78) |
| GPU "wait for rings" | 22.1 |
| CPU expert pool | 16.7 (35%) |
| — gate/up | **9.3-10.2 (~60% of the pool)** |
| — down | 5.5 |
| pool GB/s over the rows phases | 29.7-32.8 |
| decode | 36.9 tok/s |

## Phase 0 — the lookup is the single-core bound

Harness: `native_expert_parity --synthetic <GU/DOWN>`, one thread, weights L2-resident, **pinned with `taskset -c 2`**. Real geometry H=2560, FF=640.

Gate+up, one token, one thread (AVX2 `iq256_gu_rows` vs ggml-cpu per-token `vec_dot`):

| gate/up type | AVX2 | ggml |
| --- | ---: | ---: |
| IQ3_XXS | 245 us, 5.13 GB/s | 243 us, 5.17 GB/s |
| IQ3_S | 334 us, 4.21 GB/s | 372 us, 3.72 GB/s |
| IQ2_XXS | 193 us, 4.38 GB/s | 208 us, 4.06 GB/s |
| IQ2_XS | 223 us, 4.26 GB/s | 229 us, 4.13 GB/s |
| IQ2_S | 241 us, 4.36 GB/s | 199 us, 5.26 GB/s |

At one token AVX2 ≈ ggml (the documented "no faster at one"); the win is decode-once amortization (`nt`=3 is ~2.2x/token for IQ3_XXS). A scratch build substituting a constant for the grid vector (correctness ignored, reverted) nearly doubles the two 32-bit-grid formats:

| format | full | grid bypassed | speedup |
| --- | ---: | ---: | ---: |
| IQ3_XXS | 248 us, 5.1 GB/s | 130 us, 9.7 GB/s | **1.91x** |
| IQ3_S | 357 us, 4.0 GB/s | 140 us, 10.0 GB/s | **2.55x** |
| IQ2_XXS | 191 us | 138 us | 1.38x |
| IQ2_S | 242 us | 198 us | 1.22x |

So the grid lookup is ~half the single-core kernel for the 32-bit-grid formats (18/21 — 6+13 of 48 layers, ~58% of the gate/up lookup work). A 64-bit-index AVX2 gather for 16/17/22 measured a wash/loss and was dropped.

## H1 — AVX2 gather (bit-exact, ~1.3x isolated, +1.0% e2e)

`STRATA_IQ256_GATHER` selects `_mm256_i32gather_epi32(iq3xxs_grid, cvtepu8_epi32(q), 4)` (same for `iq3s_grid`, 9th bit built with `srlv/add`). Opt-in, default off — the same shape as the AVX-512 kernels' `STRATA_IQ_GATHER` (gather throughput is core-dependent).

- **Parity:** a direct same-binary output dump, scalar vs gather, is bit-identical for all five gate/up formats (`cmp` clean); `native_expert_parity` stays green (same `rel ≈ 3.1e-08` vs ggml's `vec_dot`); width invariance 0 rows; AVX-512 untouched.
- **Isolated, pinned:** IQ3_XXS 245→185 us (1.32x), IQ3_S 334→250 us (1.34x).

End to end (`bench/e2e.sh`, same binary, switch on/off, interleaved, 5 pairs, `decode_tok_s` from `/metrics`):

| rep | off | on |
| ---: | ---: | ---: |
| 1 | 36.4 | 36.8 |
| 2 | 36.9 | 37.0 |
| 3 | 36.5 | 36.9 |
| 4 | 37.1 | 37.1 |
| 5 | 35.9 | 36.9 |
| median | 36.5 | 36.9 |
| mean | 36.56 | 36.94 |

**+1.0% mean, on ≥ off in 5/5.** Deterministic `--stats` is flat (gate/up off median 9.756 / on 9.773 ms/round), i.e. the pool does not move, but the gather is never worse in the engine.

Why it does not move the whole token: at 15 workers the i-quant gate/up is not decode-bound enough. `STRATA_NO_IQ256=1` (ggml per-token `vec_dot`) raises gate/up to 12.8-13.3 ms/round, so the AVX2 multi-token kernel is ~1.35x at pool scale — the pool responds to the kernel — but substituting the *whole* lookup assembly (a superset of any micro-tweak) moves almost nothing: the SMT/E-core concurrency eats most of the single-core decode win.

## Decision

- Kept as opt-in `STRATA_IQ256_GATHER` (default off), branch `perf/iq-gateup`. No numbers change; the engine is unchanged unless the variable is set.
- H2/H3 not pursued: H2 is smaller than the whole-lookup substitution already measured near zero, and H3 needs a new table/pack (out of scope).
- The remaining *large* gate/up lever is **fewer bytes per lookup (pack format)** or the **GPU** (P6/#610), not the runtime decode.
- Measurement caveat for anyone repeating this: pin the harness thread or use the engine A/B. Unpinned isolated runs are bimodal (the same gather run reads 188 or 1028 us). Note the repo's only 3% is `tools/calibrate.py`'s launch-setting `MIN_GAIN`, not a kernel-change bar.

Refs #494 (P4 pool profile), #610 (P6 GPU floor), #485 (PCIe probe).