Views
No views yet
tpa_pre/ + tpa_dec/ (tp_0.pt + tp_1.pt = one per core), loaded via neuronx_distributed.trace.parallel_model_load.input_output_aliases) makes decode compute-bound; TP=2 then splits that compute across both cores. q/gate/up ColumnParallel, o/down RowParallel, k/v replicated (MQA).optb_server_tp.py (stdlib HTTP; parallel_model_load at startup, per-request KV reseed). Prebuilt image: docker pull xbill9/gemma4-optb:tp2-2048 (run --device /dev/neuron0 --ipc=host -p 8080:8080).temperature/top_k/top_p), SSE streaming (stream:true), /metrics (Prometheus), spot-drain + bounded queue (429) + per-request timeout + graceful SIGTERM — validated on Inferentia. Endpoints: /generate /v1/chat/completions /v1/completions /v1/models /health /metrics.parallel_model_load's lazy move_trace_to_device (fires on the first forward) replaces each states._parameters[...] with a new on-device tensor, so state params captured before that first forward become orphaned CPU refs — the per-request reseed then wrote to dead tensors and only the first request after startup was correct (later ones were contaminated by prior KV). Fixed in optb_server_tp.py by re-fetching the live state params each request (sp_now()) and warming up at startup so the device-move happens before serving. Verified: sequential unrelated prompts all correct.optb_server_tp_slim.py / image xbill9/gemma4-optb:tp2-slim. Same TP=2 neffs, but the host embeddings load slim (meta build + bf16 non-decoder safetensors + fp32 activation cast) so it fits a 16GB host (steady ~9GB, needs swap for the neff-load peak; grow root EBS to ~200GB for the image). Validated: correct output, ~61 tok/s on both cores (2.5x the 24 tok/s single-core :slim). Run --device /dev/neuron0 --ipc=host -p 8080:8080.kv_pre_2048.pt + kv_dec_2048.pt extend the context window to 2048 tokens (prompt bucket 512). Run with KV_MAX=2048 KV_BUCKET=512 KV_PRE_OUT=kv_pre_2048.pt KV_DEC_OUT=kv_dec_2048.pt.docker pull xbill9/gemma4-optb:2048-512 (now also :latest).google/gemma-4-E2B-it coherently and
fast on a single AWS Inferentia2 device (inf2.8xlarge), at ~44 tokens/sec.NxD / optimum-neuron stack
cannot currently do (there is no Gemma-4 in optimum-neuron, and NxD can't represent the
KV-sharing graph).| Base model | google/gemma-4-E2B-it (~5B params, 2B effective via MatFormer + Per-Layer Embeddings), Apache-2.0 |
| Hardware | AWS Inferentia2 — inf2.8xlarge, one logical NeuronCore (cores 0–1) |
| Precision | bf16 (fp32 neffs overflow the 16 GB core) |
| Throughput | ~44 tok/s (~23 ms/token), measured |
| Context | max_total_tokens=512, max_prompt_tokens=128 (baked into the neffs) |
| Cold start | ~100 s to load both neffs onto the core, then instant |
| Output parity | Greedy decode is token-for-token identical to the CPU reference |
torch_neuronx.trace()s the Hugging Face transformers (5.13) Gemma-4 text forward
pass directly, so KV-sharing traces as ordinary live graph dependencies — exactly like
it does on TPU/XLA.DynamicCache at trace time; a fixed static KV buffer at decode timekv_pre_512.pt) — consumes a padded prompt (≤ KV_BUCKET = 128 tokens),
returns the 15 non-shared layers' K/V (shared layers never write cache).kv_dec_512.pt) — single-token forward against a fixed KV_MAX = 512-length
KV buffer. Each step writes the new K/V via a one-hot masked scatter
(buf*(1-oh)+k*oh, pure arithmetic → trace-safe); KV tensors are graph inputs/outputs.KV_MAX = maximum total tokens (buffer length, baked into both neffs).
KV_BUCKET = maximum prompt tokens (baked into the prefill neff).| File | What it is |
|---|---|
kv_pre_512.pt | Prefill neff (TorchScript, bf16) |
kv_dec_512.pt | Decode neff (TorchScript, bf16) |
optb_kv.py | Builds/compiles both neffs (cpu = reference check, trace = compile) |
optb_server.py | Stdlib-only HTTP server (full model on host); loads both neffs once, then serves |
optb_server_slim.py | Low-RAM server — loads only the embedding/PLE tables on the host (bf16, ~6 GB) so it fits inf2.xlarge (16 GB). Same API. |
optb_gen.py | Minimal standalone greedy-generation example |
Dockerfile / Dockerfile.slim | Reproducible runtime images (full / slim) |
google/gemma-4-E2B-it — see License below.1docker pull xbill9/gemma4-optb:latest
2# on an AWS inf2 instance:
3docker run --rm -p 8080:8080 --device=/dev/neuron0 xbill9/gemma4-optb:latest
4# then: curl -s localhost:8080/healthdocker.io/xbill9/gemma4-optb (~16 GB each, Apache-2.0):latest / 512-128 — full server, for inf2.8xlarge (128 GB host RAM), ~44 tok/s.slim — low-RAM server for inf2.xlarge (16 GB host RAM), ~24 tok/s.1sudo fallocate -l 16G /swapfile && sudo chmod 600 /swapfile && sudo mkswap /swapfile && sudo swapon /swapfile
2docker run --rm -p 8080:8080 --device=/dev/neuron0 xbill9/gemma4-optb:slim/metrics endpoint
(requests, tokens, tokens/sec, errors, resident memory).inf2 instance with the Neuron runtime, plus:transformers==5.13.0
torch-neuronx==2.8.0.2.12.22436
neuronx-cc==2.23.6484.0
libneuronxla==2.2.15515.01export KV_MAX=512 KV_BUCKET=128 \
2 KV_PRE_OUT=./kv_pre_512.pt KV_DEC_OUT=./kv_dec_512.pt PORT=8080
3python optb_server.py # ~100 s warmup, then serves on :80801curl -s localhost:8080/health
2curl -s -X POST localhost:8080/generate \
3 -H 'content-type: application/json' \
4 -d '{"prompt":"What is AWS Inferentia?","max_tokens":64}'
5# also: /v1/chat/completions, /v1/completions, /v1/models1KV_MAX=1024 KV_BUCKET=256 \
2KV_PRE_OUT=./kv_pre_1024.pt KV_DEC_OUT=./kv_dec_1024.pt \
3python optb_kv.py traceneuronx-cc) and does not need a NeuronCore, so you can
recompile on any box. Larger buffers cost more device memory — fp32 neffs already exceed
the 16 GB core, which is why these ship as bf16.google/gemma-4-E2B-it, © Google, Apache-2.0.LICENSE/NOTICE when redistributing.