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: every kernel is bit-audited against a CPU reference, speculative and graph-replay decode are gated token-identical to plain decode, and batched serving is gated byte-identical to the same request served alone — speed never changes what the model says.
curl -fsSL https://raw.githubusercontent.com/avifenesh/memra/main/tools/install.sh | sh # prebuilt, checksum-verified
run-gen hf:owner/repo:Q4_K_M --prompt "hi" # auto-downloads from Hugging FaceAlso on crates.io (cargo install memra-server — compiles the CUDA kernels locally) and
as tarballs on each release — full
installation and 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).
Prebuilt binaries (recommended — self-contained, kernels embedded, no CUDA toolkit needed):
curl -fsSL https://raw.githubusercontent.com/avifenesh/memra/main/tools/install.sh | shThe script picks your GPU arch via nvidia-smi, verifies the sha256 against the
release's SHA256SUMS, and installs memra-server, run-gen, run-spec, and
kernel-check to ~/.local/bin (MEMRA_INSTALL_DIR overrides). Same binaries, other
routes: cargo binstall memra-server, or grab a tarball from
releases.
From source / crates.io (requires the CUDA 13.1 toolkit — build.rs compiles the kernels with nvcc, arch auto-detected from your GPU):
cargo install memra-server # or: git clone && cargo build --releaseRequirements table:
| Prebuilt binaries | cargo install / source |
|
|---|---|---|
| GPU | RTX 50-series (sm_120a, tuned target), H100 (sm_90a), Ada (sm_89, untuned) | + B200 (sm_100a, compile-gated) |
| NVIDIA driver | >= 580 (CUDA 13 runtime support) | same |
| CUDA | runtime libraries only (cudart, cublas, cublasLt) | full CUDA 13.1 toolkit (nvcc) |
| OS / libc | Linux x86_64, glibc >= 2.35 | Linux x86_64 |
| Rust | not needed | stable, edition 2024 (rust-version 1.85) |
Model weights are not bundled — they pull from Hugging Face on first use (below).
Three commands from install to an OpenAI-compatible endpoint:
curl -fsSL https://raw.githubusercontent.com/avifenesh/memra/main/tools/install.sh | sh
kernel-check # every kernel vs CPU reference — must end ALL GREEN
MEMRA_MODELS="chat=hf:owner/repo:Q4_K_M" memra-server # downloads the GGUF, serves /v1 on :8080Or generate directly without the server:
MEMRA_CHAT=1 run-gen hf:owner/repo:Q4_K_M --prompt "Explain KV caches."
MEMRA_SPEC_K=3 run-spec /path/to/qwen36-27b.gguf # MTP speculative decodingrun-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). The llama.cpp column is a frozen reference recorded through 2026-08-03 (benching stopped that day). 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 on every prompt class on the measuring rig. Cells below it are tracked as open gaps, never hidden. The llama.cpp denominators are frozen reference points recorded through 2026-08-03 — head-to-head benching stopped that day (owner call) and forward work is self-competition, so treat the ratios above as regression anchors rather than a live scoreboard. 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 — available under constraint, at 99.4% of unconstrained speed on the
plain decode lane (the speculative lane pays more: 79%). 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).
Multi-key bearer auth maps keys to tenants (file-backed keyring + CLI, hot-reload, per-key
rate limits and lane class) so cache isolation and metering follow real tenant identity;
session affinity resumes a conversation whose history the client rewrote (think-stripping
agent clients) — identity nominates, bytes decide, so time-to-first-token stays flat as the
conversation grows instead of re-priming everything each turn.
The gateway-facing surface is complete for OpenRouter-style listing: /v1/models with
per-model metadata, X-RateLimit-* headers on every response, and graceful drain on
SIGTERM (in-flight requests finish, new ones get a 503 + Retry-After). Official FP8
safetensors checkpoints serve directly — block-128 weights load bit-exact (dequantized onto
the Q8_0 arm; the win is a 2.89x faster load, not faster decode) and speculative decoding
runs out of the box on the checkpoint's own MTP head, no GGUF conversion step. GGUF remains
the primary delivery format and Q8_0 the shipping 8-bit arm; the FP8-native compute paths
are implemented, bit-exact, and default off until they clear the deployment bar
(docs/FLAGS.md).
Streaming is round-cadence (one flush per spec-round commit, not one per burst) and a
request arriving mid-burst preempts at the next round boundary, so felt latency no longer
scales with the speculative burst size: solo first text ~0.12 s with ~27 ms chunk gaps, and
contended first text 0.12–0.15 s — the solo class — instead of waiting a full burst out
(docs/SERVING.md).
Admission is VRAM-aware and 64 concurrent speculative clients is a gated property, not
an aspiration: tools/serve-stress-gate.sh runs in local CI and asserts 64/64 well-formed
streams on a 24GB card (peak 23.1 GB), with a teeth arm that forces a broken headroom
reserve and proves the gate still catches it. Requests should send an explicit max_tokens
— omitting it sizes the KV ladder from the context ceiling and strands measurable VRAM
(docs/SERVING.md).
The contract: greedy serving is isolated-identical under concurrent load — a request's
output tokens are byte-identical whether it arrives alone or inside a full batch, gated by
replaying the same prompts at c=1 and c=16 against the same server and byte-comparing every
stream — and chunk-size-invariant: chunked prefill produces bit-identical logits across
MEMRA_PRIME_CHUNK values (one canonical greedy output per prompt, gated by the chunkinv
battery arm). (That is the scoped claim; it is not an identity claim against a single-token
reference decode — the batched-plain path has a documented, bounded near-tie flip class, and
speculative decode is gated self-consistent per run-spec K=1..8 on MTP-capable artifacts.)
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).
- The NVFP4 speculative serve path trails its bare-CLI twin (−8.66%, pre-fix measurement) — the spec tier's burst loop is a separate path from the solo fast path that closed the plain-serve c=1 gap (serve c=1 now runs the same m=1 fused trunk as the CLI, +5–8% decode-only, bit-identical; docs/SERVING.md). Open lane.
- Cold time-to-first-token on short turns trailed llama.cpp on the local laptop rig (0.53 s vs 0.19 s, same model file, N=5 interleaved — a pre-fix measurement) alongside short-agentic decode and raw prefill, while long-generation sampled decode leads by +17%. Most of that latency stack is since fixed: the per-request session realloc under VRAM pressure (evict-first + right-size ladder), rewritten-history resume via session affinity (TTFT flat at any depth), per-burst SSE emission (round-cadence streaming: solo first text 0.41 → 0.12 s), and the contended admission wait (0.15 s at any burst size instead of scaling with it). Still open: the prefill wall and short-context acceptance (0.55 vs 0.73). The head-to-head number itself has not been re-measured (competitor benching is stopped by doctrine).
- NVIDIA RTX 50-series (sm_120a — the primary optimized target; the tracked 5090 boards are measured on an RTX 5090 Laptop, and RTX PRO 6000 Blackwell Workstation numbers come from rented pods), 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.