mirror of https://github.com/deepseek-ai/DeepInfer.git synced 2026-10-28 23:07:09 +08:00

whn09 d8f7a2e898 Neuron: PR #147 OpenAI server (+ /v1/completions, Neuron worker setup), offline bench, MoE fixes (#26 , #27 , #29 )

- serve/: the minimal OpenAI-compatible server of deepseek-ai/DeepInfer#147,
  with a Neuron worker setup (basic mode env, per-rank NeuronCore, fixed KV
  pool via --num-kv-blocks, no MTP lookahead) and a /v1/completions endpoint
  (prompt text or token ids, no chat template): what vllm bench serve
  --backend openai drives, as vllm-neuron#71's run_bench.py does.
- examples/neuron/bench_neuron.py: the same TTFT/TPOT matrix offline.
- static MoE: one packed fp32 all-gather instead of three (each 64-rank
  collective is ~2 ms of fixed latency a layer); a capacity-grouped expert
  path for prefill chunks (DEEPINFER_NEURON_MOE_CAPACITY, dense rerun on
  overflow; not yet validated on device).
- Sub-region timings at split() boundaries under DEEPINFER_NEURON_REGION_TIMING.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

2026-10-23 16:23:28 +08:00

4.4 KiB

Raw Blame History

Serving DeepInfer + host prefix cache (local additions, not part of upstream)

Pieces

serve/openai_server.py — OpenAI-compatible /v1/chat/completions (streaming, tools, usage) on top of the offline engine: one engine process per GPU (attention DP + MoE EP), prompts rendered with the checkpoint's own encoding.py. serve/server.sh start|stop runs it as a process group (stop reaps the engine helpers, which otherwise keep the GPUs).
deepinfer/hicache/ — host-memory prefix cache (SGLang HiCache style): pinned host pools organized in pages (one page = one 128-token KV block) plus SWA state slots, hash-chain index, LRU eviction, write-through at the end of every prefill, "kernel" transfers (a CUDA kernel copying pages through UVA pointers on the engine stream, page_copy.cu, JIT-built on first use). Needs the read-only Sequence.block_table binding added in csrc/engine/pybind_sequence.cpp (rebuild deepinfer._C with bash scripts/build.sh build).
tests/hicache/test_host_cache.py — cache logic against a fake engine (needs one GPU).
serve/check_prefix_cache.py — A/B on a live server: each turn sent with "prefix_cache": false and with the cache, greedy; compares outputs and TTFT.

Run

export DEEPINFER_ACCELERATOR=nvidia DEEPINFER_HF_MODEL_DIR=/path/to/DeepSeek-V4.1-Flash-1017
serve/server.sh start --dp 4 --port 8888                       # cache on: 32 GiB pages + 4096 states per rank
serve/server.sh start --dp 4 --port 8888 --hicache-size-gb 0   # cache off
python serve/check_prefix_cache.py --port 8888 --turns 3
curl -s localhost:8888/deepinfer/metrics | grep hicache         # per-rank hit/store counters
serve/server.sh stop

Knobs: --hicache-size-gb, --hicache-state-slots, --routing affinity|least-loaded (affinity keeps a conversation's turns on one rank, keyed by the first user message; required for hits since the cache is per rank). The cache requires decoder bounded replay (--yoco-keep-tokens > 0, the default 140).

EPLB expert rearrangement is off by default (--eplb-interval 0). Every rearrangement stalls all ranks for about 0.3 s, and with --eplb-after-prefill 1 it fires after nearly every step under multi-turn traffic; enable it with --eplb-interval 30 only for decode-heavy workloads whose expert load is skewed.

Prompt tokenization

Agentic clients resend the whole conversation every turn, and tokenizer.encode needs ~1.7 ms per 1K tokens on a GB300 Grace core (p50 278 ms per request on replayed AgentX prompts of p50 163K tokens), which dominates TTFT once the prefix cache hits. serve/prompt_tokenizer.py gives the same ids faster; --tokenizer-backend picks one:

fastokens (pip install fastokens; x86-64 and aarch64 wheels): Rust reimplementation of the HF pipeline, built from the loaded tokenizer. 3.5 ms p50 on the same prompts.
segment (no dependency): splits the prompt at added tokens, where HF tokenizers already encode the pieces independently, and caches each piece's ids (LRU, 512 MiB), so a turn only encodes what it adds. 18 ms p50. Enabled only when the tokenizer config makes the split exact (no single_word tokens, no prepend_scheme="first").
hf: tokenizer.encode.

auto (default) takes the first of these that matches tokenizer.encode on startup probes built from the tokenizer's added tokens. --tokenizer-verify-rate (default 0.005) re-encodes a sample of prompts with tokenizer.encode in the background; a mismatch switches to hf. The backend, encode latency and check counters are in /deepinfer/metrics (deepinfer_tokenizer:*). Tests: tests/serve/test_prompt_tokenizer.py.

How a hit works

A new prompt's page hashes are matched from the start; the match is capped so that at least keep_tokens tokens remain to be recomputed (the decoder SWA state is rebuilt from that suffix by bounded replay), then snapped down to the nearest 128-aligned position that has a stored SWA state. The sequence is allocated up front, pages and state are copied in on the engine stream, num_cached_tokens is set, and the sequence is enqueued; the scheduler prefills only the suffix. Generated tokens are not stored at decode end (they are stored when they come back as part of the next turn's prompt).

AgentX

The server speaks what InferenceX's AgentX replay (aiperf, --endpoint /v1/chat/completions --streaming --use-server-token-count) needs; prefill-any-interval matters there, since a single lane keeps most ranks idle.

4.4 KiB Raw Blame History