The VERIFY slate, per served row width, on all seven 27B geometries
===================================================================

WHAT THIS DECIDES, and what the earlier receipt beside it could not.
`int8-verify-small-m-2026-09-05.txt` set BOTH the `auto` threshold
(`exl3_int8_verify_min_rows`) and the tile band edge (`INT8_VERIFY_BAND_ROWS`)
from a sweep of ONE linear -- out_proj -- at M in {32,40,48,56,64}; the other
two it timed (gate_proj, down_proj) were measured only at {8,16,32,64}, and the
four remaining geometries of the served checkpoint were never measured at all.
This is the same question asked at every slate width BOTH shipped recipes
capture, on ALL SEVEN geometries, against the leg the engine actually serves.

Three things move.

  1. THE INCUMBENT WAS THE WRONG LEG.  The old receipt's `legA` arm calls
     `exl3_gemm` with `force_shape_idx = -1`, which is the UNPINNED dispatch --
     the GEMV special case below nine rows and the autotuner above it, free to
     draw a member of ANY family.  No boot serves that: `resolve_kernel_shape_pins`
     freezes one (TILESIZE_K, TILESIZE_N) family per geometry before the first
     forward and every served call passes `_pin_for_rows(ladder, rows)`.  The
     two agree exactly up to 32 rows and diverge above, because the shipped
     families cap at TILESIZE_M 32 while the autotuner may draw K16N128's
     48/64/80/96.  `--lega-shape` adds the `legAp` arm, which is the served
     member.  Every ratio below is against `legAp`.

  2. THE TILE THAT WINS IS A FUNCTION OF OUTPUT WIDTH, not of the row band.
     A wide-N geometry already has enough n-tiles to fill 128 SMs and wants the
     DEEP m-tile (s85, TILE_M 64) so the slate stops paying padding; a narrow-N
     one is grid-starved and wants the shallow tile (s98, TILE_M 16).  The band
     ladder (s98 to 48 rows, s94 above) fits out_proj and down_proj -- the two
     it was measured on -- and is wrong on the other five.

  3. THE CROSSOVER IS BELOW 40 ROWS, at every geometry.  Against the SERVED
     leg the int8 chain wins at M = 36 on all seven; `min_rows = 40` leaves
     that width -- `compose.serve.yaml`'s B = 6 rung, K = 5, S = 6 -- on a leg
     it beats by 1.4-1.6x.  NOT CHANGED HERE: moving the threshold moves a
     slate onto different arithmetic and owes an acceptance A/B, which no
     kernel bench can pay.  Recorded as a finding, not as a default.

WHY THE CLIFF IS THERE AT ALL, structurally and without a stopwatch.
`exl3_gemm` walks M in TILESIZE_M steps with a `grid.sync()` between them and
re-stages and re-decodes the whole packed weight inside each step, so a linear
at M rows under a pin whose deepest legal tile is T reads its trellis
`ceil(M / T)` times.  Deeper members exist ONLY at TILESIZE_K 16 --
`exl3_kernel_map.cuh` says why: "TILESIZE_K 16 keeps blockDim at 256, where the
register cap is 255 rather than 128" -- and of the four families the 27B's
geometries are pinned to, three cap at 32 rows and one (K32N256, q_proj) at 16.
So the whole verify forward reads its weights TWICE from 33 rows up:

verify-forward trellis bytes (400 linears): 11.328 GiB

   M  B(K=7)  B(K=5)  trellis GiB read  passes(wtd)  vs M<=16
   1       -       -            11.328        1.000     1.000
   2       -       -            11.328        1.000     1.000
   4       -       -            11.328        1.000     1.000
   6       -     1.0            11.328        1.000     1.000
   8     1.0       -            11.328        1.000     1.000
  12       -     2.0            11.328        1.000     1.000
  16     2.0       -            11.328        1.000     1.000
  24     3.0     4.0            11.797        1.041     1.041
  30       -     5.0            11.797        1.041     1.041
  32     4.0       -            11.797        1.041     1.041
  36       -     6.0            23.125        2.041     2.041
  40     5.0       -            23.125        2.041     2.041
  42       -     7.0            23.125        2.041     2.041
  48     6.0     8.0            23.125        2.041     2.041
  56     7.0       -            23.594        2.083     2.083
  64     8.0       -            23.594        2.083     2.083

