6 Commits

Author SHA1 Message Date
whn09
1dce2c4995 Neuron bench: vllm-neuron#71's sonnet prompts, prompt/output dumps; DSpark draft ground truth
- bench_neuron.py --dataset vllm-sonnet: vLLM's SonnetDataset as #71's bench
  builds it for the completions backend ("Pick as many lines..." + random poem
  lines, chat-formatted; --chat_tokenizer for a template the checkpoint lacks).
  --dataset sonnet: a contiguous text window. --dump: prompt/output token ids.
- reference_dspark.py: the official inference/model.py forward_spec on the host
  CPU (tiny checkpoint, teacher-forced on reference_golden.py tokens): each
  decode step's 5 drafts and their logits.
- run_deepinfer_tiny.py RUN_TINY_DRAFTS=1 prints every draft block.

DeepInfer's drafts match the reference on host eager, on device and with
attention TP: every divergence is a near tie (top-1/top-2 gap 0.001-0.055).
With #71's prompt format the acceptance is #71's: 2.8 tokens a step at 512
(first position 65-79%); a raw text window gave 1.4 (continuations not in the
context), random token ids 3.85 (repetitive outputs).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-26 03:16:54 +08:00
whn09
b2d94e16b2 Neuron DSpark (#15): greedy verification and MTP ops, CP followers run the drafts
- fused_sample reject mode for greedy rows (accept iff the draft is the
  argmax), get_accept_result / finalize_verification (static final rows),
  store_mtp_info (static pool writes, slot 0 kept), confidence_head_fused,
  dflash_input_store, get_fused_dflash_get_input: the frontend refs with the
  kernels' argument lists.
- CP followers run the draft forward (its MoE has collectives); the driver's
  accepted offsets reach them through one all-reduce on Neuron (no eager
  broadcast), the request first rows through the graph metadata.
- run_deepinfer_tiny.py --nextn: free-running greedy with drafts. Host tiny,
  nextn=5, eager: short 9/9 golden tokens; long matches up to its known near
  tie at step 5.
- bench17: the two-level MoE all-reduce does not move TPOT (433 ms).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-24 11:31:59 +08:00
whn09
f9cbf6f2e5 Neuron: CP group as a replicated-decode group; MoE all-reduce mode (#28)
- 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>
2026-10-23 22:46:04 +08:00
whn09
bf447622c2 Neuron: EP > 1 with real neuron collectives (issue #6)
- 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>
2026-10-21 11:46:59 +08:00
whn09
380e160b4b Neuron flash_mla: round attention P to bf16 like the CUDA kernels
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>
2026-10-20 20:48:36 +08:00
whn09
8a5209ee65 Neuron basic mode — part 3: end-to-end check vs the official reference
- 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>
2026-10-20 20:20:02 +08:00