test(MODEL-MM-QWEN4-EXP): the CUDA router tie-break is stable, and the geometry that proves it had never been run

[localai-org-maint-bot]

test(MODEL-MM-QWEN4-EXP): the CUDA router tie-break is stable, and the geometry that proves it had never been run

Wave MOEDIV (#2552) measured that 32.9% of this model's routing boundaries are
exact bf16 ties, and concluded that the CPU/CUDA expert flips are tie-break
order rather than a defect. It verified that conclusion BY SOURCE READ of both
kernels. A source read establishes the INTENDED tie-break; it does not establish
that a parallel argmax realises it. This wave tests the part a source read
cannot reach.

The answer

The CUDA top-k realises the lowest-index tie-break it declares. On
bit-identical logits whose top-k boundary is an EXACT bf16 tie, at E = 512
k = 10 -- the geometry qwen4_exp routes -- the dispatched kernel returns the
closed-form correct selection, agrees byte-for-byte with the serial GPU oracle
and with the CPU reference, repeats identically over 32 launches, and gives the
same answer for a row run alone and as each of 257 rows of a batch.

Measured on thor:gpu0 (job 5f50e5f7-c730-463f-93cc-3293835ed007, pod
rc-worker-n8smh, NVIDIA Thor, driver 595.78, cc 11.0, nvcc cuda_13.0.r13.0,
sm_110): CMAKE rc=0, BUILD rc=0 objects=586, 41 *.cu.o, libcudart.so.13
and libcublasLt.so.13 in the test binary's own ldd, and
test_moe_router_tie_stability reading 4 cases | 4652 assertions | 0 failed |
rc=0.

So MOEDIV's reading stands, the remaining CPU-vs-CUDA token divergence stays
attributed to arithmetic, and the session goal's three disagreeing ids do not
move -- because nothing moved them. No product code changed. No fix was made,
because none was needed.

What the wave found anyway

The tree could not have answered this question before. MoeRouterWarpValuesPerThread
returns 0 for E = 512, so qwen4_exp routes through MoeRouterTopKKernel<Tin,false>
at 2 experts per thread, and the standing adversarial sweep in
tests/vt/test_ops_moe_grouped.cpp stops at E = 256 -- where a thread holds ONE
expert and the per-thread strided compare can never see a tie at all.

Flipping that compare to highest-index-wins (cuda_moe.cu, ONE site, asserted
n == 1, MUT BUILD rc=0 objects=1) is caught by 104 of the new assertions
and by ZERO assertions of the standing sweep, whose counts are byte-identical
before and after: test_ops_moe_grouped read 1907 | 1906 passed | 1 failed
both times, that one failure being the pre-existing #962 Marlin NVFP4 block-size
disagreement on this arch, and test_ops_moe read 33451 | 33451 both times.
Restoring returned 4652/4652.

geometry	pattern	par == r.expect failures under the mutation
E = 512	3, 5	12, 12
E = 1024	3, 5	12, 12
E = 256	any	0

Zero at E = 256 is the finding: the compare the mutation breaks does not exist
below E > kBlock.

Two negatives kept rather than dropped

The mutant stayed perfectly deterministic. All 279 repeat comparisons and
2313 batch-position comparisons passed with the wrong tie-break compiled in. A
repeatability check alone would have called it clean, so determinism is recorded
as necessary and never sufficient.

A corroborating orin:gpu0 lease registered no CUDA backend, skipped every
device case, and read 488/488 GREEN with the broken tie-break in the binary.
The assertion count is what distinguishes it from Thor -- 488 against 4652 --
but the assertion count ALONE separates only two of three states, because a
CPU-only build reads 488 as well. The CASE count is the other half:

state	cases	assertions	rc
CPU-only build -- the device cases are #ifdef'd out	1	488	0
CUDA build, no device -- they skip and return	4	488	0
CUDA build with a device	4	4652	0

Every table here therefore carries both counts and never a status line, and the
spec's Gates section now teaches both.

How it is asserted

Every row has a closed-form correct answer -- h strict winners, then a tied set
S sharing one bf16 bit pattern, then a low floor -- so the top-k is the h winners
followed by the k-h LOWEST INDICES OF S. That is upstream's own rule at the
parity pin 5559679229, not a local convention:
csrc/libtorch_stable/moe/topk_softmax_kernels.cu:536-537 ("We want lower
indices to "win" in every thread so we break ties this way"), :515-517
("only updated if > (not >=)"), :707-708 (case 512: LAUNCH_TOPK(512, ...),
so E = 512 is a registered topkGating width upstream), :186,222,225
(cub::ArgMax in the fallback) and :465. The gate therefore measures the
semantics and cannot pass by both arms being wrong the same way.

Indices are compared with == and the selection as an ordered vector; there is
no tolerance anywhere in the file, because a discrete selection has bimodal
error. The counted property is tied_seen -- rows whose boundary carries the
same bf16 bit pattern on both sides, read off the LOGITS rather than any kernel
-- CHECKed equal to the row count, so a builder that stopped producing ties
cannot read green.

Red-first on the CPU half, on the dev box, CPU-only build at this branch head:
flipping the greedy argmax's strict > to >= at ONE site --
src/vt/cpu/cpu_ops.cpp:2985, grep -c = 1 -- turns every selection
assertion in the CPU case red. Baseline 1 case | 488 | 0 failed | rc 0; mutant
1 case | 488 | 380 passed | 108 failed | rc 1, all 108 of them
CHECK( ids == r.expect ) and none of them the counted property; restored
488/488. 108 is the whole population and not a subset: h < k is REQUIREd,
so each of the 54 rows selects at least one member of the tied set, and >=
mis-selects on every row in both dtypes -- 54 x 2 = 108.

(An earlier revision of this body said 90. That figure came from a
five-pattern draft of the file -- 5 x 3 x 3 x 2 -- and the shipped file has
six patterns. It was re-measured rather than carried forward, because
squash_merge_commit_message = PR_BODY makes this text the commit message and a
wrong measured number in it cannot be repaired afterwards.)

Owed

Two gaps the fresh review found are RECORDED here rather than fixed, each
with an issue and a ## Owed entry in the row's spec, because neither belongs
in a test-only change.

• #2603 -- the permissive
  HasCuda() skip lets a device-less CUDA build read green. That is the
  orin:gpu0 row above, and it is a tree-wide shape:
  grep -rl 'no CUDA backend registered; skipping' tests/ finds 14 files, 13 of
  them not this one. The tree already carries the idiom that refuses instead
  (tests/parity/test_qwen27n_fp8_tower_paged_engine.cpp:149 reads
  VT_REQUIRE_27N_FP8_GATE and turns absence into a FAILURE), so adopting it is
  a row of its own.
• #2604 -- half the device
  sweep is a duplicate above E = 256. MoeRouterWarpValuesPerThread returns 0
  for every E outside {32,64,128,256} and LaunchRouterWarp
  (cuda_moe.cu:564) returns false on vpt == 0, so at E = 512 and E = 1024
  -- qwen4_exp's own geometry -- both VT_MOE_ROUTER_WARP arms dispatch the
  same block kernel. The counts stay honest, because every assertion really
  executes; but 36 of the 72 rows at E > 256 are byte-identical repeats, the
  12 per cell above is 2 arms x 3 h x 2 dtypes of which 6 are a second
  count of one disagreement, and COVERAGE at E > 256 is half what the doubled
  figures imply. Only E = 256 compares two structures.

The gate against vLLM itself stays owed for the reason the row's
spec already records under ## Owed: it needs an artifact vLLM can load, and
vLLM's GGUF support is an out-of-tree plugin while every safetensors arm of this
model exceeds the largest fleet box. This change says our two arms agree with
each other AND with upstream's stated rule; it does not run upstream's kernel on
upstream's weights.

Evidence: docs/bench-evidence/qwen4exp-moe-tiebreak-stability-20260902.md.
Spec: .agents/specs/qwen4-exp-flash-next.md "## Wave TIEBREAK".

Closes #2586

FOLLOWING_AGENTS_PROTOCOL

Following-Agents-Protocol: true
AI-Assisted: true
Assisted-by: AGENT:claude-opus-5 [claude-code]