per-geometry passes  (family, legal tiles -> deepest)
        geometry    n    family                  tiles    1    2    4    6    8   12   16   24   30   32   36   40   42   48   56   64
   5120x17408xK4  128   K16N512               [16, 32]    1    1    1    1    1    1    1    1    1    1    2    2    2    2    2    2
   17408x5120xK4   64   K32N128               [16, 32]    1    1    1    1    1    1    1    1    1    1    2    2    2    2    2    2
    6144x5120xK4   64   K32N128               [16, 32]    1    1    1    1    1    1    1    1    1    1    2    2    2    2    2    2
   5120x10240xK4   48   K32N128               [16, 32]    1    1    1    1    1    1    1    1    1    1    2    2    2    2    2    2
    5120x6144xK4   48   K32N128               [16, 32]    1    1    1    1    1    1    1    1    1    1    2    2    2    2    2    2
    5120x1024xK4   32   K32N128               [16, 32]    1    1    1    1    1    1    1    1    1    1    2    2    2    2    2    2
   5120x12288xK4   16   K32N256                   [16]    1    1    1    1    1    1    1    2    2    2    3    3    3    3    4    4
  5120x248320xK6    1   K16N512               [16, 32]    1    1    1    1    1    1    1    1    1    1    2    2    2    2    2    2

  The family pin is raced (and RECORDED, in shipped_pins.json) over
  `_PIN_SELECT_ROWS = (1, 16, 32)`.  All three are on the flat side of that
  step, so the selection cannot see the doubling it causes -- and it is a
  doubling of the quantity the verify forward is bound by.

