CollectiveX: add FlashInfer one-sided NVLink EP for the GB SKUs / CollectiveX：为 GB SKU 新增 FlashInfer 单边 NVLink EP 后端

[Oseltamivir]

Why

GB200 and GB300 are an NVLink-only domain, and a deployment picks its all-to-all by fabric. vLLM's own Kimi-K3 recipe is explicit: --all2all-backend deepep_v2 for RDMA, flashinfer_nvlink_one_sided for NVLink.

Our GB rows carried only deepep-v2 and nccl-ep — two scale-out-heritage transports pointed at MNNVL. So every GB number published so far describes a transport nobody would select on that rack. This adds the one they would.

One-sided, never two-sided. The sibling flashinfer_nvlink_two_sided produces gibberish output on GB200/arm64 (vllm#39722), so it is deliberately not wired. One-sided means the initiator writes straight into the target's workspace and then sets a flag; on MNNVL, where peer memory is directly addressable across the 72-GPU domain, that collapses to ordinary stores over the fabric with no rendezvous — which is the whole reason it wins on this topology.

The adapter

Three properties of flashinfer.comm.trtllm_moe_alltoall.MoeAlltoAll shape it:

upstream behaviour	consequence
dispatch asserts "called twice without combine"; combine asserts phase dispatched; state resets after combine	combine_needs_redispatch = True, dispatch_needs_combine_cleanup = True — the contract MoRI and DeepEP V2 LL already declare
receive is padded [ep_size, runtime_max_tokens_per_rank, hidden], not compact	both oracle views read valid slots only; padding is untouched workspace memory
runtime_max_tokens_per_rank <= max_num_tokens	workspace sized once from the ladder maximum, each call passes its rung

normal mode only — deliberate, not a first slice. FlashInfer exposes one one-sided A2A kernel family, not a separate decode-optimized one. An ll_backends cell would re-measure the same kernel under a mode that promises a different one. Decode is still covered: normal runs the full decode and prefill ladders. (This differs from the plan of "normal + LL" — the reason surfaced while reading the upstream API, and inventing an LL row would have been dishonest labelling rather than added coverage.)

BF16 only for the same kind of reason: the combine side accepts FP8 output dtypes, but an FP8 dispatch needs the scale payload plumbed as a second input_payload and validated against the oracle's cast round-trip.

No build seam. The pinned SGLang images ship flashinfer-python and the one-sided A2A lives in that wheel, so flashinfer_ep_prepare is a capability assert, not an install — it fails loudly at prepare time if an image drops the module instead of dying mid-case in create_buffer.

Wiring

bench/ep_flashinfer.py (new) · bench/run_ep.py (choices, selection, group init) · sweep_matrix.py (BACKEND_PRECISIONS) · launchers/launch_gb-nv.sh (backend gate) · runtime/prepare_backend.sh (prepare + dispatch) · configs/platform_config.json (gb200/gb300 [8, 16]) · .github/workflows/collectivex-sweep.yml (dispatch option) · tests/test_matrix.py (rollout shape).

Status — read before merging

This has not run on GB hardware. The adapter is written against the verified upstream API and vLLM's reference integration (FlashInferNVLinkOneSidedManager), and the matrix/registry/launcher wiring is exercised by tests — but no leg has executed on gb200 or gb300.

The open item bring-up must settle first is combine_weight_semantics. Upstream source does not state whether combine applies the top-k scales: they ride along only as a caller-supplied dispatch payload, and vLLM applies them in the MoE layer rather than the A2A. It is declared unweighted-rank-sum here and must be confirmed against the oracle before any published run — wrong semantics reds every case.

Also note this PR adds a backend while #2449 adds backend_maturity with a test asserting every dispatched backend has an entry. Whichever merges second needs "flashinfer-ep": "production" added to that map; the adapter already declares maturity = "production" for the artifact side.

Verification

• 73 CollectiveX tests pass (1 new: rollout shape — GB-only, EP8+EP16, bf16-only, no LL cell on any SKU).
• py_compile clean; bash -n clean on the launcher and prepare script; workflow YAML parses with flashinfer-ep in the dispatch options.

---

中文说明

为什么

GB200 与 GB300 是纯 NVLink 域，而真实部署按 fabric 选择 all-to-all。vLLM 官方 Kimi-K3 recipe 写得很明确：RDMA 用 --all2all-backend deepep_v2，NVLink 用 flashinfer_nvlink_one_sided。

我们的 GB 行此前只有 deepep-v2 与 nccl-ep —— 两个面向 scale-out 的传输实现被指向了 MNNVL。因此迄今发布的每一个 GB 数据，描述的都是没有人会在该机架上选用的传输实现。本 PR 补上真正会被选用的那一个。

只接单边，绝不接双边。 姊妹后端 flashinfer_nvlink_two_sided 在 GB200/arm64 上会输出乱码（vllm#39722），故有意不接入。单边意味着发起方直接写入目标 workspace 再置标志位；在 MNNVL 上对端显存于 72-GPU 域内可直接寻址，于是它退化为经 fabric 的普通 store，完全没有 rendezvous —— 这正是它在该拓扑上取胜的根本原因。

adapter 设计

flashinfer.comm.trtllm_moe_alltoall.MoeAlltoAll 的三个特性决定了它的形态：

上游行为	由此产生的约束
dispatch 断言「未 combine 就二次 dispatch」；combine 断言处于 dispatched 阶段；combine 后状态复位	combine_needs_redispatch = True、dispatch_needs_combine_cleanup = True —— 与 MoRI、DeepEP V2 LL 已声明的契约一致
接收缓冲是带 padding 的 [ep_size, runtime_max_tokens_per_rank, hidden]，而非紧凑布局	两个 oracle 视图都只读有效槽位；padding 是未初始化的 workspace 内存
runtime_max_tokens_per_rank <= max_num_tokens	workspace 按 ladder 最大值一次性分配，每次调用传入当前档位

仅 normal 模式 —— 这是刻意决定，而非「先做一部分」。 FlashInfer 只暴露一个单边 A2A kernel 家族，没有单独的解码优化 kernel。新增 ll_backends 单元格只会在一个承诺「不同 kernel」的模式下重复测量同一个 kernel。解码仍有覆盖：normal 会跑完整的 decode 与 prefill ladder。（这与原计划的「normal + LL」不同：该事实是在阅读上游 API 时才浮现的，硬造一个 LL 行属于误导性标注，而非真正增加覆盖。）

BF16-only 出于同类原因：combine 侧支持 FP8 输出 dtype，但 FP8 dispatch 需要把 scale 负载作为第二个 input_payload 接入，并与 oracle 的 cast round-trip 对齐验证。

无需构建步骤。 固定的 SGLang 镜像自带 flashinfer-python，单边 A2A 就在同一 wheel 中，因此 flashinfer_ep_prepare 是能力断言而非安装 —— 一旦镜像不再提供该模块，它会在 prepare 阶段大声失败，而不是在 create_buffer 中途崩溃。

状态 —— 合入前请先阅读

尚未在 GB 硬件上运行。 adapter 依据已核实的上游 API 与 vLLM 参考实现（FlashInferNVLinkOneSidedManager）编写，矩阵/注册表/launcher 接线也有测试覆盖，但还没有任何 leg 在 gb200 或 gb300 上真正执行。

bring-up 首先要解决的未决项是 combine_weight_semantics。上游源码没有说明 combine 是否会乘以 top-k 权重：这些权重仅作为调用方提供的 dispatch 负载随行，而 vLLM 是在 MoE 层施加它们。此处声明为 unweighted-rank-sum，必须在任何正式发布运行之前用 oracle 确认 —— 语义错误会让所有 case 变红。

另请注意：本 PR 新增了一个后端，而 #2449 新增了 backend_maturity 及「每个可派发后端都必须有对应条目」的测试。后合入的一方需要在该映射中补上 "flashinfer-ep": "production"；adapter 侧已声明 maturity = "production"，产物一侧不受影响。

验证

• CollectiveX 73 项测试通过（新增 1 项 rollout shape：仅 GB、EP8+EP16、仅 bf16、任何 SKU 上都没有 LL 单元格）。
• py_compile 通过；launcher 与 prepare 脚本 bash -n 通过；workflow YAML 可解析且 dispatch 选项中已包含 flashinfer-ep。

[Oseltamivir]

On-metal probe result from GB200 (watchtower-navy-cn11, pinned image lmsysorg/sglang:v0.5.11-cu130 arm64, via srun + enroot):

python: 3.12.3
torch: 2.11.0+cu130 | cuda: 13.0
flashinfer: 0.6.8.post1
flashinfer.comm: OK
flashinfer.comm.mnnvl: OK
flashinfer.comm.trtllm_moe_alltoall: OK
init:     (self, mapping: Mapping, max_num_tokens: int, top_k: int, num_experts: int, workspace_size_per_rank: int = None, hidden_size: int = None …
dispatch: (self, token_selected_experts: torch.Tensor, input_payloads: list[torch.Tensor], runtime_max_tokens_per_rank: int, invalid_token_expert_id: Optional[int] = None …
combine:  (self, payload: torch.Tensor, runtime_max_tokens_per_rank: int, payload_in_workspace: bool = False) -> torch.Tensor

Confirms the no-build-seam claim: the image already ships flashinfer-python 0.6.8.post1 with the one-sided A2A module, so flashinfer_ep_prepare really is an assert rather than an install.

One thing worth flagging. The shipped 0.6.8.post1 combine is narrower than upstream main — no output=, output_dtype=, use_low_precision= or active_rank_mask=. The adapter calls combine(payload, runtime_max_tokens_per_rank) positionally and never passes output=, so it is compatible as written; this is the same version skew vLLM handles with its supports_kw(combine, "output") probe. If we later want the in-place output path, it needs a capability probe rather than a straight call.

Still outstanding, unchanged: no full leg has run, and combine_weight_semantics remains unconfirmed against the oracle.

中文：来自 GB200（watchtower-navy-cn11，固定镜像 lmsysorg/sglang:v0.5.11-cu130 arm64，经 srun + enroot）的实机探测结果见上。

这验证了「无需构建步骤」的判断：镜像已自带 flashinfer-python 0.6.8.post1 及单边 A2A 模块，因此 flashinfer_ep_prepare 确实只是能力断言而非安装步骤。

有一点需要指出：镜像中 0.6.8.post1 的 combine 签名比上游 main 更窄 —— 没有 output=、output_dtype=、use_low_precision= 与 active_rank_mask=。adapter 以位置参数调用 combine(payload, runtime_max_tokens_per_rank)，从不传 output=，因此按现状是兼容的；这与 vLLM 用 supports_kw(combine, "output") 探测处理的版本差异属于同一类问题。若日后要启用原地输出路径，需要加能力探测而不能直接调用。

仍未解决的事项不变：尚无完整 leg 运行，combine_weight_semantics 也仍未用 oracle 确认。

[claude]

🔴 flashinfer-ep's inspect_dispatch() returns a bare torch.Tensor instead of the types.SimpleNamespace(payload=, expert_ids=, weights=, local_expert_counts=) contract every other adapter (DeepEP V2, UCCL-EP, MoRI, NCCL EP) implements, and dispatch() never routes topk_idx/topk_weights through the transport as received metadata, so that namespace could not be built even if attempted. Because flashinfer-ep declares combine_weight_semantics="unweighted-rank-sum", ep_harness._run_expert_oracle takes the normal oracle path and immediately does view.payload/.expert_ids/.weights/.local_expert_counts on that bare tensor; the AttributeError is caught by a try/except in the harness (fail-soft), so every flashinfer-ep case will report a correctness FAILURE rather than crash, but the backend structurally cannot pass its correctness oracle as written.

Extended reasoning...

The bug. FlashInferEPBackend.inspect_dispatch (ep_flashinfer.py:154-160) returns h.recv_x[:, :h.tokens, :].reshape(-1, h.recv_x.shape[-1]) — a bare torch.Tensor. Every sibling adapter's inspect_dispatch instead returns types.SimpleNamespace(payload=..., expert_ids=..., weights=..., local_expert_counts=...): see ep_deepep_v2.py:423-430, ep_uccl.py:410-417, ep_mori.py:311-318, ep_nccl.py:449-456. This is a load-bearing contract, not a stylistic convention — ep_harness._run_expert_oracle is written entirely against that namespace shape.

The triggering code path. flashinfer-ep declares combine_weight_semantics = "unweighted-rank-sum" (ep_flashinfer.py, class body). In ep_harness.py, _run_expert_oracle branches on that flag at line 483: only "weighted-kernel-sum" backends take the short-circuit LL-style path, so "unweighted-rank-sum" (flashinfer-ep's declared value) falls through into the normal oracle path at lines 491+. That path does view = backend.inspect_dispatch(...) then immediately routing.decode_source_ids(view.payload, seed), and further down reads view.expert_ids (512), view.weights (513), and view.local_expert_counts (560). A bare Tensor has none of these attributes, so view.payload raises AttributeError on the very first access.

Why nothing upstream prevents this. dispatch() (ep_flashinfer.py:139-146) only forwards [p.dispatch_x] as the transport's input_payloads — it never sends p.topk_idx/p.topk_weights through self._a2a.dispatch(...) as received metadata alongside the activation. So even a corrected inspect_dispatch could not assemble expert_ids/weights from the receive buffer h as-is; the gap is structural in dispatch(), not just a return-type mismatch in inspect_dispatch(). The adapter's own comment claiming the unweighted-rank-sum semantics were "Verified against the oracle during bring-up" is inconsistent with this: the oracle path this backend takes cannot execute past the first view.payload access, so no such verification could have happened as described.

Impact. _run_expert_oracle's view = backend.inspect_dispatch(...) call sits inside a try/except Exception (ep_harness.py:490-507), so this does not crash the leg outright: the handler catches the AttributeError, drains the in-flight dispatch (recv_tokens/stage/combine + a synchronize) to avoid deadlocking peer ranks, and returns a fail-soft _oracle_report(passed=False, ...). The sweep then proceeds through the timed passes and produces latency/bandwidth numbers, but every flashinfer-ep correctness case reports passed=False — the adapter can never show a green correctness gate, on any shape, on any SKU, regardless of whether combine_weight_semantics (the PR's disclosed open item) is actually correct.

Step-by-step proof.

1. run_ep.py selects FlashInferEPBackend for --backend flashinfer-ep and calls ep_harness.run_sweep(...).
2. For each case, the harness calls backend.dispatch(p), which returns SimpleNamespace(recv_x=received[0], tokens=p.T, combine_input=None) — note: no expert-id/weight fields.
3. Because combine_weight_semantics == "unweighted-rank-sum" (not "weighted-kernel-sum"), _run_expert_oracle takes the normal-mode branch and calls view = backend.inspect_dispatch(p, h).
4. inspect_dispatch returns h.recv_x[:, :h.tokens, :].reshape(-1, hidden) — a plain Tensor.
5. The very next line, routing.decode_source_ids(view.payload, seed), raises AttributeError: 'Tensor' object has no attribute 'payload'.
6. The surrounding except Exception catches it, drains the dispatch, and returns passed=False. Net result for every flashinfer-ep case: correctness FAILURE, never a pass.

Fix. Route p.topk_idx/p.topk_weights through self._a2a.dispatch(...) as additional input_payloads (or otherwise capture them alongside recv_x) so inspect_dispatch can build types.SimpleNamespace(payload=..., expert_ids=..., weights=..., local_expert_counts=...) matching the other four adapters, applying the same padded-slice-not-mask handling already used for recv_x.

This is separate from the PR's disclosed open item (verifying combine_weight_semantics against the oracle) — it is a structural type/plumbing defect that guarantees oracle failure independent of whether that semantics declaration turns out to be correct.

[claude]

🟡 FlashInferEPBackend doesn't override _topk_idx_dtype() (default torch.int64 in ep_backend.py), so dispatch() re-casts p.topk_idx to int32 with .to() on every call instead of once at problem construction like DeepEP V2 and UCCL-EP do. Since dispatch() runs inside the timed window for benchmark_dispatch and benchmark_roundtrip, this adds a per-iteration device cast that no other GB-comparison backend pays, slightly inflating flashinfer-ep's measured latency. Fix: override _topk_idx_dtype to return torch.int32 and drop the .to() call in dispatch().

Extended reasoning...

ep_backend.py's make_problem() (line ~294) builds topk_idx=idx.to(self._topk_idx_dtype()) exactly once, untimed, at problem-construction time — this is called from Pass 1 of run_sweep, well outside any timed region. The default _topk_idx_dtype() (ep_backend.py:301-304) returns torch.int64. Every sibling adapter that needs a different dtype overrides this hook so the cast happens once: ep_deepep_v2.py overrides it to deep_ep.topk_idx_t, and ep_uccl.py overrides it as well. ep_mori.py similarly casts to int32 once inside its own make_problem-equivalent, untimed. FlashInferEPBackend (this PR) does none of that, so problem.topk_idx stays int64, and dispatch() (ep_flashinfer.py:132) does p.topk_idx.to(torch.int32) — a genuine int64->int32 device cast that allocates a new tensor and launches a kernel — on every single call.

dispatch() runs inside the timed region for two of the three benchmark components. benchmark_dispatch times lambda: self.dispatch(p) directly (ep_backend.py ~378-391), and benchmark_roundtrip/run_roundtrip call self.dispatch(problem) inside the timed callable (ep_backend.py ~334-376). time_us records a CUDA start event, runs fn(), then records the end event, so the cast kernel launch lands squarely inside the measured window for both dispatch and roundtrip, on every iteration of every trial. (The combine component is unaffected: when combine_needs_redispatch=True, dispatch() runs inside pre=prep_combine, which is untimed for that component.)

No other GB-comparison backend pays this. This PR's stated purpose is a fair transport comparison on GB200/GB300 (deepep-v2, nccl-ep, and now flashinfer-ep), and the existing adapters demonstrate the codebase already has the intended seam for avoiding exactly this cost — _topk_idx_dtype() — which this new adapter simply forgot to use.

Step-by-step proof of the defect:

1. run_sweep Pass 1 calls backend.make_problem(T, idx, weights, x) once per shape, before any timing starts.
2. Inside make_problem, problem.topk_idx = idx.to(self._topk_idx_dtype()). Since FlashInferEPBackend does not override _topk_idx_dtype, this resolves to the base class's torch.int64, so problem.topk_idx is an int64 tensor on device.
3. Pass 2 begins timing. benchmark_dispatch (or run_roundtrip via benchmark_roundtrip) starts a CUDA event, then calls self.dispatch(problem) — for every one of the configured warmup and measured iterations.
4. Inside ep_flashinfer.py's dispatch(): p.topk_idx.to(torch.int32) executes. Since the source dtype (int64) differs from the destination (int32), this is not a no-op — PyTorch allocates a new int32 tensor and launches a device conversion kernel, then that new tensor is passed into self._a2a.dispatch(...).
5. The CUDA end event is recorded after dispatch() returns, so the cast kernel's launch (and its position in the stream) is captured inside the measured s->e interval.
6. Repeat for every iteration — unlike DeepEP V2 and UCCL-EP, which instead resolved problem.topk_idx to the right dtype back in step 2, so their dispatch() bodies pass p.topk_idx straight into the kernel with zero extra ops.

The fix mirrors the existing pattern exactly: add def _topk_idx_dtype(self): import torch; return torch.int32 to FlashInferEPBackend, and change dispatch() to pass p.topk_idx directly instead of p.topk_idx.to(torch.int32).

Severity: this is a real, reproducible measurement-fidelity bug, but the cast itself is a small kernel launch over a [T, topk] index tensor, negligible next to a multi-GPU MNNVL all-to-all collective and likely within run-to-run measurement noise. Nothing crashes, no incorrect result is produced, and the PR explicitly states this backend has not yet run on GB hardware, so no biased numbers have been published yet. It's a nit worth fixing before any published GB comparison, not a merge-blocker.

[Oseltamivir]

Correction: this HAS now run on GB hardware, and it fails

Superseding the "not yet run on GB hardware" note in the description. A full leg ran on gb200 (2 nodes x 4 GPU, EP8, normal, bf16, decode) through the real gb-nv launcher. The case fails at first kernel execution.

File "/ix/experimental/CollectiveX/bench/ep_harness.py", line 855, in run_sweep
    backend.warm(problem, CONDITIONING_ROUNDS_PER_SHAPE)
File "/ix/experimental/CollectiveX/bench/ep_backend.py", line 340, in warm
    torch.cuda.synchronize()
torch.AcceleratorError: CUDA error: unspecified launch failure

then Cuda failure 719 on rank 6 and a cascade of NCCL teardown failures across both nodes.

What this does and does not tell us. Two things did work: flashinfer_ep_prepare passed on the allocation (FlashInfer 0.6.8.post1 one-sided A2A available, rc=0), and create_buffer completed — so Mapping, MnnvlConfig and MoeAlltoAll(...) were all constructed and accepted. The fault is in executing the first dispatch/combine, and 719 is sticky and asynchronous, so the traceback points at the sync rather than the offending launch.

Prime suspect: the _TorchDistCommunicator shim. I wrote it to bridge MnnvlConfig to the harness's existing process group, guessing at the interface (Get_rank/Get_size/allgather/Split) from vLLM's CustomCommunicator. If its allgather semantics do not match what the MNNVL handle exchange expects, peers resolve to wrong pointers and the kernel writes out of bounds — which is exactly a 719. That shim is the least validated thing in this PR and the first thing to instrument.

Marked draft. Merging a backend that faults on the only hardware it targets would publish a leg that cannot green. #2449 is unaffected and still ready.

Next steps, in order: rerun with CUDA_LAUNCH_BLOCKING=1 to localize the launch; compare the shim against vLLM's CustomCommunicator contract; verify the workspace assert workspace.size(0) == moe_ep_size actually holds for a 2-node EP8 mapping.

---

中文：更正 —— 本 PR 现已在 GB 硬件上运行，并且失败

此说明取代描述中「尚未在 GB 硬件上运行」的表述。已在 gb200（2 节点 × 4 GPU、EP8、normal、bf16、decode）上通过真实的 gb-nv launcher 跑了一个完整 leg。该 case 在首次 kernel 执行时失败。

随后 rank 6 报 Cuda failure 719，并在两个节点上引发一连串 NCCL 拆除失败。

这说明了什么、又没说明什么。 有两点确实成功了：flashinfer_ep_prepare 在该分配上通过（FlashInfer 0.6.8.post1 one-sided A2A available，rc=0）；create_buffer 顺利完成 —— 因此 Mapping、MnnvlConfig 与 MoeAlltoAll(...) 均已成功构造并被接受。故障发生在首次 dispatch/combine 的执行阶段；719 是粘滞且异步的错误，因此 traceback 指向的是 sync 调用而非真正出错的 launch。

首要嫌疑：_TorchDistCommunicator 适配层。 它是我为把 MnnvlConfig 桥接到 harness 既有进程组而写的，接口（Get_rank/Get_size/allgather/Split）是参照 vLLM 的 CustomCommunicator 推测出来的。若其 allgather 语义与 MNNVL 句柄交换的预期不符，对端就会解析到错误指针、kernel 越界写入 —— 这正好对应 719。该适配层是本 PR 中验证最少的部分，应作为第一个插桩排查对象。

已转为 draft。 合入一个在其唯一目标硬件上就会崩溃的后端，等于发布一条永远无法变绿的 leg。#2449 不受影响，仍可合入。

后续步骤（按顺序）：用 CUDA_LAUNCH_BLOCKING=1 重跑以定位具体 launch；将该适配层与 vLLM CustomCommunicator 的契约逐项比对；确认 2 节点 EP8 映射下 workspace.size(0) == moe_ep_size 这一断言确实成立。

[Oseltamivir]

Bring-up progress: transport correct, blocked on combine accumulation precision

Five legs on gb200 (2 nodes x 4 GPU, EP8, normal, bf16). Each failure was a distinct real defect, all now fixed:

run	stopped at	defect	fix
1	first kernel, CUDA 719	shim missing barrier/bcast — wrote to unmapped peer memory	subclass CommBackend, full contract + EP barrier
2	correctness gate	inspect_dispatch returned a bare tensor; oracle fell back to its default report (max_rel=0.0, recv=0)	dispatch routing metadata as payloads 2/3, build the real view
3	oracle setup	view missing local_expert_counts	read every view.* access; also found expert_ids needed non-local masking
4	combine_transformed	scatter buffer sized from ladder max, mask from actual receive	zeros_like on the real receive
5	correctness gate	— see below	—

Run 5: the transport is correct

T=1      passed=True   max_rel=0.024414  recv=45      rt_p50=222.8us
T=16     passed=True   max_rel=0.027466  recv=681     rt_p50=225.1us
T=32     passed=False  max_rel=0.036621  recv=1338    rt_p50=224.4us
T=512    passed=False  max_rel=0.048828  recv=21640   rt_p50=259.7us
T=8192   passed=False  max_rel=0.062561  recv=346930  rt_p50=1881.0us

Receive counts are exactly right (~5.3 ranks per token for topk=8 over 8 ranks), routing identity holds, and small rungs pass the tight gate. Data is moving correctly.

`combine_weight_semantics = "unweighted-rank-sum" is confirmed by measurement. This was the open question no reading of upstream source could settle. If combine applied the top-k gates (~0.125 each), the error would be ~87%, not 2%. It does not. Closed.

The blocker

Errors are 6-16 ulps against the 8 * 2^-8 = 0.03125 gate, present at every rung — T=1 is already 6 ulps. They only cross the gate as the max is taken over more elements, which is why small T passes and large T does not. That is a systematic precision difference, not noise.

The oracle models single-domain combine as FP32_sum(BF16(per-rank FP32 aggregate)) — ranks inside one scale-up domain reduce in FP32. FlashInfer's one-sided combine appears coarser. The shipped 0.6.8.post1 combine(payload, runtime_max_tokens_per_rank, payload_in_workspace) exposes no accumulation-precision knob (upstream main has output_dtype/use_low_precision, this wheel does not).

Not loosening the gate. It is the thing that catches real corruption, and widening it to fit one backend would silently widen it for every backend. The next step is to determine the kernel's actual reduction empirically — an on-device FP32 reference against the kernel output for one small case — and then either declare a third reduction model in the oracle if FlashInfer's is legitimately different, or take it upstream if it is a kernel precision bug.

Staying draft.

---

中文：bring-up 进展 —— 传输已正确，卡在 combine 累加精度

在 gb200（2 节点 × 4 GPU、EP8、normal、bf16）上跑了五个 leg。每次失败都是一个独立的真实缺陷，均已修复（见上表）。

第 5 次运行：传输本身是正确的

接收计数完全正确（topk=8、8 个 rank 下每 token 约 5.3 个 rank），路由一致性成立，小档位通过了严格门禁。数据搬运正确。

`combine_weight_semantics = "unweighted-rank-sum" 已由实测确认。 这是任何源码阅读都无法确定的未决问题。若 combine 施加了 top-k 权重（每个约 0.125），误差会是约 87% 而非 2%。它没有。该问题就此关闭。

阻塞点

误差为 6–16 ulp，门禁为 8 * 2^-8 = 0.03125，且每个档位都存在 —— T=1 时已达 6 ulp。只是随着取 max 的元素数量增加才越过门禁，因此小 T 通过、大 T 不通过。这是系统性精度差异，而非噪声。

oracle 对单域 combine 的建模是 FP32_sum(BF16(每个 rank 的 FP32 聚合)) —— 同一 scale-up 域内的 rank 以 FP32 归约。FlashInfer 的单边 combine 似乎更粗。镜像中 0.6.8.post1 的 combine 没有暴露累加精度选项（上游 main 有 output_dtype/use_low_precision，该 wheel 没有）。

不会放宽门禁。 它正是用来捕捉真实数据损坏的机制，为迁就某一个后端而放宽，等于对所有后端同时放宽。下一步是用实验确定该 kernel 的真实归约方式 —— 对一个小 case 做设备端 FP32 参考值与 kernel 输出的比对 —— 然后要么在 oracle 中声明第三种归约模型（若 FlashInfer 的确实不同），要么作为 kernel 精度问题上报上游。

保持 draft 状态。

[Oseltamivir]

Instrumented diagnostic: the residual is a 1-ulp contribution mismatch, not rounding mode

Ran the oracle comparison with instrumentation (scratch branch, never merged). Max-error element per rung, decode EP8:

actual=0.032226562   expected=0.031738281   absdiff=2^-11    ulps=3.94
actual=0.037597656   expected=0.037109375   absdiff=2^-11    ulps=3.37
actual=-0.06640625   expected=-0.065429688  absdiff=2^-10    ulps=3.82
actual=-0.079101562  expected=-0.078125     absdiff=2^-10    ulps=3.20
actual=-0.016601562  expected=-0.016967773  absdiff=3*2^-13  ulps=5.53
expected->trunc == expected->nearest == expected     (already BF16-exact)

This corrects two earlier readings on this PR.

The rounding-mode work does not apply to these elements. expected is already BF16-representable, so truncate and nearest agree exactly. The combine_output_rounding model still describes the kernel correctly (the standalone probe stands: 129 ulps where nearest gives 130) and it did move T=1 below the gate, but it is not what fails the larger rungs.

The failures are not on zero-expected elements. I inferred that from the identical max_rel values across models; the instrumentation shows expected values of ~0.03 to ~0.08. That inference was wrong.

What the numbers say. Every absdiff is a clean power of two equal to one BF16 ulp at a per-rank contribution's own magnitude (a contribution near 0.125 has ulp 2^-11 = 0.00048828125, exactly the observed difference). The sign varies. So a single contribution lands one BF16 step away from the oracle's — not accumulation drift, not a systematic bias, not the output conversion.

Most likely cause: FP32 association order when one rank hosts several of a token's top-k experts. The oracle sums the top-k axis in problem order via (gate * scale).sum(dim=1); the adapter's staged aggregate is built from routing metadata read back off the wire. Both are defensible arithmetic, and they can land on adjacent BF16 values after the cast. Backends whose received rows carry a single expert each (the low-latency padded layouts) never hit this.

That is the next thing to test: make the adapter's staged per-rank aggregate reproduce the oracle's summation order exactly, and see whether the residual collapses. It is a concrete, bounded experiment — not a search.

Still draft. COMBINE_REL_TOL untouched throughout.

---

中文：插桩诊断结果 —— 残差是 1 ulp 的单项贡献偏差，而非舍入模式问题

带插桩运行 oracle 比对（scratch 分支，从未合入），各档位最大误差元素见上。

这更正了本 PR 此前的两处判断。

舍入模式的改动对这些元素并不适用。 expected 本身已可被 BF16 精确表示，因此截断与就近舍入结果完全相同。combine_output_rounding 模型对 kernel 的描述依然正确（独立探测结论成立：129 ulp，而就近舍入为 130），也确实把 T=1 压到了门禁之下，但它不是大档位失败的原因。

失败元素并非期望值为零的元素。 我此前根据不同模型下 max_rel 数值完全相同推断为零，插桩显示期望值在 ~0.03 到 ~0.08 之间。该推断是错的。

这些数字说明了什么。 每个 absdiff 都是干净的 2 的幂，恰好等于某个 per-rank 贡献在其自身量级下的一个 BF16 ulp（量级约 0.125 的贡献其 ulp 为 2^-11 = 0.00048828125，与实测完全一致），且符号有正有负。因此是单个贡献落在了与 oracle 相差一个 BF16 步长的位置 —— 不是累加漂移，不是系统性偏置，也不是输出转换。

最可能的原因：当同一个 rank 承载某 token 的多个 top-k 专家时，FP32 结合顺序不同。oracle 通过 (gate * scale).sum(dim=1) 按 problem 顺序在 top-k 轴上求和；而 adapter 暂存的聚合值是用从线上读回的路由元数据构造的。两者算术都站得住脚，但转换为 BF16 后可能落在相邻取值上。那些接收行每行只含一个专家的后端（低延迟 padded 布局）不会遇到这个问题。

下一步要验证的正是这一点：让 adapter 暂存的 per-rank 聚合完全复现 oracle 的求和顺序，看残差是否消失。这是一个有界的具体实验，而不是漫无目的的搜索。

仍保持 draft。全程未改动 COMBINE_REL_TOL。

[Oseltamivir]

Final bring-up state: the residual is invariant to the oracle, so it is not an oracle problem

Three materially different expectation models were tested on gb200 hardware. The failing elements returned identical correctness to six decimal places every time:

oracle model	decode max_rel
run 5 — unweighted-rank-sum, nearest	0.024414 … 0.036621, 0.038147, 0.048828
run 8 — truncating output conversion	0.018311 … 0.036621, 0.038147, 0.048828
run 14 — canonical top-k summation order	0.018311 … 0.036621, 0.038147, 0.048828

Only T=1 moved (0.024414 -> 0.018311, below the gate) and only for the rounding-mode change. Everything from T=32 up is bit-for-bit unchanged across all three.

That invariance is the conclusion. An error that does not respond to the expectation is not produced by the expectation, and every lever available from outside the kernel acts on the expectation. Instrumentation localised it precisely — each failing element differs by exactly one BF16 ulp of a per-rank contribution (2^-11 near 0.125, 2^-10 near 0.25), sign varying, on expected values that are already BF16-exact — but nothing constructible in the oracle reaches those elements.

What is established and worth keeping:

• The transport is correct. Receive counts match routing exactly, routing identity holds, a standalone probe confirms every contribution is summed (2^rank across 8 ranks -> exactly 255.0), and the tight gate passes at T=1.
• combine_weight_semantics = "unweighted-rank-sum" — measured, not assumed. A gated combine would be ~87% off with topk=8; the error is 2%.
• The kernel truncates its FP32 -> BF16 output (129 ulps where nearest gives 130) — from a controlled probe, useful upstream regardless of what happens here.
• Six real defects found and fixed en route: missing barrier/bcast in the CommBackend shim (the CUDA 719), bare-tensor oracle view, missing local_expert_counts, unmasked non-local expert_ids, scatter buffer sized from ladder max, and the rounding mode.

Reverted along the way as measured dead ends: the bf16-rank-sum reduction model (built on a confounded probe, made agreement worse) and the canonical-ordering change (exact no-op). Both reversions are in the history with their evidence.

COMBINE_REL_TOL was never touched.

Recommendation unchanged: mark the GB rows an unsupported coverage row — the treatment b300 low-latency got in #2449 — which lands the adapter, the six fixes and all findings without publishing numbers that cannot reproduce. Closing it properly needs FlashInfer kernel internals; the probe results above are a ready-made upstream report.

---

中文：bring-up 最终状态 —— 残差对 oracle 不敏感，因此并非 oracle 的问题

在 gb200 硬件上测试了三种本质不同的期望值模型，失败元素每次都返回到小数点后六位完全相同的正确性结果（见上表）。只有 T=1 发生了变化（0.024414 -> 0.018311，降到门禁之下），且仅由舍入模式改动带来；T=32 及以上在三种模型下逐位不变。

这种不变性本身就是结论。 不随期望值变化的误差，不是由期望值产生的；而从 kernel 外部能用的每一个手段，作用的都是期望值。插桩已经精确定位：每个失败元素恰好相差某个 per-rank 贡献的一个 BF16 ulp（0.125 附近为 2^-11，0.25 附近为 2^-10），符号不定，且期望值本身已可被 BF16 精确表示 —— 但 oracle 侧能构造的任何模型都触及不到这些元素。

已确立并值得保留的结论、以及沿途修复的六个真实缺陷，见上文英文部分。作为实测死路而回退的两项改动（基于被混淆探测的 bf16-rank-sum 模型、以及完全空操作的规范化排序）都连同证据保留在提交历史中。

COMBINE_REL_TOL 全程未被改动。

建议维持不变：把 GB 行标记为 unsupported coverage row —— 即 #2449 中 b300 低延迟所采用的处理方式 —— 这样可以合入 adapter、六项修复与全部发现，而不发布无法复现的数据。要真正解决它需要 FlashInfer 的 kernel 内部信息；上述探测结果已经是一份可直接提交上游的报告。

[Oseltamivir]

Anatomy of the failing element: cancellation amplifying a one-ulp message

Instrumented the worst element directly (scratch branch, deleted):

rank=4 tok=350 col=200 rel=0.048828
  actual=0.0185546875   expected=0.017578125   absdiff=2^-10
  n_contributing_ranks=5   dests=[0,1,2,3,7]

Five per-rank contributions of magnitude ~0.25 summing to ~0.018 — heavy cancellation. The absdiff is exactly one BF16 ulp at 0.25, so a single contribution is one step off and the cancellation turns that into 5.5% of the result. Across 160 sampled elements the worst one always has 5-7 contributing ranks (66x six, 62x five, 32x seven), never one.

This corrects my previous comment, which guessed the worst elements were single-contributor. They are not.

It also explains the invariance that four expectation models could not budge. The defect is in one per-rank message, not in how the messages are combined — so reduction order, associativity and rounding mode are all irrelevant to it, exactly as observed.

Both sides compute that message from the same inputs through mathematically equivalent expressions: the adapter via _expert_transform(view.payload, ...), the oracle via semantic_x.float() * (gate*scale).sum(...) + .... Same coefficients, same fp32 arithmetic, same single narrowing. They agree bit-for-bit on the vast majority of elements and disagree by one step on a minority — the signature of two equivalent expressions evaluated through different tensor layouts or fusion choices, amplified here by cancellation and the 0.02 denominator clamp.

Where that leaves it: this is not a FlashInfer defect and not something a reduction model can fix. It is a last-bit disagreement between two harness-side expressions that only becomes visible when a backend's received rows carry multi-expert aggregates AND the combine cancels. Closing it means making the adapter's staged message and the oracle's expected message bit-identical by construction — computing both through the same code path — rather than relying on two independent formulations agreeing to the last bit.

That is a real, bounded change, but it touches shared oracle code and needs a full-matrix sweep to prove no other backend shifts. I am not landing it speculatively at the end of a long session.

---

中文：失败元素的构成 —— 相消放大了一个 ulp 的单条消息

直接对最差元素做了插桩（scratch 分支，已删除），数据见上。

五个量级约 0.25 的 per-rank 贡献相加得到约 0.018 —— 存在严重相消。absdiff 恰好是 0.25 处的一个 BF16 ulp，即单条贡献偏了一步，而相消把它放大为结果的 5.5%。在采样的 160 个元素中，最差元素的贡献 rank 数始终是 5 到 7 个（6 个出现 66 次，5 个 62 次，7 个 32 次），从来不是 1 个。

这更正了我上一条评论中「最差元素只有单个贡献者」的猜测 —— 事实并非如此。

这也解释了四种期望值模型都无法撼动的不变性：缺陷在于单条 per-rank 消息，而不在于这些消息如何被合并，因此求和顺序、结合方式与舍入模式对它统统无关，正如实测所示。

两侧都用相同输入、通过数学上等价的表达式计算这条消息，绝大多数元素上逐位一致，少数元素相差一步 —— 这正是两个等价表达式在不同张量布局或算子融合下求值的特征，并在此被相消与 0.02 分母下限放大。

结论：这不是 FlashInfer 的缺陷，也不是归约模型能解决的问题。要真正解决，需要让 adapter 暂存的消息与 oracle 期望的消息在构造上逐位相同（走同一条代码路径），而不是依赖两个独立表述在最后一位上碰巧一致。这是一个有界的真实改动，但它触及共享 oracle 代码，需要全矩阵扫描证明其他后端不受影响 —— 我不会在长时间会话的末尾贸然合入。

[Oseltamivir]

The oracle was never wrong — the kernel is not summing what the adapter stages

Summed every rank's staged contribution for one token (global token 0, col 40, T=1):

rank 0  0.0490722656      rank 4  0.0620117188
rank 1  0.0402832031      rank 5  0.0317382812
rank 2  0.0322265625      rank 6  0.0446777344
rank 3  (no row)          rank 7 -0.0044250488
                          --------------------
Σ staged        = 0.2555847168
oracle expected = 0.255859375    = BF16-nearest(Σ), exact
kernel actual   = 0.2578125      = Σ + ~2.3 ulps

The expectation is the correctly-rounded sum of the messages the adapter actually staged — bit-exact. The kernel's output is not that sum. So the correctness gap is not a modelling problem, and the five expectation models tried on this PR were all fixing something that was never broken. That also explains why every one of them was byte-identical at the failing rungs: they were varying a quantity that already agreed.

Where the bug is. combine_transformed scatters the oracle's rows back with
padded.view(-1, hidden)[self._valid_rows(h)] = transformed, relying on that mask enumerating slots in the same order inspect_dispatch used to build the view. If those two orders differ for any row — or if the kernel's combine reads the payload through topk_send_indices rather than the recv layout the mask describes — the kernel sums correct values from the wrong places, which is exactly a small, data-dependent discrepancy that grows with token count.

That is a concrete adapter bug with an address, not an open question about kernel arithmetic. It is the next thing to fix, and it invalidates the "needs FlashInfer internals" framing in my earlier comments: the internals are fine, the scatter is suspect.

---

中文：oracle 从来没有错 —— 是 kernel 没有对 adapter 暂存的值求和

对单个 token（全局 token 0、col 40、T=1）汇总各 rank 暂存的贡献，结果见上。

期望值正是 adapter 实际暂存消息之和的正确舍入结果 —— 逐位精确。而 kernel 的输出并不等于这个和。因此正确性差距并非建模问题；本 PR 上尝试过的五种期望值模型，修的都是本就没坏的东西。这也解释了为什么它们在失败档位上全部逐位相同：它们改动的那个量本来就已经一致。

缺陷位置。 combine_transformed 用 padded.view(-1, hidden)[self._valid_rows(h)] = transformed 把 oracle 的行散回去，依赖该掩码枚举槽位的顺序与 inspect_dispatch 构造视图时一致。若两者顺序在任何一行上不同 —— 或 kernel 的 combine 实际是通过 topk_send_indices 而非该掩码所描述的 recv 布局来读取 payload —— 那么 kernel 就会从错误的位置读取正确的数值，这恰好会表现为随 token 数增长的、依赖数据的小幅偏差。

这是一个有明确位置的 adapter 缺陷，而不是关于 kernel 算术的开放问题。它是下一步要修的地方，同时也推翻了我此前评论中「需要 FlashInfer 内部信息」的判断：内部实现没有问题，可疑的是 scatter。

[Oseltamivir]

Kernel arithmetic read from source: identical to the oracle's, so the inputs differ

cast_store on vec_t<float, N> -> bfloat16 resolves to vec_cast<nv_bfloat16, float>:

((nv_bfloat162*)dst)[i] = __float22bfloat162_rn(((float2*)src)[i]);
store: ((int4*)ptr)[i] = data[i];

Round-to-nearest-even, no scaling, no bias, no FMA anywhere in the store path. Combined with vectorized_combine_impl (fp32 accumulators per top-k slot, duplicates zeroed, pairwise tree, single narrowing), the kernel computes exactly RN(Σ fp32 of the messages it reads) — which is precisely what the oracle computes.

This closes the arithmetic question and invalidates my own earlier combine_output_rounding = "truncate" (already reverted). The probe that suggested truncation shared the top-k slicing confound described above.

So the diagnosis is now airtight rather than inferred: the arithmetic on both sides is the same, the oracle's expectation is bit-exactly BF16-nearest(Σ staged) (0.2555847168 -> 0.255859375, verified by summing every rank's staged contribution), and the kernel returns 0.2578125. The kernel is summing different values than the adapter stages.

Eliminated, each independently

candidate	how ruled out
reduction arithmetic	source: fp32 accumulate + RN, no extra terms — same as the oracle. Five models tested on hardware, all byte-identical
output rounding mode	source: __float22bfloat162_rn, not truncation
scatter index arithmetic	round-trips exactly (local test, no hardware)
buffer copy / layout	in-place staging into the receive tensor: identical results
per-rank slot mapping	probe: 2^rank across 8 ranks returns exactly 255.0
within-rank row mapping	at T=1 each plane holds one row — no ambiguity possible, yet the discrepancy persists

What is left

The values the kernel reads at the combine slots are not the values written there. moe_a2a_dispatch returns a combine_payload_offset and combine consumes it from internal state, so the combine payload occupies a separate workspace region from the dispatch receive — the copy into that region is the one step between "written" and "read" that remains unaudited. recv_counters and topk_send_indices are not exposed on the Python object, so that region's layout cannot be confirmed from the adapter.

That is the single question left, and it is bounded. Everything else on the list above is closed.

---

中文：从源码读到的 kernel 算术与 oracle 完全一致，因此差异出在输入

vec_t<float, N> -> bfloat16 的 cast_store 解析到 vec_cast<nv_bfloat16, float>（见上）：就近偶数舍入，存储路径中没有任何缩放、偏置或 FMA。结合 vectorized_combine_impl（每个 top-k slot 一个 fp32 累加器、duplicate 置零、成对树形归约、末尾单次 narrowing），kernel 计算的正是 RN(它所读取消息的 fp32 之和) —— 与 oracle 的计算完全相同。

这关闭了算术层面的疑问，也证明我此前的 combine_output_rounding = "truncate" 是错的（已回退）；提示截断的那次探测存在上文所述的 top-k 切片混淆。

因此结论现在是确凿的而非推断：两侧算术相同，oracle 的期望值逐位等于 BF16-nearest(Σ 暂存消息)（0.2555847168 -> 0.255859375，通过汇总各 rank 暂存贡献验证），而 kernel 返回 0.2578125。kernel 求和的并不是 adapter 暂存的那些值。

已逐项排除的候选见上表。

剩下的问题

kernel 在 combine 槽位读到的值，不是写入那里的值。moe_a2a_dispatch 会返回 combine_payload_offset，combine 从内部状态取用它，因此 combine 的 payload 位于与 dispatch 接收区不同的 workspace 区域 —— 从「写入」到「读取」之间，唯一尚未审计的就是向该区域的拷贝。而 recv_counters 与 topk_send_indices 并未暴露在 Python 对象上，因此无法从 adapter 侧确认该区域的布局。

这就是仅剩的一个问题，且边界清晰；上表中的其他项均已关闭。

[Oseltamivir]

Correction: an unexplained slot-order observation, not a clean retraction

I reported a layout probe as "self-consistent, no bug". That retraction was too confident and I want it on the record properly, because it points at the one thing still worth testing.

The probe wrote value = slot_index into every combine slot, with TOPK = world = 8, T = 2:

plane test: token0 = 8.000   expect 8.000    <- source-rank/plane mapping is correct
slot  test: token0 = 3.000   expect 0.000
            token1 = 5.000   expect 8.000

I explained the slot result away as send-index placement — dispatch putting a token at its send index rather than its token index, which combine then inverts. That explanation does not hold here. With TOPK = world, every token routes to every rank, so nothing is skipped and no compaction can occur; the send index should equal the token index on all eight ranks. Instead three ranks returned the two tokens' values swapped relative to the other five.

Totals are conserved (3 + 5 = 8, every contribution accounted for) and the plane mapping is exact, so this is not lost or duplicated data — it is an ordering question inside a plane.

I cannot currently tell whether that is real kernel behaviour or a third flaw in my probe, and after getting two probe expectations wrong today I am not going to assert either. But it is the only unexplained observation left, it is cheap to re-test, and it is adjacent to the residual: a per-rank slot-order difference is exactly the shape of defect that would make the kernel sum correct values from the wrong places while conserving totals — which is precisely what the arithmetic evidence says is happening (kernel 0.2578125 vs BF16-nearest of the staged sum, 0.255859375).

Next test, if someone picks this up: repeat with T = 4 and distinct per-slot values on a single source plane, and read back which slot each token receives. That distinguishes a genuine per-rank ordering difference from a probe artefact in one run.

---

中文：更正 —— 这是一个尚未解释的 slot 顺序现象，而非可以干脆撤回的结论

我此前把一次布局探测判定为「自洽、无缺陷」。这个撤回下得过于武断，有必要如实记录，因为它恰好指向唯一还值得测试的地方。

该探测在每个 combine 槽位写入 value = slot_index，参数为 TOPK = world = 8、T = 2，结果见上。

我当时用「send-index 放置」来解释 slot 结果 —— 即 dispatch 按 send index 而非 token index 放置，combine 再将其反转。但这个解释在此处不成立。 当 TOPK = world 时，每个 token 都会路由到每一个 rank，不存在跳过、也就不会发生压缩；send index 在全部八个 rank 上都应等于 token index。而实际结果是其中三个 rank 相对另外五个把两个 token 的值互换了。

总量是守恒的（3 + 5 = 8，所有贡献都在），plane 映射也精确无误，因此这不是数据丢失或重复，而是 plane 内部的顺序问题。

我目前无法判断这究竟是 kernel 的真实行为，还是我的探测第三次出错；在今天已经有两次探测预期出错之后，我不会对任何一种下断言。但这是仅剩的未解释现象，复测成本很低，而且它与残差高度相关：per-rank 的 slot 顺序差异，正好会导致 kernel 从错误位置读取正确数值、同时保持总量守恒 —— 而这恰恰就是算术证据所指向的情形（kernel 为 0.2578125，而暂存和的 BF16 就近舍入为 0.255859375）。

后续测试建议：改用 T = 4，在单个 source plane 上写入互不相同的 per-slot 值，回读每个 token 收到的是哪个槽位。一次运行即可区分「真实的 per-rank 顺序差异」与「探测本身的假象」。

[Oseltamivir]

Verified: dispatch slot index is a per-rank permutation of tokens, not token order

Ran the discriminating test (T=4, TOPK = world = 8, so every token routes to every rank and no compaction is possible; per-slot values 0..3, all exact in bf16):

got    : [17.0, 23.0, 5.0, 3.0]
expect : [ 0.0,  8.0, 16.0, 24.0]     (if slot order were token order on every rank)
sums   :  48   ==   48                 every contribution present, none lost or duplicated

Each rank contributes 0+1+2+3 = 6, times eight ranks = 48, and the total is conserved exactly. So the eight ranks each place the four tokens into slots in a different permutation. This is a real kernel property, reproducible, and it holds even when nothing is skipped.

This supersedes my earlier retraction on this thread, where I explained a similar T=2 result away as send-index placement and called it self-consistent. That was wrong: with TOPK = world there is nothing to compact, so send-index placement cannot produce a per-rank reordering. The observation was real; my dismissal was not.

It does not explain the correctness residual. The adapter is slot-agnostic end to end: inspect_dispatch reads payload, expert ids and weights from the same rows; the harness establishes token identity by decoding the payload rather than trusting slot position; and combine_transformed writes each transformed row back to the slot it was read from. A permutation that is consistent between dispatch and combine round-trips correctly, which is why the totals conserve here.

Recording it because it is a trap for anyone who touches this code: any change that assumes recv[r][t] is "rank r's token t" is wrong. That assumption is false on this kernel, and it fails silently — totals stay right while individual tokens get each other's values.

---

中文：已验证 —— dispatch 的 slot 索引是各 rank 各自的 token 排列，而非 token 顺序

运行了具有判别力的测试（T=4，TOPK = world = 8，因此每个 token 都会路由到每个 rank，不可能发生压缩；per-slot 取值 0..3，在 bf16 下均可精确表示），结果见上。

每个 rank 贡献 0+1+2+3 = 6，八个 rank 合计 48，总量精确守恒。也就是说，八个 rank 各自以不同的排列把四个 token 放入槽位。这是真实、可复现的 kernel 特性，且在没有任何跳过的情况下依然成立。

这一结论取代了我此前在本线程中的撤回 —— 当时我把类似的 T=2 结果解释为 send-index 放置并称其「自洽」。那是错的：当 TOPK = world 时没有任何内容需要压缩，send-index 放置不可能产生 per-rank 的重排。该现象是真实的，而我的否定不成立。

但它并不能解释正确性残差。 adapter 自始至终与 slot 无关：inspect_dispatch 从相同的行读取 payload、expert id 与权重；harness 通过解码 payload 而非依赖槽位来确定 token 身份；combine_transformed 把每个变换后的行写回它被读取的那个槽位。只要该排列在 dispatch 与 combine 之间保持一致，往返就是正确的 —— 这也正是此处总量守恒的原因。

之所以记录下来，是因为它对后续改动是一个陷阱：任何假设 recv[r][t] 就是「rank r 的第 t 个 token」的改动都是错的。 该假设在这个 kernel 上不成立，而且会静默失败 —— 总量看起来正确，个别 token 却拿到了彼此的值。

[Oseltamivir]

Best remaining lead: a rank that should contribute but has no row

Arithmetic from the cross-rank dump:

Σ staged (7 ranks with a row) = 0.2555847168
oracle expected               = 0.255859375   = BF16-nearest(Σ)
kernel actual                 = 0.2578125
actual - Σ staged             = 0.0022277832

Rank 3 had has_gtok0=False — no row for that token, so it staged nothing and its slot was written as zero. If rank 3 actually hosts one of that token's experts, the kernel reads its slot anyway, and a single missing contribution of 0.0022277832 makes the kernel's sum exactly correct while making the staged sum — and therefore the oracle's expectation, which was computed to match it — wrong.

That inverts the framing I have been using all evening. I assumed the kernel was returning a wrong value. This says the kernel may be right and the adapter may be failing to stage one contribution.

It fits every observation:

• Deterministic from the first invocation. An unwritten combine slot retains dispatch-era content, which is deterministic — unlike leaked state from a previous pass, which staleness testing already ruled out.
• Small and data-dependent, growing with token count as more slots are involved.
• Totals conserved in the layout probes, because those wrote every slot unconditionally and so never exercised this path.

The suspect is _valid_rows: it decides which slots are filled by testing the expert-id payload for the sentinel (.any(id != -1)). That is the one place this adapter infers state the kernel owns. If the kernel stamps the sentinel into a slot it nonetheless reads — the source describes it as marking "duplicate or dead-rank" per top-k slot, not strictly per unfilled slot — then a filled slot gets dropped from the view, never receives a staged message, and contributes whatever the workspace already held.

This is also why recv_counters not being exposed on the Python object matters: it is the authoritative per-source count that would replace the heuristic outright.

Next test, one run: print, for a single token, each rank's rank_in_dests (from the problem routing) alongside has_row (from _valid_rows). Any rank where those disagree is the bug. If they agree everywhere, this lead is dead too and the residual is genuinely kernel-side.

---

中文：最有希望的剩余线索 —— 某个本应贡献的 rank 却没有对应行

跨 rank 数据的算术关系见上。

rank 3 的 has_gtok0=False —— 该 token 在它上面没有对应行，因此它没有暂存任何消息，其槽位被写为零。如果 rank 3 实际上承载了该 token 的某个专家，kernel 仍会读取它的槽位；那么单独缺失的 0.0022277832 这一项，恰好会让 kernel 的求和完全正确，而让暂存和（以及据此构造的 oracle 期望值）成为错误的一方。

这颠覆了我整晚使用的框架：我一直假设 kernel 返回了错误的值；而这条线索说明 kernel 可能是对的，是 adapter 漏掉了一项贡献。

它与全部观测一致：从第一次调用起就是确定性的（未写入的 combine 槽位保留 dispatch 时期的内容，是确定性的，而非已被排除的跨 pass 泄漏）；数值小且依赖数据，并随 token 数增长；此前的布局探测总量守恒，因为它们无条件写入了每个槽位，从未触及这条路径。

嫌疑点是 _valid_rows：它通过检测 expert-id 负载中的哨兵（.any(id != -1)）来判断哪些槽位被填充 —— 这是本 adapter 唯一一处推断本应由 kernel 掌握的状态。若 kernel 会把哨兵写入某个它仍会读取的槽位（源码把它描述为按 top-k slot 标记「duplicate 或 dead-rank」，而非严格标记未填充槽位），那么一个已填充的槽位就会被排除在视图之外、永远收不到暂存消息，并把 workspace 中原有内容贡献出去。

这也正是 recv_counters 未暴露在 Python 对象上之所以重要的原因：它是能够直接取代该启发式判断的权威每源计数。

下一步测试，一次运行即可： 针对单个 token，打印每个 rank 的 rank_in_dests（来自 problem 路由）与 has_row（来自 _valid_rows）。两者不一致的 rank 就是缺陷所在；若处处一致，则这条线索同样告吹，残差确实在 kernel 侧。

[Oseltamivir]

Lead tested and dead: _valid_rows is correct, the residual is kernel-side

Ran the discriminating test from the previous comment — for one token, each rank's hosts_expert (from the problem routing) against has_row (from the adapter's _valid_rows sentinel):

rank=3 hosts_expert=False has_row=False  AGREE
rank=0 hosts_expert=True  has_row=True   AGREE
rank=1 hosts_expert=True  has_row=True   AGREE
rank=2 hosts_expert=True  has_row=True   AGREE
rank=4 hosts_expert=True  has_row=True   AGREE
rank=5 hosts_expert=True  has_row=True   AGREE
rank=6 hosts_expert=True  has_row=True   AGREE
rank=7 hosts_expert=True  has_row=True   AGREE

disagreements across all 160 sampled elements: 0

Rank 3 legitimately does not host an expert for that token, so staging nothing is correct. The sentinel heuristic agrees with the routing everywhere — it is not dropping filled slots, and recv_counters would tell it nothing new.

So the previous comment's inversion is withdrawn: the seven contributing ranks are the correct seven, each stages a correct message, Σ staged = 0.2555847168, and the oracle's 0.255859375 is the correctly-rounded sum of exactly the right contributions. The kernel's 0.2578125 remains unexplained.

Candidates eliminated, each by direct measurement or source

candidate	how
reduction arithmetic	source: fp32 accumulate, pairwise tree, no extra terms; five models tested, all byte-identical
output rounding	source: __float22bfloat162_rn, round-to-nearest, no truncation
scatter index arithmetic	round-trips exactly (local test)
buffer copy / zeros_like layout	in-place staging into the receive tensor: identical results
per-rank plane mapping	probe: 2^rank across 8 ranks returns exactly 255.0
within-rank slot mapping	per-rank permutation verified, and shown symbolically to cancel
combine-region layout	plane test exact; totals conserved
stale workspace state	discrepancy bit-identical on the FIRST invocation
dropped contribution (_valid_rows)	this test: 0 disagreements in 160 samples

Nine candidates closed. Everything reachable from the adapter is verified correct, and the kernel deterministically returns a value that is not the sum of the messages staged into its combine slots. Closing this needs someone who can instrument inside moe_a2a_combine itself.

---

中文：线索已验证并排除 —— _valid_rows 是正确的，残差在 kernel 侧

按上一条评论中给出的判别测试执行：针对单个 token，对比各 rank 的 hosts_expert（来自 problem 路由）与 has_row（来自 adapter 的 _valid_rows 哨兵判断），结果见上：160 个采样元素中零处不一致。

rank 3 确实不承载该 token 的任何专家，因此不暂存任何消息是正确的。哨兵启发式与路由处处一致 —— 它没有丢弃已填充的槽位，recv_counters 也不会带来新信息。

因此撤回上一条评论中的「反转」结论：贡献的七个 rank 就是正确的七个，各自暂存了正确的消息，Σ staged = 0.2555847168，oracle 的 0.255859375 正是这些正确贡献之和的正确舍入。kernel 的 0.2578125 仍无法解释。

已排除的九个候选及其依据见上表。凡是能从 adapter 侧触及的部分均已验证正确，而 kernel 仍确定性地返回一个并非其 combine 槽位中所暂存消息之和的值。要继续推进，需要能够在 moe_a2a_combine 内部插桩的人。

[Oseltamivir]

Synthesis: the residual is fp32 association order at BF16 rounding boundaries

Twelve structural candidates are now eliminated by direct measurement on gb200. The transport is correct on every axis I can construct:

tested	result
uniform payload	exact
per-rank distinct (2^rank)	exactly 255.0
per-slot distinct	per-rank permutation, shown symbolically to cancel
per-hidden-element distinct	exact, including across the 64-element vector boundary
full routing (TOPK = world)	exact
partial routing + compaction	exact (2.0 per token, two destinations)
duplicate experts on one rank	exact (1.0 — read once, dedup works)
plane / source-rank mapping	exact
scatter round-trip	exact (local test)
buffer copy vs in-place staging	identical
stale workspace state	discrepancy bit-identical on the FIRST invocation
_valid_rows vs problem routing	0 disagreements in 160 samples

What every passing probe has in common: small exact integers. The workload uses arbitrary BF16 values, and that is the whole difference.

I dismissed fp32 association order early on the grounds that tree-vs-sequential moves a sum by ~1e-9 while one BF16 step at these magnitudes is ~7e-5. That reasoning is only valid away from a rounding boundary. At a boundary, an arbitrarily small fp32 difference flips the narrowed result by exactly one ulp. Across 7168 elements per token a few always sit there, and where the combine has cancelled (five contributions near 0.25 summing to 0.018) one ulp is 3-6% relative — over the 8 * 2^-8 gate.

This is the only hypothesis consistent with all twelve eliminations and the residual:

• sparse — only elements near a boundary diverge, which is why most of the tensor agrees bit-for-bit;
• exactly 1-2 ulps, never more;
• nearly insensitive to model changes, because only the exact association matters. The tree model tried earlier reproduced the shape of the kernel's reduction but not its precise unrolled TOP_K == 8 pattern (acc[0]+=acc[1]; acc[2]+=acc[3]; acc[4]+=acc[5]; acc[6]+=acc[7]; acc[0]+=acc[2]; acc[4]+=acc[6]; acc[0]+=acc[4]), and anything short of bit-exact association is as wrong as sequential;
• invisible to every probe above, since exact small integers never land near a boundary.

What this means for the gate

Passing requires the oracle to reproduce the kernel's fp32 association bit-exactly — not approximately. That is achievable in principle (the pattern is in the source) but it makes the expectation a mirror of one kernel's unrolling, which is a real design question for this suite rather than a bug fix: every other backend passes because its association already matches the oracle's, not because the oracle is association-agnostic.

The alternative reading is that a tight relative gate with a 0.02 magnitude floor is measuring floating-point association on cancelling elements rather than transport correctness. That is a judgement call for whoever owns the correctness contract; I am not making it unilaterally, and I have not touched COMBINE_REL_TOL.

---

中文：综合结论 —— 残差源于 BF16 舍入边界处的 fp32 结合顺序

已在 gb200 上通过直接测量排除十二项结构性候选（见上表）。传输在我能构造的每一个维度上都是正确的。

所有通过的探测有一个共同点：使用的都是小整数且可精确表示。 而实际负载使用的是任意 BF16 数值 —— 这正是全部差别所在。

我早先以「树形与顺序累加的 fp32 差异约为 1e-9，而该量级下一个 BF16 步长约为 7e-5」为由排除了结合顺序。该论证仅在远离舍入边界时成立。在边界处，任意微小的 fp32 差异都会使 narrowing 结果整整偏移一个 ulp。 每个 token 有 7168 个元素，总有若干落在边界上；而在 combine 发生相消处（五个约 0.25 的贡献相加得到 0.018），一个 ulp 就是 3-6% 的相对误差，超过 8 * 2^-8 门禁。

这是唯一与十二项排除结论以及残差都相容的假设，理由见上（稀疏、恒为 1-2 ulp、对模型改动几乎不敏感、且对上述所有探测不可见）。

这对门禁意味着什么

要通过，oracle 必须逐位复现 kernel 的 fp32 结合顺序 —— 近似不行。原则上可行（该展开模式就在源码里），但这会让期望值变成某一个 kernel 展开方式的镜像，这对本套件而言是一个真实的设计问题，而非缺陷修复：其他后端之所以通过，是因为它们的结合顺序本就与 oracle 一致，而不是因为 oracle 与结合顺序无关。

另一种解读是：带 0.02 量级下限的严格相对门禁，在相消元素上度量的其实是浮点结合顺序，而非传输正确性。这属于正确性契约所有者的判断范畴；我不会单方面决定，也始终未改动 COMBINE_REL_TOL。

[Oseltamivir]

Correction to the synthesis above

I wrote that the tree model tested earlier "reproduced the shape of the kernel's reduction but not its precise unrolled TOP_K == 8 pattern". That is wrong, and it was the load-bearing claim.

The kernel unrolls:

acc[0]+=acc[1]; acc[2]+=acc[3]; acc[4]+=acc[5]; acc[6]+=acc[7];
acc[0]+=acc[2]; acc[4]+=acc[6];
acc[0]+=acc[4];

which is ((0+1)+(2+3)) + ((4+5)+(6+7)). The generic pairwise halving in that experiment produces exactly the same association for eight elements. So run 15 did match the kernel's association bit-for-bit, and the result was still byte-identical to the sequential baseline.

So the association-at-rounding-boundaries theory is not confirmed — it was tested and did not move the numbers. It remains the only mechanism I know of that is consistent with all twelve structural eliminations, but consistency is not evidence, and my own experiment argues against it. Either the per-slot values fed into that tree differed from what the kernel accumulates in some way I have not identified, or the mechanism is something else entirely.

Correcting it because a tidy theory left standing at the end of a long investigation is worse than an open question: it would send the next person to verify association order, which is the one thing already ruled out by experiment.

Accurate final state of the diagnosis: twelve structural candidates eliminated by measurement; kernel and oracle arithmetic proven identical from source; the oracle's expectation proven bit-exactly the correctly-rounded sum of the staged messages; the kernel deterministically returns a different value; no confirmed mechanism.

---

中文：对上一条综合结论的更正

我此前写道，先前测试的树形模型「复现了 kernel 归约的形态，但未复现其精确的 TOP_K == 8 展开模式」。这是错的，而且它正是整个论证的支点。

kernel 的展开（见上）等价于 ((0+1)+(2+3)) + ((4+5)+(6+7))。那次实验中使用的通用成对折半，对八个元素产生的正是完全相同的结合方式。因此第 15 次运行确实逐位匹配了 kernel 的结合顺序，而结果仍与顺序累加基线逐位相同。

所以「舍入边界处的结合顺序」这一理论并未得到证实 —— 它已被测试且没有改变任何数值。它仍是我所知唯一与十二项结构性排除结论相容的机制，但相容并不等于证据，而我自己的实验恰恰与之相悖。要么是喂入该树形归约的 per-slot 数值在某种我尚未识别的方式上与 kernel 实际累加的不同，要么机制完全另有其因。

之所以更正，是因为在一场漫长调查的末尾留下一个看似圆满的理论，比留下一个公开问题更糟：它会把下一个人引去验证结合顺序，而那恰恰是已被实验排除的那一项。

诊断的准确最终状态： 十二项结构性候选已通过测量排除；kernel 与 oracle 的算术已从源码证明一致；oracle 的期望值已证明逐位等于暂存消息之和的正确舍入；kernel 确定性地返回不同的值；尚无已确认的机制。

[Oseltamivir]

Update — the correctness failures are resolved

flashinfer-ep is now green on gb200 EP8: decode status=success (10/10 rungs), prefill
status=success (4/4), zero correct=False on any rank.

What was wrong: the oracle, not the adapter. FlashInfer's one-sided combine does not
accumulate in FP32 the way every other backend here does. Through wheel 0.6.15,
moeAlltoAllKernels.cu:556 is

template <typename T, int ELEMS_PER_VEC>
__device__ void accumulate_vec(T* dst, T const* src) { dst[j] += src[j]; }   // T = payload dtype

over acc[TOP_K] held in the payload dtype and reduced by a hand-unrolled pairwise tree —
(a0+a1)(a2+a3)(a4+a5)(a6+a7); (a0+a2)(a4+a6); (a0+a4) — then stored raw with no final cast.
That is three BF16 roundings where the oracle modelled one, which is exactly the few-ulp gap
that was failing the 0.03125 gate from T=32 up.

The operands sit at their original top-k slot: the dispatch kernel blanks duplicate-rank
slots in place rather than compacting, so the tree's shape depends on the routing, not just on
how many ranks contributed.

Fix: backends declare combine_reduction; "domain-fp32" (default, unchanged for every
existing backend) or "topk-slot-tree". The gate is not widened.

Version-gated, because upstream fixed this: 0.6.16 ("Port the TensorRT-LLM one-sided A2A
optimizations") rewrote the accumulator to FP32 with a single narrowing store. The adapter
reads flashinfer.__version__ and switches models at that boundary, so a wheel bump does not
silently re-break the oracle.

Discriminating evidence. Contributions {1.0, 7 × 2⁻⁹} give three different answers:
bf16-tree 1.0078125, fp32-then-narrow 1.015625, sequential bf16 1.0. gb200 returns
1.0078125. That value was measured early in this branch and misread as evidence of a
truncating store — it is not; it is the tree. Pinned in TopkSlotTreeReductionTests, which
asserts the match and the two non-matches.

Also in this update: combine now submits its payload through
get_combine_payload_tensor_in_workspace(...) + payload_in_workspace=True, the sanctioned
zero-copy form vLLM uses. This did not affect correctness (measured — still red at T=32
before the oracle fix); it is here for measurement fidelity, because it takes the kernel's
staging copy out of the number we publish. Combine p50 at T=1 fell 131.6 → 82.3 µs. The copy
is still executed and still reported, as stage.

---

中文：正确性失败已解决

flashinfer-ep 在 gb200 EP8 上已全绿：decode status=success（10/10），prefill
status=success（4/4），各 rank 均无 correct=False。

问题出在 oracle，而非 adapter。 FlashInfer 单边 combine 并不像本仓库其他后端那样以 FP32
累加。在 0.6.15 及更早的 wheel 中，accumulate_vec 的 T 就是 payload 精度，acc[TOP_K]
亦以该精度保存，并由手工展开的成对树归约后原样存储、无最终转换 —— 即三次 BF16 舍入，而 oracle
只建模了一次。这正是 T=32 起突破 0.03125 阈值的那几个 ulp。

操作数停留在其原始 top-k 槽位：dispatch kernel 就地将重复 rank 的槽位置空而不压缩，因此树的
形状取决于路由本身，而非贡献 rank 的数量。

修复： 后端声明 combine_reduction，取 "domain-fp32"（默认，现有后端行为不变）或
"topk-slot-tree"。阈值不放宽。

按版本切换，因为上游已修复：0.6.16 将累加器改为 FP32 加单次收窄存储。adapter 读取
flashinfer.__version__ 并在该边界切换模型，避免 wheel 升级后 oracle 再次静默失配。

判别性证据： 贡献 {1.0, 7 × 2⁻⁹} 在三种模型下给出三个不同结果：bf16 树 1.0078125、
fp32 后收窄 1.015625、bf16 顺序累加 1.0。gb200 返回 1.0078125。该数值本分支早期即已测得，却被
误读为截断存储的证据 —— 它其实是树。TopkSlotTreeReductionTests 同时锁定这一匹配与两项不匹配。

本次一并包含： combine 改用 get_combine_payload_tensor_in_workspace(...) +
payload_in_workspace=True 提交载荷，即 vLLM 所用的官方零拷贝形式。此项未改变正确性
（已实测，oracle 修复前 T=32 仍为红），保留是为了测量保真：它把 kernel 的暂存拷贝移出我们发布的
combine 数值。T=1 的 combine p50 由 131.6 降至 82.3 µs。该拷贝仍会执行并作为 stage 上报。