7 Commits

Author SHA1 Message Date
whn09
1ecd3d9859 neuron: norms feeding only NKI kernels skip their dead fp8 copy (TTFT 1387 -> 1229 ms, step 457 -> 450 ms, tokens identical); capacity MoE wired into the layer region graph with its overflow flag riding the MoE all-reduce (still off: per-layer host sync makes prefill slower, 1823 ms)
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-28 05:29:40 +08:00
whn09
cedd237772 neuron: fp8-activation linears take the NKI kernel (Engram projection); drop unused fp8 hidden copies before region splits; log weights without NKI scales
An fp8 tuple input fell back to the deep_gemm shim, which dequantized the whole fp8 weight in the graph on
every call (Engram fused_kv_projection: 9 ms of an Engram layer's prefill graph, GpSimd POW for the scales).
Now decoded to bf16 and run through the NKI kernel. TTFT 1.71 -> 1.66 s (bucket 512), tokens unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-27 11:50:34 +08:00
whn09
c4c50f1ea5 neuron: shared expert sliced over ranks, summed in the MoE all-reduce (DEEPINFER_NEURON_SHARED_SLICE=1)
18 of 64 ranks compute a 128-column slice of the 2304-wide shared expert (the rest add zeros, same graph);
the partial joins the routed experts' sum before the existing all-reduce. Decode step 591 -> 581 ms.
Summation order changes, so greedy outputs differ from before (coherent; tokens/step 2.70/3.10/3.34,
TPOT p50 184 ms). Tiny golden 9/9. Also DEEPINFER_NEURON_HOST_TIMING (host time per region call).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-27 01:31:41 +08:00
whn09
3c38fd5829 Neuron MoE: skip unrouted experts inside the fp4 kernels at decode sizes (DEEPINFER_NEURON_GATED_EXPERTS, default on)
The fp4 expert GEMMs dequantize the whole weight to bf16 in HBM on every call
(device profile: the MoE graph reads 756 MB a layer), and a rank computes all
6 of its experts while a 6-row step routes to about one. fp4_linear_2core_gated
takes a per-expert flag (any row routed to it) and runs the dequant and GEMM in
an nl.fori_loop bounded by it, so an unrouted expert costs a zero fill. Decode
sizes only (<= 16 rows); prompt chunks route to every expert anyway.

bench70 (vllm-sonnet 512, TP=64, nextn=5): TPOT 258.9 -> 248.4 ms, step
738 -> 711 ms, tokens per step identical. Tiny cp=2 nextn=5 on device: 9/9
golden tokens, drafts unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-26 06:36:48 +08:00
whn09
e846322c0c Neuron: block-quant GEMMs on both physical cores (grid 2); sync experiment knob (#27)
At LNC=2 a logical NeuronCore is two physical cores and a kernel launched at
grid 1 only has code for one of them. The vllm-neuron#71 _2core kernels split
the output columns by program_id; launched at grid 2 they are 1.7-1.9x faster
on every decode GEMM shape and bitwise equal on one core (fp8 1x1280x36864
2.22 -> 1.17 ms, fp4 64x5120x4608 0.68 -> 0.40 ms). Used when n % 256 == 0;
DEEPINFER_NEURON_BQ_2CORE=0 turns it off.

Device profile of a decode attention region (neuron-explorer): 12.9 ms, of
which the q_b_proj fp8 GEMM alone is ~1.8 ms of vector-engine dequant.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-23 11:08:05 +08:00
whn09
c24574bfd6 Neuron: full DeepSeek-V4.1-Flash load + decode region graphs at EP=64 (issues #21, #25)
- Layer-major KV / state layout so per-layer views are contiguous graph inputs
  (the executor copied non-contiguous inputs every call).
- One compiled entry per (class, method), rank as a registered buffer, no
  layer_idx guard: every layer reuses the same graphs and loaded NEFFs.
- fp8 weight rescale on the host (numpy LUT, ~2 min instead of 35+ min of
  eager per-shape compiles); engram sf packing on the host.
- int16 ue8m0 packing in fp32 (int32 interleaved slices narrowed to int16
  crash neuronx-cc, NCC_IBCG902).
- Token embedding and the text model's vision tower stay in host memory on
  Neuron (2.5 GiB of HBM a rank).
- offline example: --max_num_seqs.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-23 01:39:22 +08:00
whn09
9c8c85d83d Neuron: per-layer region graphs + vllm-neuron#71 NKI kernels (issues #10, #11, #13, #25)
- Decode runs eagerly and compiles per-layer regions (attention half, MoE) with stage boundaries
  (neuron_ops/regions.py, DEEPINFER_NEURON_GRAPH_MODE=regions); layer KV views are sliced from the
  roots inside the graph (_LazyKVCache).
- vllm-neuron#71 NKI kernels (Apache-2.0, vendored under neuron_ops/nki/ with attribution), called via
  torch_neuronx.nki_hop.wrap_nki: blockquant (fp8/fp4 GEMMs incl. grouped einsum, MoE and shared
  expert on bf16 activations), sparse_attn / index_score / compress_pool (nki_attention.py),
  hc_mixes / hc_expand / hc_collapse / hc_attn_out / rmsnorm / swiglu (nki_hc.py, nki_norm.py).
- Compile-safe formulations: row-major scales, no dtype bitcasts of graph inputs (uint8 weight
  tensors), whole-row KV stores, byte assembly in int32 narrowed once at the store, arithmetic int16
  scale packing, plain fma/rsqrt and bf16 F.linear under compile; default backend deepinfer_neuron
  (torch fp8 carried as uint8, avoiding NCC_EOCP001 with the kernels' OCP fp8).
- Host: tiny model 9/9 vs golden on all prompts; device: kernel tests pass (13 attention, 88 hc/norm);
  region decode graph compiles except the attention projection stage (being reduced).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
2026-10-22 17:36:54 +08:00