From-scratch LLM inference engine in Rust + CUDA — no frameworks, no ggml. Built for single-GPU serving on RTX 50-series Blackwell (sm_120a), every kernel written and tuned against measured hardware limits. Exactness is the contract: speculative, graph-replay, and batched serving output is gated token-identical to plain decode — speed never changes what the model says.
cargo build --release # arch auto-detected (120a / 90a / 100a / 89)
./target/release/run-gen hf:owner/repo:Q4_K_M --prompt "hi" # auto-downloads from Hugging FacePrebuilt Linux x86_64 binaries (sm_120a) ship with each release — full quick start below.
- Use memra when you serve one model on an RTX 50-series card and want measured, exactness-gated speed — tuned paths are the defaults, no flags needed.
- Use memra when your MoE doesn't fit VRAM — resident-if-fits expert residency with VRAM→RAM→NVMe spill.
- Use memra when you want published numbers you can audit — every cell is regression-tracked, re-measured on board-moving merges, raw run logs in-repo (research/).
- Use something else when you have another GPU (llama.cpp, mistral.rs) or need multi-GPU tensor-parallel serving (vLLM, SGLang).
cargo build --release
./target/release/kernel-check # every kernel vs CPU reference — must end ALL GREEN
MEMRA_CHAT=1 ./target/release/run-gen /path/to/model.gguf --prompt "Explain KV caches."
MEMRA_SPEC_K=3 ./target/release/run-spec /path/to/qwen36-27b.gguf # MTP speculative
./target/release/memra-server # OpenAI-compatible /v1run-gen prints its argmax gate (... MATCH) before generating — a MISMATCH voids every
number after it. The naked build is the tuned engine; flags exist only for runtime
parameters, machine config, and rollback seams (docs/FLAGS.md).
Performance is tracked as a regression suite, not a scoreboard: board-moving merges re-measure the tracked cells on the target rig and regenerate this section. Three representative samples — single-user decode on the RTX 5090, memra vs llama.cpp built on the same machine:
| Model / scenario | memra tok/s | llama.cpp tok/s | Ratio |
|---|---|---|---|
| Qwen3.6-35B-A3B MoE — plain decode (the flagship serving model, tg128 at 512-token context) | 187.0 | 164.9 | 1.13x |
| Qwen3.5-9B — MTP speculative, short / medium / long-agentic prompt classes (both engines at their measured best) | 281.0 / 211.7 / 187.1 | 122.2 / 121.5 / 117.7 | 2.30x / 1.74x / 1.59x |
| Qwen3.6-35B-A3B MoE — plain decode at 6.3k-token context (the short-context lead holds at depth) | 177.1 | 156.3 | 1.13x |
Measured 2026-08-02 on the RTX 5090 Laptop — same-session interleaved medians, same exact prompts; memra at its naked defaults, llama.cpp at its swept best (docs/COMPETITOR-SETUP.md). N, thermal regime, and the full boards: docs/PERFORMANCE.md; raw per-run logs: research/tune-data/.
The deployment bar for shipping a new model as supported: best-vs-best end-to-end at or above 1.1x llama.cpp on every prompt class on the target rig. Cells below it are tracked as open gaps, never hidden. Every tracked cell (wins, losses, and gaps), mechanism notes, and flip history: docs/PERFORMANCE.md.
Ten models ship supported on the 5090 — every one fully gated (kernel-check, argmax match, spec self-consistency). Per-model numbers stay in docs/PERFORMANCE.md; the tracked cells move with the tuning campaign in research/tune-data/.
| Model | Class | Quant | Drafter | Supported since |
|---|---|---|---|---|
| Qwen3.5-9B | dense | NVFP4 (5090), Q8_0 (H100) | MTP + own-gen trimmed draft | v0.1.0 |
| Qwen3.6-27B | dense | NVFP4, Q4_K_M MTP-baked | MTP + own-gen trimmed draft | v0.1.0 |
| Qwen3.6-35B-A3B | MoE | IQ4_XS | MTP + own-gen trimmed draft | v0.1.0 |
| Gemma-4 26B-A4B | MoE | QAT Q4_0 | MTP drafter | v0.23.0 |
| Gemma-4 31B | dense | QAT Q4_0 | MTP drafter | v0.35.0 |
| Gemma-4 E4B | dense | QAT Q4_0 | MTP drafter | v0.35.0 |
| Gemma-4 12B | dense | QAT Q4_0 | MTP drafter | v0.40.0 |
| Ornith-1.0-9B | dense | Q8_0 | own-gen donor-block draft | v0.63.0 |
| Ornith-1.0-35B | MoE | Q4_K_M | own-gen donor-block draft | v0.64.0 |
| Qwen-AgentWorld-35B-A3B | MoE | UD-IQ4_XS (avoid UD-Q4_K_M — its Q5_K expert mix sits outside fast-path coverage) | own-gen drafter | v0.66.0 |
In bring-up (running end-to-end, not yet over the bar): KAT-Coder-V2.5 (decode at llama parity, prefill gap open), Hy3 Layer103.5 spill overlay (5.13 tok/s served, tuning toward 10), MiniMax-M3 REAP50 (loads + generates, router tuning open). Receipts and bring-up notes: docs/PERFORMANCE.md#bring-up-notes.
OpenAI-compatible server (axum): batched decode, cross-request prefill batching,
speculative serving, /metrics. OpenAI tool calling (tools/tool_choice, streaming
tool_calls deltas, role:"tool" turns) rides the model's own chat template — zero engine
changes. Constrained decoding (response_format json_object/json_schema) applies the
grammar mask on device and keeps every fast path — device sampling, CUDA-graph decode,
speculative serving — at 99.4% of unconstrained speed. Cross-request prompt caching serves
repeated prompt prefixes without recomputing
them, reports the split as usage.prompt_tokens_details.cached_tokens on every response,
and namespaces all reuse per tenant via the request-level cache_salt field (vLLM-style).
The contract: greedy serving is isolated-identical under concurrent load — a request's
tokens never depend on its co-arrivals. Multi-GPU boxes serve as a replica fleet:
1,477 tok/s managed on 3×H100, chaos-tested (docs/SERVING.md).
- Decode kernels — NVFP4 / Q4_0 / Q8_0 / K-quant split-plane matvecs, warp-level dp4a, per-shape auto-dispatch.
- MTP speculative decoding — embedded draft head, one batched K+1 verify, adaptive depth + confidence cut.
- Hopper wgmma/TMA — FA3-class prefill attention, fused GDN chunks, grouped expert GEMM past cuBLAS.
- MoE on 24 GB — resident-if-fits residency, grouped-f16 expert prefill, SLRU spill cache, VRAM→RAM→NVMe reads.
- Quantized-KV attention — fused FlashAttention-class kernels, q8_0/q5_1/FP8-e4m3 per layer class.
- CUDA-graph decode — one replay per token, 4 bytes/token host traffic, per-session capture.
- Loaders — GGUF (memory-mapped), safetensors (modelopt NVFP4 byte-exact; official Qwen FP8 block-128 checkpoints load bit-exact).
Every kernel change passes kernel-check (CPU reference), the run-gen argmax gate, and
run-spec K=1..8 self-consistency — one command: tools/local-ci.sh (Hopper:
tools/validate-h100.sh). FP summation order is part of the contract; the perf CI
re-measures published cells on engine-touching pushes (CONTRIBUTING.md).
- 5090 prefill: 0.59–0.78x llama — llama benches NVFP4 prefill at W4A4 numerics memra's gates reject (docs/PERFORMANCE.md).
- H100 dense prefill trails vLLM per cell — the int8-GEMM dtype edge, mechanism-refuted (ARCHITECTURE-H100.md).
- Gemma plain margins are thin at the DRAM wall (1.02–1.06x); one spec cell at 0.98x.
- Hy3 spill serves at 5.13 tok/s (N=3 median), tuning toward 10 (docs/HY3-SPILL.md).
- NVIDIA RTX 50-series (sm_120a, primary target RTX 5090 Laptop), H100 (sm_90a), B200 (sm_100a, compile-gated), or Ada (sm_89, portable eval). Other GPUs compile untuned — use llama.cpp or mistral.rs there.
- H100 (sm_90a) support: the arch auto-detects at build (
MEMRA_CUDA_ARCHoverrides), Hopper-only kernel promotions are compile-gated so the naked sm_120a build stays byte-identical, andtools/validate-h100.shis the one-command gate battery. Evidence ledger: ARCHITECTURE-H100.md; tracked per-model numbers: docs/PERFORMANCE.md. - CUDA 13.1 (+12.8 dual-toolkit, ARCHITECTURE.md); Rust edition 2024; a GGUF or HF safetensors model.
- One GPU per engine process — no tensor parallelism yet (pipeline-parallel seam merged, default off); multi-GPU boxes serve as a replica fleet today.
- Moving research codebase; APIs and flags change without notice.
- docs/PERFORMANCE.md — the full tracked boards, mechanism notes, flip history, per-model receipts.
- ARCHITECTURE.md — sm_120a tech stack + feasibility ledger.
- ARCHITECTURE-H100.md — the H100 evidence ledger.
- docs/FLAGS.md — the audited flag catalog.
- docs/COMPETITOR-SETUP.md — competitor engines at their peak.
- docs/DRAFT-REGIME.md — the standard drafter pipeline.
- docs/SERVING.md — the serve surface: fleet runbook, tools, constrained decoding, prompt caching, tenant isolation.
- docs/HY3-SPILL.md — Hy3 spill runbook.
- HANDOVER.md — living state-of-work.
- research/ — every experiment as JSONL; research/benchmarks.md is the protocol.
Issues and PRs welcome — see CONTRIBUTING.md. Correctness gates run on
real GPUs (tools/local-ci.sh); GitHub CI is compile-only. Running memra on your own rig?
A hardware validation report is the
fastest way to help.
MIT — see LICENSE.