From-scratch LLM inference engine in Rust + CUDA — no frameworks, no ggml. One codebase serves two architectures, auto-detected at build time: RTX 50-series Blackwell (sm_120a), tuned single-user against llama.cpp, and H100 Hopper (sm_90a), measured model-by-model against vLLM on the same box. Every kernel is hand-written against measured hardware limits, and exactness is the contract: speculative and graph-replay output is gated token-identical to plain decode, so speed never changes what the model says.

Use memra when you serve one model on an RTX 50-series card and want measured, exactness-gated speed, or you want a single-GPU H100 engine whose wins and losses against vLLM are published per model. Use something else when you have another GPU (llama.cpp, mistral.rs) or need multi-GPU serving (vLLM, SGLang).

Standing (2026-07-31): seven supported models on the 5090, all fully gated; every MTP-spec cell is at or above 1.13x llama.cpp (up to 2.2x), plain cells sit at the DRAM wall or above. On the H100, a full per-model board against vLLM 0.26: end-to-end wins on 4 of 6 models (1.14–1.81x), decode wins on 5 of 6 (the 35B MoE decode loss flipped to a 1.07x win this release); the two e2e losses (0.92x, 0.83x) are MoE expert-prefill cells with the remaining kernel-rate rung mapped. Every number is a same-session interleaved measurement on a real-text prompt with the argmax exactness gate green; trimmed MTP drafter heads are published ready-to-use at huggingface.co/Avifenesh/memra-bench.

Running memra on your own rig? A hardware validation report is the fastest way to help.

cargo build --release   # MEMRA_CUDA_ARCH auto-detected from the GPU (120a / 90a / 100a / 89)

The build probes the GPU's compute capability and selects the arch; MEMRA_CUDA_ARCH overrides. At startup the engine verifies the binary matches the device and fails early with a rebuild hint otherwise (MEMRA_ARCH_CHECK=0 bypasses). Hopper-only promotions (wgmma/TMA kernels, graph serving defaults) are compile-gated — the naked sm_120a build is byte-for-byte the tuned 5090 engine.

The full per-model board against vLLM 0.26 on 1×H100 80GB, same box, same session (2026-07-31). One number per arm: end-to-end tok/s — 512 tokens generated on a ~2100-token real-text prompt, single request, total wall time (N=5 medians, argmax exactness gate green on every published row). Cross-artifact by design: vLLM serves what H100 users deploy (w8a8 / FP8-dynamic / bf16 HF checkpoints — it rejects these GGUFs); memra serves its GGUF artifacts.

Wins on exact math (the bf16-row wins carry a quant-advantage caveat — those vLLM arms move 4x the weight bytes). Decode alone wins 5 of 6 cells (1.07–1.85x; the 35B MoE decode loss flipped to a win when its per-layer shared-expert gate left cuBLASLt). The two e2e losses are MoE expert-prefill cells: both MoE models already route their expert GEMMs through an int8-MMA expert kernel, and that kernel's ~16 TF rate — 60x off the CUTLASS int8 roofline — is the single mapped rung between here and those rows flipping. The dense-model prefill gaps are the int8-GEMM dtype edge — a Q8_0-exact int8 GEMM is mechanism-refuted on Hopper (per-32-block rescale costs 5.4x naive / 17x pipelined; ptxas serializes cross-bank GMMA register reads), so crossing it means w8a8-class numerics that change model outputs — an accuracy-bar decision with measured receipts, not an engineering unknown.

Shipped on this build: FA3-class prefill attention (TMA swizzled ring + wgmma, 4.8x the mma kernel), fused wgmma GDN chunk kernels with varlen twins, Q4_0/Q8_0→fp16 prefill mirrors on the cuBLASLt lane, cross-request prefill batching, per-session CUDA-graph decode, and the Hopper wgmma toolkit (cu/wgmma_common.cuh) — canonical core-matrix pairings probed for bf16/tf32/s8: one byte-geometry, three MMA kinds.

• Evidence ledger (every verdict and refutation): ARCHITECTURE-H100.md
• Flags + promoted defaults: docs/FLAGS.md §7
• One-command battery: tools/validate-h100.sh <model.gguf> [--quick]

Prebuilt Linux x86_64 binaries (sm_120a) ship with each release — or build from source:

cargo build --release
./target/release/kernel-check                     # every kernel vs CPU reference
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/run-gen hf:owner/repo:Q4_K_M --prompt "hi"        # auto-download from HF
./target/release/memra-server                      # OpenAI-compatible /v1

kernel-check must end with ALL GREEN, and run-gen prints its argmax gate (... MATCH) before generating — a MISMATCH voids every number after it. Tuned paths are the defaults; flags exist only for runtime parameters, machine config, and rollback seams (docs/FLAGS.md).

Measured 2026-07-27 on the target rig (RTX 5090 Laptop, N=2+ medians, both engines interleaved in the same thermal window on the same rig, same exact prompts, no flags (tuned paths are defaults); plain/depth rows from the 2026-07-09 validity-gated cold-start rebaseline, spec rows re-paired 2026-07-18, Gemma card rows from the 2026-07-15 best-vs-best re-audit. Full per-run logs: research/tune-data/ (Qwen) and research/gemma4-bringup/ (Gemma) — every win and every loss; Gemma 12B plain row from the 2026-07-24 official N=5 cell stamp (research/gemma4-bringup/g12tg-cellstamp.log)) against llama.cpp built on the same machine, same exact prompts, both engines re-baselined the same day. Boards move with the tuning campaign — research/tune-data/rig5090.jsonl is the running record; the README is refreshed with every board-moving merge.

