int8 leg vs the SERVED trellis pin at DECODE-shaped M, 1 to 2048 rows
=======================================================================

WHAT THIS DECIDES.  Whether the exl3 int8 GEMM should serve the DECODE class
(and, by the same numbers, a VERIFY slate or a prefill sub-chunk) at ANY row
count a deployment can present, and where the crossover against the leg the
engine serves sits on THIS build.  Until this receipt DECODE was refused by a
categorical step-class label ("off-class"), on the strength of
docs/receipts/int8-verify-small-m-2026-09-05.txt -- a file that scopes itself
to --max-batch 8 in its own words, was measured before four tile-selection
commits (f45571161, 93f6cbad0, 14a14b881, ba4e44170), timed the int8 leg on a
tile chosen for another class, and had already watched its own crossover move
64 -> 40 -> 36.  No decode-shaped M above 32 had ever been timed on a
class-appropriate tile against the SERVED member.  This is that measurement.

VERDICT, IN ONE LINE.  GO.  On all three linears the int8 chain loses to the
served trellis member through M = 32 (1.20x / 0.83x / 1.36x at 32 rows on
out/gate/down) and WINS from M = 36 on (0.66x / 0.56x / 0.71x), and it keeps
winning through 2048 rows -- 0.20x / 0.21x / 0.23x of the trellis leg at 2048,
and 0.65x / 0.62x / 0.68x of leg B, which is the leg the engine actually
serves above auto_reconstruct_threshold (280 rows).  The crossover is the
served pin's DEEPEST TILE + 1: every shipped 27B family caps at TILESIZE_M 32,
exl3_gemm tiles M serially with a grid.sync between tiles and re-reads the
whole packed weight per tile, so at 33 rows the trellis leg pays a second
full weight pass (legA us doubles: 28.7 -> 52.6 out_proj, 67.3 -> 120.0
gate_proj, 66.0 -> 126.2 down_proj between M = 32 and 36) while the int8
chain still pays one.  That edge is the threshold the engine now derives
(int8_policy.int8_decode_min_rows_from_tiles), not a number typed here.

