=== the residual was the OPERAND boundary, not the output cast (#1861) ===
10.2.0.206 GPU 0, sole tenant, registry.arbi.work/arbi-serve:test-latest,
Qwen3.8-27B-exl3-4.0bpw, shape 27, 6 prompts x 16384 tok, --ignore-eos.

--- 1. the instrument had to be fixed before the accounting meant anything ---

(a) UNEQUAL WORK. Greedy decode stops on EOS and the arms diverge in numerics,
    so the profiled prompt ran to "stop" on INT8F and to "length" on INT8P --
    in the previous receipt and in the first two reproductions here. That
    removes most of the decode phase from one side of the diff: leg A's
    exl3_gemm_kernel read n=8 on INT8F against n=3216 on INT8P, which is not a
    2.5-order-of-magnitude regression, it is one arm not having decoded.
    Fixed with --ignore-eos plus a refusal that compares generated-token counts
    across arms and voids the tables when they differ.

(b) NAME MATCHING. Inductor numbers kernels per graph and folds the SURROUNDING
    origin nodes into the name, so the same work is called ..._mul_silu_10 in
    one arm and ..._mlp_down_10 in the other and a name-matched diff reports
    renames as deltas. Kernels are now bucketed on the symbol that names the
    WORK (a_prep_kernel<5, exl3_i8_gemm_kernel, had_hf_r_128, ...).

(c) A DENOMINATOR THAT IS COUNTED IN THE SAME WINDOW, in eager code: every MLP
    tail enters the op (large-M fuses, small-M takes the shipped leg inside the
    same op), so fused+smallM is the exact tail count for the profiled prompt.
    Counting it in the PATCHED FORWARD instead cost one boot: a global-dict
    mutation inside the fullgraph region installs a value guard, recompiles
    every call, and takes down the readiness gate as FailOnRecompileLimitHit.

--- 2. the 2x was real. it is not an output cast. ---

Per-kernel, ONE prompt, both arms verified to have generated 8 tokens and
finished "length":

  MLP tails through the op, INT8P .............. 1024  (512 fused + 512 small-M)
  INT8F  MLP glue silu*mul ......... n=1024  = 1.00x / tail   108.27 ms
  INT8P  "..._mlp_down_*" .......... n=2048  = 2.00x / tail   155.14 ms

The generated Inductor code says what those 2048 launches are:

    triton_poi_fused__to_copy_mlp_down_10.run(buf31, buf35, ...)   # xnumel=17408
    triton_poi_fused__to_copy_mlp_down_10.run(buf34, buf36, ...)   # xnumel=17408
    buf37 = torch.ops.prologue_e2e.mlp_down.default(buf35, buf36, ...)

They run BEFORE the op and produce its two INPUTS, at the intermediate width
K=17408. They are `gate` and `up` being materialised in the engine dtype. The
name says mlp_down because Inductor names a kernel after the node DOWNSTREAM of
it, which is why they read as an output cast that mysteriously runs twice.

WHY THE SHIPPED ARM DOES NOT PAY THEM. EXL3LinearBase._exl3_forward ends in
y.to(orig_dtype), so gate/up are bf16 as graph VALUES; their only consumer in
the shipped MLP is the glue, and Inductor folds both conversions into it. The
bf16 tensors are never written. An opaque op where the glue was has nothing to
fuse into, so both conversions become real passes over the widest tensor in the
block: 2 operands x 4 B/elem = 8.91 MB/token, against the 4.456 MB/token the
prologue saves by not materialising silu(gate)*up. The prologue wins its
traffic and hands back twice as much at the boundary.

--- 3. the fix: hand the op the projections' OWN fp16 ---

_gemm_f16() is _exl3_forward minus its trailing y.to(orig_dtype). The quant
kernels emit fp16 and the prep wants fp16, so no conversion exists to fuse or
to pay. The single output cast stays in the traced graph, now spelled off x
(the engine dtype) rather than off gate, which is fp16 here.

Refuses rather than serving something else: a projection carrying a dense bias
or Hadamard pad/trim falls back and is counted (f16_operands=0 on this run).

SENTINEL, not a neutral value: the A-prep's own template instantiation reports
which dtype it actually read.

  before   a_prep_kernel<5, 32, true, SrcSiluMul<__nv_bfloat16> >   n=512  121.58 ms
  after    a_prep_kernel<5, 32, true, SrcSiluMul<__half> >          n=512  110.89 ms

--- 4. device time by bucket, equal work, same prompt ---

bucket                    INT8F        INT8P (defect)   INT8P (fixed)
int8 GEMM (leg B sub)   1713.63/3200   1711.99/3200    1718.27/3200
other                    647.81/16061   651.81/19645    656.53/19133
output Hadamard          243.87/3200    243.59/3200     245.63/3200
other triton             209.24/9984    194.71/9984     178.65/9984
leg A GEMM (shipped)     144.88/3216    145.35/3216     145.37/3216
other A-prep             109.94/2688    109.09/2688     109.00/2688
MLP glue silu*mul        108.29/1024      0.00/0          0.00/0
operand casts (mlp_down)   0.00/0       155.14/2048       0.00/0
down_proj A-prep          61.66/512     121.58/512      110.89/512
                        ------------   -------------   -------------
total                    3239.32        3333.26         3164.34
vs INT8F                                +94.80 (+2.9%)  -74.98 (-2.3%)

The down_proj site alone, fixed: glue -108.29, A-prep +49.23, casts 0
                              = -59.06 ms, i.e. the 4.456 MB/token it was
                                supposed to win, delivered.

--- 5. end to end, with a null control ---

medians over prompts 1-5 (prompt 0 is compile-bearing and excluded):

  REF    4.910   [4.89, 5.04, 4.90, 4.91, 4.91]
  REF2   4.910   [4.89, 4.90, 4.91, 4.92, 4.91]
  INT8F  3.400   [3.39, 3.39, 3.40, 3.70, 3.41]
  INT8P  3.340   [3.35, 3.34, 3.34, 3.49, 3.34]

  NULL CONTROL  REF2 vs REF ....... +0.00%   (two identical arms)
  INT8P vs INT8F ................... -1.76%   (1.0180x)
  REF/INT8F 1.444x    REF/INT8P 1.470x

The same harness measured this arm at +3.1% (previous receipt), +4.1%, +3.8%
and +3.9% before the fix. It is now a win, and the null control says the
harness resolves 0.00% between two arms that are the same code.

Accuracy, same run (first-prefill-position KL vs the bf16 reference):

  REF    0.000e+00   REF2 0.000e+00     (null control: bit-identical)
  INT8F  4.627e-04   median 7.072e-05   max 2.456e-03
  INT8P  6.492e-04   median 5.765e-05   max 3.620e-03

Same band. INT8P no longer rounds gate/up through bf16 before the prep, which
is what the shipped fused glue does too (emulate_precision_casts is off, so
Inductor keeps that chain in fp32 and never applies the bf16 rounding).

--- 6. confirmed on the merged tree (this PR's base + this PR) ---

Re-run on fix/exl3-prologue-double-cast after merging feat/exl3-int8-gemm, so
the numbers belong to the code that lands, not to the branch point. All four
arms verified to have generated 8 tokens and finished "length".

  REF    4.920   [4.89, 4.90, 4.92, 4.92, 4.92]
  REF2   4.920   [4.90, 4.91, 4.92, 4.92, 4.92]
  INT8F  3.410   [3.41, 3.41, 3.41, 3.42, 3.42]
  INT8P  3.360   [3.36, 3.36, 3.36, 3.36, 3.38]

  NULL CONTROL  REF2 vs REF ....... +0.00%
  INT8P vs INT8F ................... -1.47%   (1.0149x)
  REF/INT8F 1.443x    REF/INT8P 1.464x

  prologue calls=3136  fallbacks: not_serveable=0 lora=0 pad=0 act=0
                                  noncontig=0 f16_operands=0
  small-M on the shipped leg A=3200   dynamo unique graphs=14

  bucket                    INT8F        INT8P
  int8 GEMM (leg B sub)   1706.85/3200  1712.83/3200
  other                    648.29/16061  656.90/19133
  output Hadamard          244.29/3200   245.64/3200
  other triton             209.43/9984   178.42/9984
  leg A GEMM (shipped)     144.91/3216   145.38/3216
  other A-prep             109.57/2688   109.00/2688
  MLP glue silu*mul        108.27/1024     0.00/0
  operand casts (mlp_down)   0.00/0        0.00/0
  down_proj A-prep          61.69/512    110.86/512

--- 7. and again on the exact shipping commit ---

  REF    4.910   [4.91, 4.91, 4.91, 4.93, 4.94]
  REF2   4.920   [4.91, 4.92, 4.92, 4.93, 4.95]
  INT8F  3.420   [3.41, 3.42, 3.42, 3.52, 3.44]
  INT8P  3.370   [3.36, 3.37, 3.37, 3.37, 3.40]

  NULL CONTROL  REF2 vs REF ....... +0.20%
  INT8P vs INT8F ................... -1.46%   (1.0148x)

  a_prep_kernel<5, 32, true, SrcSiluMul<__half> >   n=512   110.88 ms
  kernels matching "mlp_down": 0
  fallbacks: not_serveable=0 lora=0 pad=0 act=0 noncontig=0 f16_operands=0

And once more on the branch tip, after the TP refusal was added to _install:

  REF    4.910   [4.89, 4.90, 4.91, 4.92, 4.92]
  REF2   4.910   [4.90, 4.90, 4.91, 4.92, 4.92]
  INT8F  3.410   [3.40, 3.41, 3.41, 3.42, 3.42]
  INT8P  3.360   [3.36, 3.36, 3.38, 3.36, 3.37]

  NULL CONTROL  REF2 vs REF ....... +0.00%
  INT8P vs INT8F ................... -1.47%   (1.0149x)
  kernels matching "mlp_down": 0     fallbacks all 0

Four runs: -1.76%, -1.47%, -1.46%, -1.47%. The null control read +0.00%,
+0.00%, +0.20% and +0.00%, so the harness's floor between two arms that are the
same code is ~0.2% and the win is about seven times it. Quoted as a floor, not
as zero.

--- 8. the defect reproduces in seconds, without a checkpoint ---

tools/prologue_fusion/probe_output_cast.py, triton launches per op call:

  BASE      3.00x  expected 3   glue absorbs both operand conversions
  P16_BF16  4.00x  expected 4   triton_poi_fused__to_copy_down_1 at 2.00x/call
  P16_F16   2.00x  expected 2   no conversion exists

The middle arm reproduces both the launch count AND the misleading name.
