Uneven all_to_all on the neuron process group is unreliable at scale (EP=64:
"Invalid NEFF" in prefill; test_ep at 64 ranks hung after 9/18 modes).
EPBuffer now moves tokens without all_to_all:
- dispatch: one all_gather of x / sf / weights / topk_idx packed as uint8
rows (row count = global max via a 0-dim MAX all_reduce, or the static
num_max_tokens_per_rank with DEEPINFER_NEURON_EP_AG_ROWS=cap); each rank
selects the rows routed to its local experts in the same (source rank,
source token) order as before, so layout / metadata / cached dispatch are
unchanged.
- combine: per-token fp32 partials over local experts (bf16-rounded as
before) + weights, one fp32 SUM reduce_scatter (all_reduce fallback with
DEEPINFER_NEURON_EP_REDUCE=ar), chunked under
DEEPINFER_NEURON_EP_CHUNK_BYTES (64 MiB).
- The a2a path stays behind DEEPINFER_NEURON_EP_ALGO=a2a.
test_ep passes on the neuron group at 2/4/8/16/32/64 ranks (combine
bit-identical to the reference in practice); new test_collectives.py found
that int32 all_gather with full-range values is wrong on a 16-rank neuron
group for >= 1 MiB outputs (uint8 / fp32 fine) -- payloads move as bytes.
Tiny model EP=8 on NeuronCores: 25/27 argmax vs the reference, same as EP=1.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- deep_ep comm: torch-neuronx picks the all_to_all lowering per rank from
that rank's own splits (even / NKI a2a-v / 1-D-only padding fallback), so
ranks could diverge. All-gather the split matrix first and choose one path
for all ranks: even a2a, NKI a2a-v for groups of 4/8, otherwise a padded
even a2a (2 ranks, 16 ranks where a2a-v fails on runtime 2.35).
- PlatformNeuron: LOCAL_RANK=0 / LOCAL_WORLD_SIZE=1 so every one-core
process is neuron:0 of a 1-device local world (otherwise rank > 0 died in
_lazy_init at EP/DP > 1).
- tests/neuron/deep_ep: references on host copies with the gloo group,
buffers on the device; ep/engram/bucket pass at world size 2/4/8/16 on the
neuron process group, and 1/2/4 on gloo.
- run_deepinfer_tiny.py --ep N: one process per rank / NeuronCore; tiny
model EP=2 runs end to end (CPU: within the noise floor).
- run_all.sh: per-core busy check, bucket + world size 16 device stages.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- tests/neuron/setup_trn2.sh: one-shot setup of a fresh Trn2 (venv layered
on the DLAMI torch-neuronx venv, C++ build, official reference + tokenizer,
tiny checkpoint, CPU golden).
- tests/neuron/run_all.sh cpu|device|all [stage...]: vendor kernel tests,
tilelang_ops interface tests (basic-mode stubs reported as skips via
nie_skip), deep_ep at world size 1/2/4 (gloo) and 2/4/8 on the neuron
process group, device primitive checks, end-to-end vs golden on CPU and
on a NeuronCore, and the device diff; PASS/FAIL summary.
- tests/neuron/device/test_device_primitives.py: torch-neuronx behaviours
the port depends on or works around (fp8 strided-view indexing, uint64,
1-rank all_to_all, inference-mode version counters, ...).
- examples/neuron/device_diff.py: re-runs every kernel call of a NeuronCore
run on host copies and ranks kernels by device/host difference.
- Unit tests written during bring-up moved from scratch dirs into
tests/neuron/{deep_gemm,flash_mla,deep_ep,tile_kernels,sampling}.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