MEASURED AGAINST
  commit      6b24e9229 on branch int8/m-keyed-routing = origin/main
              125e0de8c (#2168) + the harness/overlay commits of this branch;
              arbi_serve bind-mounted over registry.arbi.work/arbi-serve:test-latest
              (image revision 318006b38, exllamav3 cb76869b).  The int8
              kernel csrc differs from the image's by comments and a
              launch-bounds macro that expands to the identical
              __launch_bounds__ when undefined, so the JIT build is the
              served binary; built narrowed (I8_ONLY_BITS=4 I8_ONLY_CB=2 --
              the 27B is 4 bpw, mul1), which changes the build time and not
              the served instantiation (compose.int8-ab.yaml).
  tkv         NOT AN INPUT to this receipt, checked rather than assumed: the
              image bakes turbo-attn 0.53.0 against main's 0.54.0 pin
              (#2180), and nothing on the timed path reads it -- the int8
              kernel is arbi_serve/weight_quant/exl3/csrc/exl3_i8_gemm.cu,
              its A side is csrc/exl3_a_prep.cu, the trellis leg is
              exllamav3's own extension at the pinned revision, and the
              trellis tile table is read off exllamav3's header
              (kernel_shape_table.py).  `grep -l "turbo_attn\|tkv"` over
              bench_group_scale.py, ext.py, ckpt.py, gpu_measure.py,
              int8_kernel.py, a_prep.py and kernel_shape_table.py matches
              nothing.  The served-stack validation of the routing change
              (below, PR) boots with TURBO_ATTN_SRC_ROOT at a 0.54.0
              checkout and records what the engine says it loaded.
  card        RTX 4090 on 10.1.0.201, GPU 0, sm_89, 128 SM, sole tenancy on
              the card (tenancy checked before and after every run: none /
              this pid only).  NO clock pin: idle 210 MHz / 34 W between runs,
              no locked graphics clock.  Under load, sampled concurrently
              every 50 ms inside the timed blocks only: sm clock median
              2775 MHz (min 2265, max 2790), power median 329 W (out_proj
              run) and 396 W (gate/down run), 450 W limit, sw_power_cap
              asserted on 33% / 44% of samples, hw_slowdown 0.  The int8
              arms at large M are the cells that reach the cap; the trellis
              arms at small M do not, and the paired ratios are taken within
              one rep so the cap lands on both arms of a pair.
  model       Qwen3.8-27B-exl3-4.0bpw (mul1 codebook, K=4), layer 0:
              out_proj 6144x5120, gate_proj 5120x17408, down_proj 17408x5120.
              The other four 27B geometries were measured to 64 rows in
              docs/receipts/int8-verify-slate-geometry-2026-09-07.txt with the
              same cliff at 36 on every one of them; they are not re-timed
              here.
  harness     tools/int8_gemm/bench_group_scale.py at this commit: reps 7,
              burst >= 30 ms, L2 scrub 'read', ratios paired WITHIN a rep over
              a round-robin reshuffled each rep, rot 6 on out_proj and 2 on
              gate/down so the trellis copies exceed the 72 MB L2, roofline
              refusal on every arm, bit-identity witness across the timed
              tiles before every cell.  Run through the committed composes
              (compose.yaml + compose.bindmount.yaml + compose.int8-ab.yaml +
              compose.tools.yaml, `docker compose run`), never a hand-rolled
              docker run.
  legA        THE SERVED TRELLIS MEMBER: served_trellis_pin() resolves the
              shipped family under the boot's own key and walks
              _pin_for_rows() to the member the engine launches at that M --
              K32N128 s2 (TILE_M 16) at <= 16 rows and s5 (TILE_M 32) above
              for out_proj/down_proj; K16N512 s4 / s10 for gate_proj.  Never
              force_shape_idx = -1.  Below 280 rows this is the leg decode
              serves.
  legB        reconstruct_had_slice + cuBLAS hgemm, the leg the dispatch
              serves ABOVE auto_reconstruct_threshold (280 rows on these
              widths).  For M >= 512 legB, not legA, is the incumbent.
  a<sid>      the SERVED int8 chain (int8_gemm_adapter: one a_prep_frag pass
              then the GEMM with svh folded into the epilogue), on the four
              tiles the engine can pin: s111 (TILE_M 128, the prefill pin),
              s85 (64), s94 (32), s98 (16).  The per-geometry verify record
              in shipped_pins.json names s85 for gate_proj and leaves
              out_proj/down_proj on the band ladder (s98 <= 48 rows, s94
              above); both selections are in the table.

THE TABLE.  us/call medians; ratios <1.00 mean the int8 leg is faster.
legA pin is the served trellis member at that M.  best = fastest int8 tile.

linear        M legA pin         legA us legB us  a111 us   a98 us   a94 us   a85 us    s111/A   s98/A   s94/A   s85/A  best  best/A  best/B
out_proj      1 K32N128 s2 T16     22.63  169.43    81.93    33.83    40.19    51.54    3.6216  1.4951  1.7764  2.2779   s98  1.4949  0.1997
out_proj      2 K32N128 s2 T16     22.69  123.57    82.17    33.96    40.47    51.72    3.6228  1.4972  1.7834  2.2803   s98  1.4967  0.2748
out_proj      4 K32N128 s2 T16     22.73  126.91    82.28    34.03    40.45    51.70    3.6165  1.4972  1.7793  2.2744   s98  1.4971  0.2681
out_proj      8 K32N128 s2 T16     22.91  130.03    82.44    34.13    40.44    51.78    3.5986  1.4892  1.7651  2.2600   s98  1.4897  0.2625
out_proj     16 K32N128 s2 T16     23.75  130.39    82.65    34.44    40.61    51.96    3.4800  1.4501  1.7096  2.1877   s98  1.4501  0.2641
out_proj     24 K32N128 s5 T32     27.81  115.53    82.77    34.50    42.21    52.08    2.9822  1.2392  1.5182  1.8740   s98  1.2406  0.2986
out_proj     32 K32N128 s5 T32     28.74  128.79    83.15    34.43    42.33    52.31    2.8945  1.1981  1.4738  1.8216   s98  1.1980  0.2673
out_proj     36 K32N128 s5 T32     52.60  130.61    83.50    34.72    42.55    52.53    1.5842  0.6597  0.8087  0.9981   s98  0.6601  0.2658
out_proj     40 K32N128 s5 T32     52.45  118.00    83.02    34.65    42.40    52.39    1.5832  0.6602  0.8073  0.9990   s98  0.6606  0.2936
out_proj     48 K32N128 s5 T32     53.50  113.87    83.36    34.69    42.45    52.57    1.5567  0.6484  0.7926  0.9829   s98  0.6484  0.3046
out_proj     64 K32N128 s5 T32     55.16  115.04    83.69    51.70    42.60    52.78    1.5176  0.9371  0.7720  0.9571   s94  0.7723  0.3703
out_proj     96 K32N128 s5 T32     82.15  133.73    85.18    53.63    43.18    53.31    1.0384  0.6542  0.5241  0.6518   s94  0.5256  0.3229
out_proj    128 K32N128 s5 T32    108.35  136.69    85.99    74.68    62.54    53.63    0.7934  0.6895  0.5771  0.4949   s85  0.4950  0.3923
out_proj    192 K32N128 s5 T32    165.75  154.80    88.36   100.63    67.51    58.56    0.5345  0.6054  0.4073  0.3492   s85  0.3533  0.3783
out_proj    256 K32N128 s5 T32    223.27  156.79    88.36   125.14    98.95    84.10    0.3974  0.5602  0.4554  0.3752   s85  0.3767  0.5364
out_proj    512 K32N128 s5 T32    438.04  225.66   138.14   247.25   163.58   143.26    0.3104  0.5645  0.3711  0.3096  s111  0.3154  0.6122
out_proj   1024 K32N128 s5 T32    905.09  341.83   230.30   480.80   326.44   232.76    0.2529  0.5390  0.3549  0.2534  s111  0.2544  0.6737
out_proj   2048 K32N128 s5 T32   1906.43  589.43   382.42   974.82   655.76   452.19    0.2036  0.5079  0.3384  0.2389  s111  0.2006  0.6488
gate_proj     1 K16N512 s4 T16     53.27  478.38   102.87    53.28    54.50    66.31    1.9297  1.0003  1.0231  1.2451   s98  1.0002  0.1114
gate_proj     2 K16N512 s4 T16     53.03  481.38   103.11    53.27    54.58    66.46    1.9444  1.0052  1.0295  1.2529   s98  1.0045  0.1107
gate_proj     4 K16N512 s4 T16     53.61  483.28   103.13    53.33    54.56    66.45    1.9238  0.9945  1.0174  1.2392   s98  0.9948  0.1104
gate_proj     8 K16N512 s4 T16     54.45  486.38   103.25    53.46    54.70    66.69    1.8966  0.9818  1.0048  1.2246   s98  0.9818  0.1099
gate_proj    16 K16N512 s4 T16     58.61  493.15   103.58    53.72    54.98    66.87    1.7675  0.9163  0.9382  1.1417   s98  0.9166  0.1089
gate_proj    24 K16N512 s10 T32    63.23  480.00   104.34    62.30    55.33    67.14    1.6502  0.9835  0.8735  1.0619   s94  0.8751  0.1153
gate_proj    32 K16N512 s10 T32    67.28  481.00   104.40    62.41    55.63    67.50    1.5517  0.9277  0.8264  1.0034   s94  0.8268  0.1157
gate_proj    36 K16N512 s10 T32   120.04  477.30   104.46    82.36    80.45    67.70    0.8713  0.6849  0.6724  0.5654   s85  0.5640  0.1418
gate_proj    40 K16N512 s10 T32   121.92  478.14   105.03    82.34    80.39    67.88    0.8615  0.6756  0.6593  0.5570   s85  0.5568  0.1420
gate_proj    48 K16N512 s10 T32   125.42  478.18   105.13    82.88    80.51    68.07    0.8359  0.6578  0.6420  0.5428   s85  0.5427  0.1424
gate_proj    64 K16N512 s10 T32   138.24  479.77   105.71   102.00    81.01    68.93    0.7619  0.7371  0.5867  0.4952   s85  0.4986  0.1437
gate_proj    96 K16N512 s10 T32   199.69  490.56   108.66   142.06   103.57   107.64    0.5457  0.7130  0.5186  0.5334   s94  0.5187  0.2111
gate_proj   128 K16N512 s10 T32   269.61  488.56   112.67   184.56   130.37   111.99    0.4171  0.6804  0.4856  0.4113   s85  0.4154  0.2292
gate_proj   192 K16N512 s10 T32   420.51  513.46   173.88   261.84   183.87   142.63    0.4086  0.6206  0.4373  0.3356   s85  0.3392  0.2778
gate_proj   256 K16N512 s10 T32   570.60  530.45   184.47   340.91   237.92   178.91    0.3282  0.6012  0.4202  0.3136   s85  0.3135  0.3373
gate_proj   512 K16N512 s10 T32  1121.70  681.43   296.93   687.51   441.66   326.15    0.2622  0.6129  0.3937  0.2873  s111  0.2647  0.4357
gate_proj  1024 K16N512 s10 T32  2336.67  976.57   528.96  1394.69   904.64   611.20    0.2284  0.5921  0.3956  0.2667  s111  0.2264  0.5417
gate_proj  2048 K16N512 s10 T32  4710.62 1622.07  1004.95  2734.66  1868.68  1210.21    0.2128  0.5805  0.3959  0.2586  s111  0.2133  0.6195
down_proj     1 K32N128 s2 T16     54.72  480.24   224.07    88.67   106.63   138.71    4.0938  1.6211  1.9486  2.5362   s98  1.6204  0.1846
down_proj     2 K32N128 s2 T16     54.66  481.04   224.26    88.94   106.76   138.84    4.1036  1.6283  1.9545  2.5398   s98  1.6271  0.1849
down_proj     4 K32N128 s2 T16     55.21  481.96   224.03    88.94   106.92   138.69    4.0580  1.6109  1.9367  2.5125   s98  1.6109  0.1845
down_proj     8 K32N128 s2 T16     54.81  483.43   224.17    89.21   106.71   138.83    4.0911  1.6275  1.9470  2.5329   s98  1.6276  0.1845
down_proj    16 K32N128 s2 T16     56.59  486.33   224.47    89.29   107.11   138.92    3.9661  1.5780  1.8929  2.4551   s98  1.5778  0.1836
down_proj    24 K32N128 s5 T32     64.76  476.03   224.69    89.46   111.22   139.20    3.4704  1.3820  1.7165  2.1496   s98  1.3814  0.1879
down_proj    32 K32N128 s5 T32     66.00  476.97   225.21    89.55   111.30   139.41    3.4123  1.3554  1.6864  2.1111   s98  1.3568  0.1877
down_proj    36 K32N128 s5 T32    126.18  480.52   225.59    90.00   111.75   139.50    1.7857  0.7129  0.8857  1.1067   s98  0.7133  0.1873
down_proj    40 K32N128 s5 T32    126.39  480.95   225.43    90.24   111.88   139.92    1.7839  0.7135  0.8849  1.1074   s98  0.7140  0.1876
down_proj    48 K32N128 s5 T32    127.57  481.10   226.92    90.22   111.85   139.85    1.7741  0.7066  0.8773  1.0964   s98  0.7072  0.1875
down_proj    64 K32N128 s5 T32    129.59  481.67   226.15   138.24   111.98   140.22    1.7450  1.0665  0.8637  1.0820   s94  0.8641  0.2325
down_proj    96 K32N128 s5 T32    194.69  486.21   230.34   147.04   114.24   144.23    1.1832  0.7569  0.5857  0.7419   s94  0.5868  0.2350
down_proj   128 K32N128 s5 T32    260.89  488.50   229.59   212.34   170.99   143.03    0.8890  0.8133  0.6450  0.5490   s85  0.5482  0.2928
down_proj   192 K32N128 s5 T32    414.38  508.94   235.88   290.60   185.07   154.33    0.5735  0.6979  0.4519  0.3790   s85  0.3724  0.3032
down_proj   256 K32N128 s5 T32    523.09  525.08   241.56   351.02   287.85   234.84    0.4610  0.6810  0.5552  0.4473   s85  0.4489  0.4472
down_proj   512 K32N128 s5 T32   1124.50  654.57   378.88   698.95   457.86   384.36    0.3359  0.6148  0.4086  0.3485  s111  0.3369  0.5788
down_proj  1024 K32N128 s5 T32   2271.67  863.35   635.15  1369.79   929.95   649.75    0.2804  0.6116  0.4117  0.2849  s111  0.2796  0.7357
down_proj  2048 K32N128 s5 T32   4607.10 1542.23  1055.12  2799.80  1883.61  1269.93    0.2305  0.6148  0.4097  0.2769  s111  0.2290  0.6842

Paired intervals (min-max of the within-rep ratio over 7 reps) are in the run
logs beside every row; through 64 rows they are a few tenths of a percent
wide, and at 192-2048 rows -- the cells that sit on the power cap -- they
widen to 5-10% of the ratio, which is small against ratios of 0.2-0.5.

THE THREE QUESTIONS

(i) WHERE THE CROSSOVER IS NOW vs the 09-05 receipt's 40 (36 since #2150).
    Between 32 and 36 rows on every linear, on the tile the engine pins for
    those rows: out_proj 1.198 -> 0.660, gate_proj 0.827 -> 0.564, down_proj
    1.357 -> 0.713.  It has not moved since the 09-07 geometry sweep on box
    106 (which found 36 on all seven geometries against the served member);
    it HAS moved from the 09-05 receipt's 40/64, and this file explains why
    those moved: they were measured against the UNPINNED dispatch and on a
    prefill tile.  The crossover is not a stopwatch reading, it is the
    trellis leg's second pass, which begins at deepest-served-tile + 1 = 33
    rows on this build.  Nothing between 33 and 35 is a captured rung on any
    shipped recipe; at K = 5 the first rung is 36 and at K = 7 it is 40, and
    decode's own B ladder pads an intermediate B up to its next captured
    rung, so 33 admits exactly the rungs 36 admits and the threshold is
    stated as the mechanism's edge rather than as the first rung.
    gate_proj is the wide-N geometry where the int8 leg is within 2% of the
    trellis leg from 1 row (1.00 at M = 1-2, 0.98-0.92 at 8-16, 0.83 at 32)
    -- a marginal win BEFORE the cliff on the one geometry that carries the
    most weight bytes, and a wash whole-model (the 09-07 receipt's 24-32
    rows: pergeom/served 0.99-1.03).  The rule does not chase it: below the
    cliff out_proj and down_proj lose 1.2-1.6x and the sum is flat.

(ii) DOES int8 WIN DECODE AT ANY M A REAL DEPLOYMENT PRODUCES?  Yes, at
    every decode step of 36 rows or more, i.e. any deployment whose
    --max-batch is >= 36 and whose traffic reaches that concurrency.  The win
    grows with the batch: 0.56-0.71x at 36, 0.50-0.77x at 64, 0.42-0.55x at
    128, 0.31-0.45x at 256.  A --max-batch 8 deployment never presents such
    a step and is UNCHANGED by construction: at 1-8 rows the int8 leg is
    1.49x / 1.00x / 1.62x SLOWER on out/gate/down and the rule refuses by
    name.  DECODE DEVICE TIME AT THE SERVED SHAPES IS THEREFORE >= BEFORE
    ON THIS CARD: every M the rule admits is a cell where the int8 arm is
    faster than the served member in the table above, and every M it
    declines keeps the served member byte-for-byte.
    Above 280 rows (a --max-batch above 280) decode is served by leg B
    today, and the int8 leg beats that too: 0.61-0.74x at 512-2048 rows.
    A verify slate of max_batch x (K+1) = 2048 rows at --max-batch 256 is
    the same cell (0.20x of the trellis leg, 0.65x of leg B).

(iii) IS THE 0.23pp EXPECTED-ACCEPT COST THE ONLY PRICE?  Stated rather than
    re-derived: that figure (docs/receipts/int8-p16c6-verdict.md,
    memory 'int8 accuracy verdict') was measured ON ONE CHECKPOINT, and it is
    the cost of the int8 arithmetic vs the trellis arithmetic on the rows
    that take it.  This receipt does not re-measure acceptance -- it is a
    kernel receipt -- and it does not claim the figure transfers.  What it
    can say: the SERVED A/B that moved the verify threshold 40 -> 36
    (docs/receipts/int8-verify-min-rows-ab-2026-09-07.txt, 3 boots per
    arm) read acceptance FLAT (+0.006 accept_len on the treated rung against
    a 0.031 null-control spread), which is the same arithmetic on the same
    rows this rule admits for decode.  The other prices are the ones every
    int8 row already pays and this rule does not add to: ~84 MiB of scratch
    reserved at boot for the widest pinned call (unchanged -- decode M <=
    max_batch is narrower than the verify slate the scratch is already sized
    for), and a decode row's arithmetic above the threshold being the int8
    chain's rather than the trellis leg's.  That last one is the operator's
    dial (exl3_int8_decode = 0 / 1 / auto), not a hard refusal.

WHAT A BOOT-TIME PROBE WOULD COST, measured on this card
  The served set is small: decode M in the captured B ladder <= max_batch,
  verify M = B*(K+1) per captured B, prefill M = chunk.  With only the arms a
  boot would run (the four int8 tiles and the served trellis member; --arms
  a,legA), reps 3, burst >= 10 ms, one (linear, M) cell costs 0.16-0.21 s of
  wall clock (run4-probe-cost, 12 cells, 0.29 s for the first which includes
  the scrub buffer).  Seven geometries x ~10 served widths ~ 70 cells ~ 12-15
  s per boot, per card, before the checkpoint's linears are even loaded --
  affordable, and STILL NOT WHAT SHIPS.  The threshold is derived from the
  served ladder instead, at zero cost, because the crossover is the trellis
  leg's second pass and not a race: a probe would re-discover 33 on every
  boot and add a boot-to-boot variable at exactly the row where the two legs
  are 2x apart.  The probe cost is recorded so that the day a card or a
  family makes the edge a gradual crossing rather than a cliff, the next
  person knows the stopwatch is 15 s away.

WHAT SHIPS FROM THIS (PR on branch int8/m-keyed-routing)
  exl3_int8_decode = auto (0 / 1 / auto, the verify vocabulary): DECODE takes
  the int8 leg on its own declared rows from exl3_int8_decode_min_rows, which
  is `auto` = derived at the int8 arm seam from the trellis pins this boot
  resolved (deepest TILESIZE_M + 1; 33 on this build) or an operator int.
  DECODE rides the SAME rows-keyed tile ladder VERIFY serves -- every band
  shares one reduction key, so a decode row and a verify row on the leg take
  one arithmetic, which is the alignment the class channel exists to keep.
  Named refusals on its own counters: 'decode batch below min rows' (the
  rule declined), 'decode min rows unresolved' (no trellis pin to derive
  from -- a check that could not run does not report success), 'off-class'
  (mode 0).

REPRODUCING
  ARBI_IMAGE=registry.arbi.work/arbi-serve:test-latest ARBI_SRC_ROOT=<worktree> \
  ARBI_MODELS_ROOT=/mnt/k8scache/models ARBI_GPU_DEVICE=0 I8_ONLY_BITS=4 I8_ONLY_CB=2 \
  docker compose -f compose.yaml -f compose.bindmount.yaml -f compose.int8-ab.yaml \
    -f compose.tools.yaml run --rm --no-deps --entrypoint python \
    -w /work/tools/int8_gemm arbi-serve bench_group_scale.py \
    --rows 1 2 4 8 16 24 32 36 40 48 64 96 128 192 256 512 1024 2048 \
    --shapes 111 98 94 85 --linears out_proj --rot 6 --reps 7 --min-burst-ms 30
  (gate_proj down_proj with --rot 2; the probe-cost run adds
  --arms a,legA --reps 3 --min-burst-ms 10 --rows 1 8 64 512.)