Plain decode (no speculation, tg128 at 512-token context):

Depth is part of the contract: at 6.3k-token context every lead holds (1.02–1.07x).

Speculative decoding (MTP head, both engines at their measured best):

The three columns are three prompt classes: short code / medium code (greedy) / long agentic (temp 0.7, distribution-exact rejection sampling). One asterisk: the 35B short-code llama bar (236.5) rode an EOS-margin flip and is not a clean win basis. Every spec row uses one trimmed draft built by the standard regime (docs/DRAFT-REGIME.md); prebuilt drafts live in the bench repo, or build your own:

./target/release/frspec-owngen model.gguf ranks.gguf 32768        # ranks from the model's OWN generations
tools/make-trimmed-draft.sh model.gguf ranks.gguf.txt draft.gguf  # extract + trim + quantize

Same protocol, own campaign log (research/gemma4-bringup/rig5090-gemma4.jsonl); cells re-paired best-vs-best with llama's own draft depth swept per cell. Highlights (full tables and the per-cell archaeology live in the campaign log):

What buys the margins: an FP8 (e4m3) KV cache, occupancy-tuned attention tiles, wide-load Q4_0 expert dots, FR-Spec drafter trims (150 MB → 18 MB at unchanged acceptance), a per-model adaptive-K floor (+20% on 12B/31B spec at unchanged exactness), and an in-round draft-confidence cut. One near-parity cell remains (26B 1.7k spec, 0.98x). Exact prompts and llama.cpp's swept-best flags: docs/COMPETITOR-SETUP.md.

• 5090 prefill trails llama.cpp (0.59–0.78x), root-caused: llama benches NVFP4 prefill at W4A4 (FP4 activations), a numeric class memra's exactness gates reject. Output quality outranks the prefill column.
• H100 prefill trails vLLM on every model (12–70% by cell) while decode wins 5 of 6 — the e2e wins come from decode dominating the p2048/g512 shape. MoE cells are gated on the expert-GEMM kernel rate (mapped rung); dense cells on the int8-GEMM dtype edge, refused by the accuracy laws (full refutation ledger in ARCHITECTURE-H100.md).
• Gemma plain margins are thin where both engines sit at the DRAM wall (1.02–1.06x).
• Hy3 native spill serves at a 5.13 tok/s N=3 median, tuning toward 10 (docs/HY3-SPILL.md).

• NVFP4 / Q4_0 / Q8_0 decode — split-plane repacked matvecs, warp-level dp4a, int8 tensor-core prefill GEMM, per-shape auto-dispatch.
• MTP speculative decoding — embedded draft head, one batched K+1 verify, adaptive draft depth + confidence cut; K=1..8 self-consistency gate.
• Hopper wgmma/TMA kernels — FA3-class prefill attention, fused GDN chunk family, canonical descriptor pairings for bf16/tf32/s8.
• MoE on 24 GB — expert-major CSR batching, frozen SLRU expert residency, bounded host LRU, mirrored positioned reads across VRAM→RAM→NVMe.
• Quantized-KV attention — fused FlashAttention-class kernels (q8_0/q5_1 or FP8-e4m3 KV per layer class), split-K, graph-replayable device-length counters.
• CUDA-graph decode — one graph replay per token, 4 bytes/token host traffic; per-session capture for serving.
• Serving — OpenAI-compatible server with batched decode, cross-request prefill batching, KV prefix reuse, speculative serving, and a flat /metrics endpoint.
• Loaders — GGUF (memory-mapped) and safetensors (modelopt NVFP4 byte-exact).

Every kernel change passes, in order: kernel-check (CPU reference), the run-gen argmax gate, run-spec K=1..8 self-consistency — one command: tools/local-ci.sh. FP summation order is part of the contract. Exactness gates are blind to numeric shifts where decode and verify move together, so the perf CI (tools/local-ci.sh --perf) re-measures every published cell per engine-touching push and tracks speculative acceptance per cell against a rolling baseline, enforced by the pre-push hook. On Hopper, tools/validate-h100.sh is the equivalent one-command battery.

• NVIDIA GPU: RTX 50-series (sm_120a, primary target RTX 5090 Laptop), H100 (sm_90a), B200 (sm_100a, compile-gated), or Ada (sm_89, portable eval).
• CUDA 13.1 (plus 12.8 for the dual-toolkit build in ARCHITECTURE.md; MEMRA_NVCC overrides the nvcc path). Rust edition 2024, cudarc 0.19.
• A model: GGUF or HF safetensors directory.

• Tuned for two exact memory/compute ratios (5090 Laptop, H100 SXM). Other GPUs compile via the portable arch but are untuned — use llama.cpp or mistral.rs there.
• Single GPU; no tensor parallelism. Batched decode serving on both architectures.
• Moving research codebase; APIs and flags change without notice.

• ARCHITECTURE.md — sm_120a tech stack + feasibility ledger.
• ARCHITECTURE-H100.md — the H100 build's full evidence ledger.
• HANDOVER.md — living state-of-work.
• docs/FLAGS.md — the audited flag catalog.
• docs/COMPETITOR-SETUP.md — competitor engines at their peak.
• docs/HY3-SPILL.md — Hy3 spill runbook.
• research/ — every experiment as JSONL, wins and losses both; research/benchmarks.md is the measurement protocol.

Issues and PRs welcome — see CONTRIBUTING.md. Correctness gates run on real GPUs (tools/local-ci.sh); GitHub CI is compile-only.

MIT — see LICENSE.