CONFIG
  card        RTX 4090 on box 106, sm_89, 128 SM, no clock pin, 450 W limit.
              NOT sole tenancy: one IDLE co-tenant holds 6.4 GiB and sat at
              210 MHz / 33 W before and after every run (checked at both ends;
              container `arbi-verify-eff`).  Clocks under load: sm median
              2640-2760 MHz, power median 153-441 W, `sw_power_cap` asserted on
              0-79% of samples in the heaviest cells, `hw_slowdown` 0 in all.
  model       Qwen3.8-27B-exl3-4.0bpw, mcg codebook; layer 0 for the GDN and
              MLP linears, layer 3 for the full-attention ones.
  kernel      the int8 GEMM as baked into registry.arbi.work/arbi-serve:test-latest.
              `csrc/exl3_i8_gemm.cu` differs from main by comments plus an
              `EXL3_I8_MIN_BLOCKS_PER_SM` macro that expands to the identical
              `__launch_bounds__` when undefined, so the default build is the
              same binary.  `a_prep` JIT-built from the image's own source.
  harness     tools/int8_gemm/bench_group_scale.py, reps 7, burst >= 30 ms,
              L2 scrub 'read', ratios pooled WITHIN a rep over a round-robin
              reshuffled each rep, rot chosen per geometry so the trellis
              copies exceed the 72 MB L2 (the roofline guard REFUSED
              in_proj_qkv at rot 2 for exactly that reason, and it was re-run
              at rot 6).
  arms        legA  = the unpinned dispatch (the old receipt's incumbent)
              legAp = the member of this geometry's SHIPPED family that
                      `_pin_for_rows` selects at these rows: shape 2 (<=16
                      rows) / 5 (above) for K32N128, 4 / 10 for K16N512,
                      3 always for K32N256.
              a<sid> = the SERVED int8 chain: one `a_prep_frag` pass then the
                      GEMM with svh folded into its epilogue.
  witness     every cell ran the harness's bit-identity witness across the
              timed shapes before timing them; s85, s94, s98 and the prefill
              pin s111 share one `int8_reduction_key` ((TILE_K, group width) =
              (128, 128)), so which of them a geometry takes is a block count
              and not an arithmetic change.

per-geometry us/call  (legA = UNPINNED dispatch, legAp = the PINNED family member the engine serves)
      linear    n    M     legA    legAp      a98      a94      a85   best  best/legAp  shipped/legAp
   gate_proj  128    8    52.60    55.04    53.52    54.97    67.40     98      0.9724         0.9724
   gate_proj  128   16    53.64    58.82    53.82    55.17    67.75     98      0.9150         0.9150
   gate_proj  128   24    64.18    64.15    64.47    55.57    71.00     94      0.8663         1.0050
   gate_proj  128   30    67.35    67.12    65.00    55.65    68.32     94      0.8291         0.9684
   gate_proj  128   32    65.46    68.92    65.31    56.33    68.82     94      0.8173         0.9476
   gate_proj  128   36   102.13   121.40    88.40    82.95    68.68     85      0.5657         0.7282
   gate_proj  128   40   103.88   122.60    91.39    83.74    68.97     85      0.5626         0.7454
   gate_proj  128   42   101.94   125.31    90.27    85.91    70.53     85      0.5628         0.7204
   gate_proj  128   48   103.95   129.53    89.82    83.98    69.93     85      0.5399         0.6934
   gate_proj  128   56   121.03   132.14   109.85    84.49    69.89     85      0.5289         0.6394
   gate_proj  128   64   123.09   136.30   110.26    85.65    71.53     85      0.5248         0.6284

   down_proj   64    8    55.29    55.22    89.85   107.80   140.59     98      1.6271         1.6271
   down_proj   64   16    59.52    55.69    90.09   108.11   140.28     98      1.6177         1.6177
   down_proj   64   24    65.76    65.74    90.11   112.07   140.51     98      1.3707         1.3707
   down_proj   64   30    66.88    66.88    90.74   112.36   140.60     98      1.3568         1.3568
   down_proj   64   32    69.62    67.39    90.35   112.51   140.72     98      1.3407         1.3407
   down_proj   64   36   103.62   127.37    92.07   112.81   141.58     98      0.7229         0.7229
   down_proj   64   40   104.47   127.56    92.19   112.79   140.86     98      0.7227         0.7227
   down_proj   64   42   104.44   127.79    92.48   113.56   140.91     98      0.7237         0.7237
   down_proj   64   48   106.05   129.24    92.55   114.87   141.01     98      0.7161         0.7161
   down_proj   64   56   121.48   129.50   143.65   112.74   141.14     94      0.8706         0.8706
   down_proj   64   64   122.91   131.27   144.08   113.33   141.83     94      0.8633         0.8633

    out_proj   64    8    23.54    23.04    34.57    40.78    52.32     98      1.5004         1.5004
    out_proj   64   16    23.99    23.99    34.72    41.17    52.45     98      1.4473         1.4473
    out_proj   64   24    27.96    27.96    34.62    42.64    52.60     98      1.2382         1.2382
    out_proj   64   30    28.77    28.76    34.90    42.69    52.75     98      1.2135         1.2135
    out_proj   64   32    28.91    28.91    34.73    42.73    52.78     98      1.2013         1.2013
    out_proj   64   36    42.22    53.01    35.13    43.06    53.15     98      0.6627         0.6627
    out_proj   64   40    42.63    53.12    35.02    42.99    52.98     98      0.6593         0.6593
    out_proj   64   42    42.76    53.39    35.07    43.33    53.16     98      0.6569         0.6569
    out_proj   64   48    44.27    54.21    35.12    42.98    53.12     98      0.6479         0.6479
    out_proj   64   56    49.83    55.12    52.29    43.09    53.27     94      0.7817         0.7817
    out_proj   64   64    50.54    55.95    52.35    43.03    53.76     94      0.7691         0.7691

 in_proj_qkv   48    8    34.20    33.83    34.56    36.07    45.21     98      1.0216         1.0216
 in_proj_qkv   48   16    35.98    34.34    34.72    36.35    45.32     98      1.0111         1.0111
 in_proj_qkv   48   24    40.00    39.94    44.79    37.21    45.46     94      0.9316         1.1214
 in_proj_qkv   48   30    40.85    40.66    45.29    37.37    45.65     94      0.9191         1.1139
 in_proj_qkv   48   32    41.98    40.88    44.79    37.47    45.95     94      0.9166         1.0956
 in_proj_qkv   48   36    61.87    76.81    48.07    53.83    46.50     85      0.6054         0.6258
 in_proj_qkv   48   40    63.19    76.67    47.98    54.26    46.00     85      0.6000         0.6258
 in_proj_qkv   48   42    63.32    78.11    48.04    57.01    46.51     85      0.5954         0.6150
 in_proj_qkv   48   48    63.60    78.42    48.35    54.46    46.28     85      0.5902         0.6166
 in_proj_qkv   48   56    72.82    78.65    67.01    54.02    46.05     85      0.5855         0.6868
 in_proj_qkv   48   64    76.35    78.61    67.34    54.91    47.13     85      0.5995         0.6985

   in_proj_z   48    8    22.87    22.81    29.87    35.10    44.60     98      1.3095         1.3095
   in_proj_z   48   16    23.63    23.63    30.15    35.41    44.79     98      1.2759         1.2759
   in_proj_z   48   24    27.27    27.24    30.16    36.46    44.89     98      1.1072         1.1072
   in_proj_z   48   30    27.87    27.87    30.21    36.67    45.03     98      1.0840         1.0840
   in_proj_z   48   32    28.11    28.12    30.38    36.61    45.07     98      1.0804         1.0804
   in_proj_z   48   36    41.21    51.54    44.93    36.81    45.17     94      0.7142         0.8718
   in_proj_z   48   40    41.57    51.76    44.91    36.89    45.23     94      0.7127         0.8677
   in_proj_z   48   42    41.70    52.03    44.96    36.89    45.36     94      0.7090         0.8641
   in_proj_z   48   48    42.26    52.27    45.11    37.00    45.40     94      0.7079         0.8630
   in_proj_z   48   56    48.71    53.47    45.01    37.19    45.90     94      0.6955         0.6955
   in_proj_z   48   64    49.31    54.28    44.98    37.20    45.95     94      0.6853         0.6853

      k_proj   32    8    12.40    16.86    28.66    34.13    43.60     98      1.6999         1.6999
      k_proj   32   16    15.25    19.81    28.99    34.38    43.70     98      1.4634         1.4634
      k_proj   32   24    17.49    22.32    29.20    35.86    43.93     98      1.3082         1.3082
      k_proj   32   30    18.78    24.62    29.34    35.94    44.10     98      1.1917         1.1917
      k_proj   32   32    19.03    25.53    29.32    35.99    44.13     98      1.1485         1.1485
      k_proj   32   36    24.20    41.14    29.24    35.92    44.33     98      0.7107         0.7107
      k_proj   32   40    25.06    42.24    29.29    36.10    44.39     98      0.6934         0.6934
      k_proj   32   42    25.22    42.01    29.27    36.00    44.41     98      0.6967         0.6967
      k_proj   32   48    26.29    43.59    29.35    36.14    44.53     98      0.6733         0.6733
      k_proj   32   56    29.26    46.01    29.45    36.03    44.68     98      0.6401         0.7831
      k_proj   32   64    30.04    48.77    29.43    36.05    45.02     98      0.6034         0.7392

      q_proj   16    8    39.60    39.61    39.88    40.27    45.79     98      1.0068         1.0068
      q_proj   16   16    40.91    40.92    39.76    40.16    45.62     98      0.9717         0.9717
      q_proj   16   24    46.77    68.61    45.66    40.64    46.42     94      0.5923         0.6655
      q_proj   16   30    47.61    69.83    46.10    40.55    45.93     94      0.5807         0.6602
      q_proj   16   32    47.80    70.72    46.09    40.82    46.35     94      0.5772         0.6517
      q_proj   16   36    74.88    98.06    67.67    54.77    47.05     85      0.4798         0.6901
      q_proj   16   40    76.06    97.55    65.34    55.15    47.15     85      0.4833         0.6698
      q_proj   16   42    73.69    99.06    65.38    57.43    47.68     85      0.4813         0.6600
      q_proj   16   48    74.59   102.43    65.90    55.07    47.33     85      0.4621         0.6434
      q_proj   16   56    87.53   128.23    71.26    55.98    47.56     85      0.3709         0.4366
      q_proj   16   64    91.85   131.11    71.57    56.23    48.96     85      0.3734         0.4289

WHOLE-MODEL verify weight GEMM, ms per verify forward (400 linears), kernel-level sum
   M  B(K=7)  B(K=5)  legA unpinned  legAp SERVED  int8 shipped  int8 per-geom  shipped/served  pergeom/served  pergeom/shipped  cov%
   8     1.0       -         15.548        15.946        19.461         19.461          1.2205          1.2205           1.0000   100
  16     2.0       -         16.214        16.700        19.554         19.554          1.1709          1.1709           1.0000   100
  24     3.0     4.0         18.750        19.245        21.497         19.914          1.1171          1.0348           0.9263   100
  30       -     5.0         19.404        19.907        21.661         19.996          1.0881          1.0045           0.9231   100
  32     4.0       -         19.423        20.245        21.649         20.064          1.0693          0.9910           0.9268   100
  36       -     6.0         29.327        36.130        25.938         22.619          0.7179          0.6261           0.8720   100
  40     5.0       -         29.758        36.333        26.281         22.640          0.7233          0.6231           0.8615   100
  42       -     7.0         29.496        36.811        26.164         22.894          0.7108          0.6219           0.8750   100
  48     6.0     8.0         30.042        37.628        26.148         22.816          0.6949          0.6064           0.8726   100
  56     7.0       -         34.626        38.595        27.215         24.618          0.7051          0.6378           0.9046   100
  64     8.0       -         35.319        39.466        27.445         24.936          0.6954          0.6318           0.9086   100

READING THE WHOLE-MODEL TABLE
  It is a kernel-level SUM of per-linear medians weighted by the number of
  linears of each geometry in the verify forward (400 exl3 linears; the
  lm_head is eager and outside the verify graph).  It is NOT a served step.
  What licenses reading it is that its M = 8 column, 15.95 ms, reproduces the
  16.76 ms of verify-forward weight GEMM the served trace measured at c = 1
  (experiments/int8_served_ab/decode_decomp/RESULTS.md, section 2b) to within
  4.8% -- two instruments, one quantity.

  * M = 8 (B = 1, K = 7) costs 1.00x a plain decode step's weight GEMM.  At
    the single-stream operating point the whole slate is free, and the verify
    weight GEMM is at the memory floor.  There is nothing to win there.
  * The step from M = 32 to M = 36 is 1.78x (20.25 -> 36.13 ms).  That is the
    second trellis pass, and it is the largest single number anywhere in the
    verify path.
  * The int8 leg with the SHIPPED band pin already recovers most of it:
    0.69-0.72x of the served leg at M >= 36.  The old receipt reported
    0.85-0.92x for the same arms because it divided by the unpinned leg.
  * Per-geometry tiles take it to 0.61-0.64x -- a further 1.10-1.16x over the
    band pin, for a recorded field and no arithmetic change.
  * Below the threshold the leg is correctly refused: 1.17-1.22x WORSE at
    M = 8 and 16, and a wash (0.99-1.03x) at 24-32.  `auto` is right to
    decline those, and lowering the threshold below 36 buys nothing.

WHAT IS RECORDED, AND WHAT EACH ENTRY IS THE WINNER AT
  shipped_pins.json, entry `NVIDIA GeForce RTX 4090` / select_rows [1,16,32]:

    5120x17408xK4  gate_proj + up_proj  128 linears  int8_verify_shape 85
    5120x10240xK4  in_proj_qkv           48 linears  int8_verify_shape 85
    5120x12288xK4  q_proj                16 linears  int8_verify_shape 85
    5120x6144xK4   in_proj_z             48 linears  int8_verify_shape 94
    5120x1024xK4   k_proj + v_proj       32 linears  int8_verify_shape 98

  Each is the fastest tile at EVERY width `min_rows = 40` admits (40, 42, 48,
  56, 64) and also at 36, so one tile per geometry is sufficient where the
  band ladder is not.  `6144x5120xK4` (out_proj + o_proj) and `17408x5120xK4`
  (down_proj) genuinely alternate -- s98 to 48 rows, s94 above -- and are left
  unrecorded so they keep the band.

  A recorded tile is NOT the winner below the threshold: s85 loses to s94 on
  gate_proj at 24-32 rows.  A deployment that lowers `min_rows` under 36 must
  re-read this table before trusting these entries.

STILL OPEN
  * `exl3_int8_verify_min_rows` 40 -> 36.  Measured here as a 1.4-1.6x win on
    the width it admits; owes an acceptance A/B on a box that can boot the 27B
    before it moves, because it routes a slate onto different arithmetic.
  * Nothing measures the SERVED distribution of slate widths.  The band edge,
    the threshold and the family pin are all decisions about which widths
    matter, and the only instrument that records them
    (`_dbg_verify_capture_hist`, `/v1/admin/capture_hist`) is behind
    `ARBI_DEBUG_CAPTURE_LOOKUP` and off in every shipped compose.

REPRODUCING
  rsync tools/ to the box, then per geometry (rot 6 for the small ones, 2 for
  gate_proj/down_proj), inside the image with TORCH_EXTENSIONS_DIR warm:

    python bench_group_scale.py --rows 8 16 24 30 32 36 40 42 48 56 64 \
      --shapes 98 94 85 --linears gate_proj --layer 0 --rot 2 --reps 7 \
      --min-burst-ms 30 --lega-shape 10

  `--lega-shape` MUST name the member of the geometry's shipped family that
  `_pin_for_rows` picks at those rows, or `legAp` is a leg nothing serves and
  the ratios are the old receipt's again.
