- Followers claim core cp_rank (DEEPINFER_NEURON_LOCAL_RANK at spawn) and get
the graph's host metadata through shm (no in-graph broadcast on Neuron),
including the prompt-as-decode bucket and the context limit; they skip the
sampler. max_num_prefilling_tokens is not divided by cp on Neuron (prompt
chunks run replicated).
- DEEPINFER_NEURON_NUM_KV_BLOCKS: the fixed KV pool as an env var, so spawned
CP followers take it too.
- static MoE: when the EP group is the CP group the tokens are already
replicated: no all-gather, all-reduce instead of reduce-scatter (64x fewer
expert rows a rank at cp=ep=64).
- run_deepinfer_tiny.py --cp. Host tiny: cp=2 and cp=4 match the cp=1 tokens
(short 8/9 at a near tie, long 9/9).
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>
The pure-torch sparse attention followed FlashMLA's tests/ref.py (fp32 P).
The CUDA kernels (sm100 prefill phase1.cuh:238-253) and the official
reference round P = exp(s - running_max) to bf16 before the P@V GEMM in an
online softmax; do the same with 64-key blocks.
Localized with examples/neuron/probe_layer0.py (per-op reference vs DeepInfer
capture): layer 0 is now bit-exact up to two rounding points that are
DeepInfer's own GPU/Ascend behaviour (SwiGLU -> FP8 from fp32, bf16 routed
partial in the MoE combine). CPU end-to-end on the tiny checkpoint: 26/27
greedy tokens match the reference, the miss being a 0.0096-margin step the
reference flips against itself.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
- examples/neuron: tiny random V4.1 checkpoint in the real HF layout,
pure-torch replacements for the official reference kernels, golden
generation, a teacher-forced DeepInfer runner (LLM offline API) and a
comparison tool. CPU mode: 25/27 greedy tokens match, worst logits
cos 0.9985 / rel 0.056 against a reference self-noise floor of 0.032-0.060.
- model_loader: check parameters against the platform device, not is_cuda.
- deep_gemm paged MQA logits: do not dereference block-table slots past the
context.
- bind_cpu: skip in Neuron CPU mode (1 core per NeuronCore share would
serialize host compute).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
