genie-ai-runtime: a Jetson edge LLM inference engine that beats llama.cpp — designing for unified memory

A field report on building genie-ai-runtime — a from-scratch CUDA inference engine for the Jetson Orin Nano Super 8 GB — and then asking the harder question the engine can’t answer on its own: on a memory-starved edge box, should the model even be a transformer?

Table of contents

0. TL;DR

This post has two halves, because the project did.

1. The engine. genie-ai-runtime is a single-binary, single-model CUDA LLM runtime tuned for exactly one machine: the Jetson Orin Nano Super 8 GB (SM 8.7, 102 GB/s, 67 TOPS). On the home-agent workload it was built for — short prompts, one model, sharing 8 GB of unified memory with always-on voice — it runs Qwen3-4B-Q4_K_M prefill at ~38 tok/s vs llama.cpp’s ~18 tok/s on the same box: +115%. The wins are Orin-specific kernels (a tensor-core MMQ Q4_K prefill GEMM on SM 8.7), a memory-first design (INT8 KV cache, pre-allocated pools, an OOM-guard that accounts every byte before inference starts), and a persistent + in-memory prefix KV cache that makes a stable system prompt effectively free (warm-turn TTFT 444 ms, a shared-prompt re-request ~13× faster prefill).

2. The model. The engine makes a transformer fast. But the transformer’s KV cache grows linearly with context, and on an 8 GB box that’s the thing that eventually kills you. So I benchmarked the architectures head-to-head (GeniePod/genie-claw#410) — pure transformer (Qwen) vs pure SSM (Mamba) vs attention/Mamba-2 hybrid (Falcon-H1) — across 0 → 32 k context, all Q4_K_M on llama.cpp, same Jetson. The result: quality is a matter of training, not architecture; the SSM/hybrid families win long-context speed; only pure Mamba wins on memory; and KV memory scales with layers, not the “hybrid” label — so the shallow Falcon-H1-0.5B is the genuine edge pick while the 66-layer “Deep” variant is a 2.4× memory trap that OOMs at 32 k.

Both halves are public: the runtime repo (and its path-by-path ROADMAP.md), the assistant it serves, and the benchmark issue.

1. The constraint: an 8 GB box that’s already busy

genie-claw is a local, private home assistant: wake word → speech-to-text → an LLM that turns “turn the kitchen lights down” into a tool call → text-to-speech, all on-device, nothing leaving the house. The LLM is one tenant on a Jetson Orin Nano Super 8 GB that is already running parakeet STT, TTS, a denoiser, and a Home Assistant container. The model gets the leftovers of 8 GB of unified (CPU+GPU shared) memory — and that leftover is contested and bursty: a TTS spike or an STT burst can claim it at any moment.

What “unified memory” means on the Jetson — and why it changes everything. On a discrete GPU (A100, H100, RTX 4090), the GPU has its own physical memory (HBM2e, GDDR6X) that is separate from host DRAM. You allocate buffers in host RAM, explicitly copy them to the GPU over PCIe, and the GPU’s kernels run at HBM bandwidth. Two separate allocators, two separate bandwidth pools, an explicit copy step between them.

The Jetson Orin Nano has none of that. Its CPU and GPU cores share one physical LPDDR5 pool — the same DRAM chips, the same memory bus, the same 102 GB/s bandwidth. There is no PCIe link between CPU and GPU memory. There is no cudaMemcpy(H2D) step; model weights loaded by the CPU are already accessible to CUDA kernels without a copy, because there is only one memory. This is what “unified memory” means on Jetson — not the CUDA virtual unified memory API, but a single physical DRAM pool.

This has three direct implications for an inference runtime:

1. No copy overhead, but shared bandwidth. Weight loads, KV reads, audio buffers, OS page tables, and STT/TTS model reads all compete for the same 102 GB/s. A parakeet inference burst or a TTS flush steals from the same bus the GPU decode kernels are reading model weights from. On a discrete GPU this competition doesn’t exist — HBM and PCIe/system RAM are physically separate.

2. One memory allocator, one budget. There is no separate “GPU memory” that can OOM independently of “CPU memory.” When a cudaMalloc fails on Jetson it means the same pool is full that the OS and STT and TTS are all using. Any runtime that treats GPU memory as a separate isolated resource will unexpectedly compete with the rest of the system.

3. Weights are zero-copy. After the model is loaded from NVMe into memory-mapped pages, CUDA kernels can read weights directly from those pages without a host-to-device copy. The model load time (~1.3 s warm / ~30 s cold) is purely NVMe→DRAM, not NVMe→DRAM→VRAM.

The existing options don’t fit this shape:

• llama.cpp — portable, excellent, generic CUDA kernels. It’s what genie-claw runs today. But it’s Jetson-blind: no awareness of the shared-memory budget, no Orin-specific kernel path, tuned for the general case rather than short-prompt home-agent traffic.
• TensorRT-LLM — fast, but datacenter-shaped (A100/H100). Assumes the GPU has its own memory budget, separate from host RAM. Too heavy for a unified-memory iGPU on an 8 GB box.

So the thesis of genie-ai-runtime is narrow on purpose: one machine, one model, one GGUF, one shared unified-memory budget. Give up all portability and you can do things a portable runtime can’t — pin to SM 8.7, pre-allocate the KV and scratch pools from the shared pool, and account every byte across the entire system so a concurrent voice burst can’t OOM the box mid-sentence.

2. The target hardware (and the roofline that dictates strategy)

The ridge point is the whole story. At ~0.66 op/byte, the two phases of inference land on opposite sides of the roofline:

• Prefill (process the prompt, N tokens at once) is compute-bound — there’s arithmetic intensity to exploit, so the lever is the matmul kernel.
• Decode (generate one token at a time, batch 1) is bandwidth-bound — every step streams the whole model from LPDDR5, so you’re capped by 102 GB/s and there’s little to optimize beyond reading weights efficiently.

That split decides where the engineering goes: prefill is where a better kernel pays off, decode is where you’ve mostly already lost to physics — so on decode you stop chasing tok/s and instead win on memory and on not prefilling at all.

The unified memory decode ceiling — what physics actually allows

On a discrete GPU with HBM2e (A100: ~2 TB/s) the decode bandwidth ceiling is high enough that kernel quality matters a lot. On Jetson LPDDR5 at 102 GB/s shared, you can calculate the hard ceiling before writing a line of code:

Qwen3-4B Q4_K_M weights: ~2.3 GB
KV cache @ 1024 ctx (INT8): ~72 MB
Total bytes streamed per decode token: ~2.37 GB

Theoretical ceiling: 102 GB/s ÷ 2.37 GB/token ≈ 43 tok/s

The engine measures ~10 tok/s in practice — ~23 % of the theoretical ceiling. The gap is not bad kernels; it’s the structure of unified memory:

• Shared bandwidth contention. OS page table walks, STT resident process, TTS subprocess, Home Assistant container — all sharing the same 102 GB/s. In steady state, the usable GPU bandwidth for decode is closer to 70–80 GB/s than 102 GB/s.
• LPDDR5 latency vs. HBM2e. HBM is on-package with much lower access latency. LPDDR5 on Jetson has higher per-access latency that matters at batch=1 decode, where requests are fine-grained rather than bulk-streamed.
• KV cache growth. At 4096 ctx (72 MB INT8 KV) the weight-read dominates. At 16k ctx (~576 MB INT8 KV) the KV term starts materially reducing the ceiling: 102 / (2.3 + 0.576) ≈ 35 tok/s theoretical, and with contention, actual decode at 16k is closer to 6–7 tok/s.
• Activation memory traffic. Per-token activations, residual buffers, and intermediate tensors add read/write traffic on top of the weight and KV reads.

This is why speculative decoding (EAGLE-3, §9) is the correct next lever — not kernel tuning. At batch=1, the decode wall is bandwidth, not compute. Verifying K draft tokens in one forward pass means one weight-read pays for up to K generated tokens: effective decode throughput can exceed the single-token bandwidth ceiling because the per-token weight cost drops by 1/K. That’s the one optimization the physics actually allows.

It also explains why the unified memory architecture makes the prefill wins matter more than they would on a discrete GPU. At 102 GB/s, the system prompt costs real latency — a 500-token system prompt is ~2.5 GB of weight reads = ~25 ms just from bandwidth. The prefix KV cache (§5) that eliminates that re-prefill on warm turns saves a wall-clock second per command, not microseconds.

3. Beating llama.cpp on the same model: the prefill story

The runtime grew through a sequence of numbered “Paths” — each an umbrella GitHub issue with per-phase PRs, each PR posting a verified on-Jetson result before merge. The full narrative is in ROADMAP.md; here’s the prefill arc that took it from “first coherent tokens” to past llama.cpp.

(A parallel decode path, Path C — reading four packed Q4_K bytes as one uint32 on the hot GEMV kernels — lifted decode +21 %; it’s in §4, since decode is a memory story, not a prefill one.)

Net: prefill 8.2 → 38 tok/s (+363 %), and against the runtime it replaces:

One honesty caveat up front, because it matters. This is the short-prompt regime — an 18–33-token prompt, which is what a home command actually looks like, and the case genie targets. Prefill throughput rises steeply with prompt length (a bigger batch feeds the tensor cores better), so the +115% is a small-N result, not a claim that genie wins at every batch size: I tuned and measured the short prompts the product lives in, and I don’t make claims about the large-batch regime, where llama.cpp’s mature batched kernels are strong. Same short-prompt regime, same hardware and model, genie is ~2× ahead.

Why batched prefill (Path B) over-delivered

The bandwidth math under-predicted every PR in the Path B series by 3–10×. Two effects beyond raw weight bandwidth explain it:

1. Kernel-launch overhead. Per-token prefill ran ~250 launches/token × 18 tokens ≈ 4500 launches; the layer-major batched path runs ~10 launches/layer × 36 layers ≈ 360. At ~5–10 µs/launch that’s tens of milliseconds reclaimed.
2. L2 pressure. Per-token, each token streams the full weight set before the next starts, so the LPDDR5 controller re-fetches the same weight pages 18×. Layer-major touches each page once and keeps activations hot in L2 across the per-layer loop.

The tensor-core kernel (Path E) — the big win

Path E replaces the scalar Q4_K batched kernel’s CUDA-core FMAs with mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 on SM 8.7. A 16×32 FP16 staging tile is dequantized from one Q4_K block into shared memory; the 128 (d, dm) scale pairs are precomputed once per block; and 4 warps per CUDA block share one dequantized A-tile across 4 contiguous N-stripes — 32 tokens at a time vs 8 for the single-warp variant — amortizing the dequant cost 4×. ~39 registers/thread, ~407 GFLOPS aggregate on Qwen3-4B prefill shapes. The ~2.36× speedup over the scalar fallback held uniform across prompt sizes N = 33 / 88 / 235 (a 7× range).

A detail I’d flag to anyone doing this: Path E reorders float adds (the mma accumulation order differs from scalar FMAs), so it is not byte-identical to the scalar path. The bar I adopted instead was “sensibly-identical” — character-for-character the same generated text at temp=0 on the reference prompt, with drift bounded by the FP16 ULP. Insisting on byte-equality would have frozen the project on scalar kernels forever.

4. Memory-first: making every byte count

On this box, the runtime’s defining feature isn’t a kernel — it’s that memory is accounted before any inference runs. A MemoryBudget totals the model weights, the KV pool, and scratch; an OOMGuard refuses to start (cleanly) rather than crashing the Jetson when the sum won’t fit; KV and scratch are pre-allocated pools, so there is no allocation on the hot path that a concurrent voice burst could lose a race to.

Why this matters specifically on unified memory. On a discrete GPU, cudaMalloc fails independently of host memory — the GPU’s VRAM allocator and the CPU’s heap are separate, and an OOM in one doesn’t immediately affect the other. On Jetson’s unified memory, every cudaMalloc, every malloc, every OS page fault, every STT model activation all draw from the same 8 GB pool. If parakeet loads a second model for a background task, your KV cache allocation fails. If the OS buffers a large NVMe write, your scratch buffer allocation stalls. There is no isolation boundary.

The MemoryBudget accounts for this explicitly: it knows the approximate OS floor (~1.5 GB), the STT footprint (~450 MB), the TTS footprint (~100 MB), and the agent process overhead (~200 MB), and it sizes the KV pool to fit in 8 GB − fixed_overheads − weights − scratch. The pool is pre-allocated at startup — if it won’t fit, the process exits with a clean error message rather than crashing mid-inference when the first cudaMalloc on the hot path returns NULL. This is the correct design for unified memory: account the whole system budget upfront, not just the GPU “half.”

The biggest single memory lever was Path I: an INT8 KV cache (per-(layer, position, kv-head) absmax-scaled). It halves the KV body — 144 MB → 72 MB at Qwen3-4B, 1024 context (plus a ~2.25 MB scales region, ~74 MB total) — and is sensibly-identical to FP16 at temp=0 (per-element drift ~0.79 %). On unified memory, INT8 KV has two distinct wins that a discrete GPU’s separate memory accounting would only half-capture:

1. Headroom. 70 MB freed in the shared 8 GB pool is headroom for the whole system — not just a larger KV allocation but also more room for the STT model’s activation buffers, the TTS subprocess, or a longer context window. The freed memory is genuinely available to everything.
2. Bandwidth. Smaller KV means less LPDDR5 traffic per decode token. The KV read per token drops from ~144 KB/token (FP16, 1024 ctx) to ~72 KB/token (INT8) — a per-token bandwidth savings that directly raises the effective decode ceiling from the calculation above. At 16k context (576 MB → 288 MB INT8) the difference in theoretical ceiling is 102/(2.3+0.576) ≈ 35 tok/s vs 102/(2.3+0.288) ≈ 39 tok/s — a ~10% decode speedup from storage format alone, before any kernel change.

It’s been the default since alpha.12.

Here’s the v1.0 scorecard the engine ships with:

Decode did get one real kernel win — Path C, reading four packed Q4_K quants as a single uint32 load on the hot GEMV kernels (Wo, the gate/up pair, the QKV triple) and folding the residual add into the GEMV itself, for +21 % decode (7.5 → 9.1 tok/s). But that’s where decode largely stopped: it’s bandwidth-bound at 102 GB/s, there’s little kernel headroom left, and a later push (Path G) measured neutral and was paused (more in §6). So raw decode throughput is deliberately not on the brag list — at ~10 tok/s it’s competitive but not the win. The honest framing is the one the ROADMAP itself uses: the engine wins on prefill and TTFT; on decode it competes on memory and reuse, not on tok/s.

5. Making the system prompt free: persistent + prefix KV

A home agent re-sends almost the same prompt every time — a fat system prompt (tool list, device catalog, persona) plus a few new words. Prefilling that system prompt on every command is the dominant latency, and it’s pure waste. Two features kill it:

• Path F — persistent KV cache. Per-conversation KV state is saved to disk at turn end and hydrated at turn start; the file stores only the used_tokens positions per layer with the prompt token IDs inline, so a follow-up turn finds its longest common prefix and skips prefill for the matched range. Saves are atomic (.tmp + fsync + rename), truncation-checked, and fingerprinted (an FNV-1a-64 of the GGUF refuses a cache built against a different model). In a 2-turn conversation sharing 24/36 = 67 % of tokens, warm-turn TTFT dropped 857 → 444 ms (−48 %).
• In-memory prefix cache (PR #85). Requests are serialized through the engine and nothing zeroes the KV between them — so the buffers still hold the previous request’s tokens. The engine now tracks that resident token sequence and reuses the longest common prefix with the next request, with zero disk I/O, even across unrelated conversations that merely share a system prompt. Two requests sharing a 229-token system prompt: prefill 6880 ms → 526 ms, ~13×; for a ~495-token prompt the win is ~16 s → ~1 s.

Both are correct for the obvious worry: the shared prefix sits at identical absolute positions, so RoPE is consistent, and causal attention only reads [0, pos] — stale KV beyond the new sequence is never read. It’s the same invariant the cold path already relied on.

A unified memory detail that makes the persistent cache faster than on discrete. On a discrete GPU, saving a KV cache to disk requires a cudaMemcpy(D2H) to bring KV tensors from GPU VRAM to CPU memory before writing to a file. Loading requires the reverse: read from disk into CPU RAM, then cudaMemcpy(H2D) into VRAM. On Jetson unified memory neither copy exists — the KV pool is already in memory the CPU can directly write() or read() back. The save path is a direct write(fd, kv_ptr, size) — one syscall, zero GPU-CPU copy overhead. The load path hydrates directly into the pre-allocated KV pool at CPU speed and is immediately visible to CUDA kernels. The unified architecture turns a conceptually two-copy operation into a single memcpy-to-disk, which is why the 48% TTFT reduction in Path F comes with almost no CPU overhead on the warm path.

6. The negatives I kept

The path-based process made it cheap to record dead ends, and I think the dead ends are the most useful part for anyone building on the same hardware:

• Path G (decode throughput): neutral, paused. The decode-side analogue of Path E. G1 (Q6_K scale hoist) and G2-lite measured neutral-to-slightly-negative — the compiler was already hoisting those loads implicitly. Closed as no-ops; the path stays open for a different attack (a Q6_K block-layout redesign for vectorized loads). Decode is bandwidth-bound, so this is the expected hard wall.
• Path H (speculative decoding): deferred. Feasibility-audited and shelved: with Qwen3-4B as the target and no good small draft model in the Qwen3 family, the win is bounded by the draft’s quality on home-agent traffic. The EAGLE-3 plan in §9 is how I intend to come back to it — a trained draft head sidesteps the missing draft model.
• Split-K K-quant GEMV and CUDA-Graphs decode capture: slower / no-op. Both closed without merge — the residual-fused single-kernel GEMV beat split-K, and the graph capture didn’t pay for itself at this scale.
• An honest re-baseline that erased a “win.” The alpha.11 cycle re-measured a claimed --int8-kv speedup same-day on the same prompt and machine state and found it was measurement noise, not improvement — so the claim was rolled back and the broken flag was made warn-and-ignore. Re-measuring the prior release every release (instead of trusting cross-day numbers) caught it.

7. The deeper question the engine can’t answer: which architecture?

genie-ai-runtime makes a transformer fast on this box. But step back: the transformer’s KV cache grows linearly and unboundedly with context, and on a memory-starved 8 GB board the KV cache — not the weights — is what eventually OOMs you. That’s an architecture problem, not an engine problem. So before committing the edge stack to transformers forever, I ran the matchup (genie-claw#410): pure transformer vs pure SSM vs hybrid, swept across context depth.

Setup: all Q4_K_M, on llama.cpp’s CUDA backend (so this is a fair architecture comparison, not an engine one), on the same Orin Nano 8 GB, GeniePod services stopped and page-cache dropped so the GPU was uncontended. llama-bench for throughput, llama-completion for the allocator’s own memory report. A fair same-size triad anchors it: Qwen2.5-1.5B (transformer), Mamba-1.4B (pure SSM), Falcon-H1-1.5B-Instruct (parallel attention + Mamba-2 hybrid).

7.1 Decode throughput vs context depth — the headline

tg64 @ dN = decode tok/s with N tokens already in context:

Prefill (pp512) tells the same story even louder: the transformer collapses 8.7× (1264 → 145 tok/s) as attention re-scans a growing context, pure Mamba is flat (~980) on its O(1)-per-step state update, and the hybrid decays 3.4×.

7.2 Memory vs context — and the surprise

The allocator’s own report, transformer vs SSM:

At 32 k the transformer’s KV cache is 63× larger than Mamba’s entire state, and still climbing. This is the real edge argument for SSMs: not speed, memory that doesn’t grow.

The surprise was the hybrid. You’d expect “attention + Mamba-2” to inherit the SSM’s flat memory. It doesn’t — at least not at 1.5 B. Falcon-H1 runs attention heads and Mamba-2 heads in parallel inside every block, so it keeps full GQA attention in all 24 layers: its KV cache still grows ~linearly (768 MiB at 32 k), and the 73 MiB recurrent state adds back what little the layer count saved. Total at 32 k: 841 MiB vs the transformer’s 896 — only ~6 % less. The hybrid buys a flatter speed curve, not a smaller memory footprint at this scale.

7.3 KV memory scales with layers, not the “hybrid” label

Sweeping the Falcon-H1 family exposed the actual governing variable. KV ∝ layers × head_dim — so the “hybrid saves memory” story holds only when the hybrid is also shallow:

And the shallow 0.5B isn’t just the lightest — it’s the fastest decoder at every depth (54.5 → 21.2 tok/s, beating all four larger models even at 32 k), in a 300 MB file, at full quality. The Deep variant is the mirror image: depth buys reasoning-per-parameter and pays in KV memory.

7.4 Quality is training, not architecture

Four greedy prompts (trivia, arithmetic, instruction-following, explanation):

The pure Mamba fails instruction-following and loops forever — classic base-model behavior. But that’s a confound, not a verdict on SSMs: mamba-1.4b-hf is a 2023 base model; everything it’s compared against is a 2024+ instruct model. The clean read is that every instruct-tuned model — transformer and hybrid alike — scores 4/4. Quality is a function of training maturity, not of attention-vs-SSM. There simply is no well-tuned small Mamba instruct to ship today.

7.5 The ceiling on 8 GB

Going up the size ladder: Falcon-H1-3B runs and is 4/4, but OOMs trying to allocate a 32 k context (1.76 GiB weights + ~1 GB KV + compute buffer don’t fit). Falcon-H1-7B and the 7B reasoning variant won’t load at all — the ~4.2 GB weight buffer can’t be allocated contiguously on an 8 GB unified box, even with everything else stopped. The practical Falcon-H1 ceiling on this Jetson is 3B, short context.

8. The verdict

Two decisions came out of all this:

• The engine: genie-ai-runtime replaces llama.cpp for genie-claw — same model, ~2× short-prompt prefill, a memory model that won’t OOM the shared box, and a prefix cache that makes the system prompt free. The final flip from llama-server is gated on a 24-hour soak (#7); the engine is otherwise at v1.0.0.
• The model: keep Qwen3-4B as the deployed model today — best quality, and the engine makes it fast enough in the budget. If you want a long-context hybrid on this Jetson, Falcon-H1-0.5B-Instruct is the credible pick: fastest decode at every depth, smallest footprint, 4/4 quality, most headroom alongside the voice stack. Avoid the Deep variant at the edge — its depth makes it the most memory-hungry thing tested. Pure Mamba owns the efficiency extremes (flat decode, 63× less memory) but has no shippable small instruct, so you can’t actually deploy that combination yet.

The meta-point: an inference engine and a model architecture are two different optimization surfaces. The engine got Qwen3-4B from 8 to 38 tok/s on a board it was never designed for — but no kernel can fix the fact that a transformer’s KV cache grows with every token. That ceiling is architectural, and on the edge it’s the one that actually decides what you can run.

9. What’s next

v1.0 is the floor, not the ceiling. Three threads are queued, each aimed squarely at a wall this post ran into.

Speculative decoding via EAGLE-3, trained for the home

§6 deferred speculative decoding for a concrete reason: with Qwen3-4B as the target and no good small draft model in the Qwen3 family, the speedup was bounded by a draft I didn’t have. EAGLE-3 removes that blocker — it doesn’t need a separate draft model; it trains a lightweight autoregressive draft head on top of the target’s own hidden features (multi-layer feature fusion), so the “draft” reuses the model already resident in memory. That’s the one lever that actually beats the decode bandwidth wall from §4: verify K speculated tokens in a single weight-read forward pass instead of K of them — exactly the right trade on a 102 GB/s box where decode is otherwise pinned at ~10 tok/s and the GPU is idle waiting on HBM.

The edge twist is to train the head on smart-home context. A home assistant’s traffic is a narrow distribution — a fixed system prompt, a known device catalog, a small grammar of tool calls. A draft head fit to that distribution should accept far more aggressively than a general-purpose draft, because the next token is usually highly predictable (turn, the, a {tool_name}, a device id straight out of the catalog). And because the agent runs greedy (top-k = 1) — deterministic tool calls, not creative prose — verification collapses to exact-match on the target’s argmax: the simplest, highest-acceptance regime there is. Domain-tuned head + greedy verify is the most promising route to finally moving the decode number this engine deliberately didn’t chase. The head itself will live in genie-ai-model.

Two more model targets: Falcon-H1 and Gemma 4 E2B

genie-ai-runtime is transformer-only today, and the next two models stretch it in different directions — one in architecture, one in modality — while both stay inside the same 8 GB memory discipline.

• Falcon-H1 — the long-context play. §7’s benchmark named the long-context winners: the SSM/hybrid families, with Falcon-H1-0.5B the fastest decoder at every depth and the lightest deployable footprint — but measured on llama.cpp. Bringing it onto genie-ai-runtime means adding Mamba-2 / hybrid-block kernels (parallel attention + SSM state) to a runtime that only knows transformers. The payoff: a flatter decode curve and 16 k+ context inside the budget, on our Orin-tuned kernels — turning the §7 “credible alternative” into a deployable one.
• Gemma 4 E2B — multimodal input, structured output, on-device footprint. A different reason entirely, and the one I’m most excited about. Today’s engine is text-only; Gemma 4 E2B is the path to three things a home agent actually needs:
  • Multimodal input. It’s a vision-language model, so a camera frame becomes a query — “is the garage door open?”, “read this label”, “who’s at the door?” — which a text-only Qwen3-4B simply can’t answer. That turns the assistant from voice-only into voice-plus-sight.
  • Structured output. The agent’s whole job is emitting tool calls. A model whose JSON / function-call output is reliable by default — not salvaged by a post-hoc parser — makes the grammar of the home (turn on, set to, query state) land cleanly, which is where small models usually wobble.
  • On-device efficiency. This is why it’s E2B, not the full model: the “effective-2B” line keeps resident memory near a 2 B model’s despite a larger raw parameter count — exactly the footprint discipline an 8 GB box already shared with ASR + TTS demands. The one thing to engineer around is its large (~262k) vocabulary, which makes the final projection memory-bandwidth-bound on Orin — a natural fit for the engine’s on-device-output / vocab-reduction paths.

Different directions, one constraint: every new model still has to earn its place inside the 8 GB budget the voice stack already shares.

10. The process that made it work

Two habits I’d keep for any constrained-hardware project:

1. Path-based umbrella issues. Every numbered Path was one issue with a phase table and risks up front, and every per-phase PR posted a verified on-device result before merge. It made negative results cheap to absorb without losing momentum.
2. “Sensibly-identical,” not byte-identical. Once tensor-core kernels reorder float adds, byte-equality is gone. Adopting FP16-ULP-bounded, character-identical output as the bar is what unlocked Path E and INT8 KV.

References

• genie-ai-runtime — the runtime; see ROADMAP.md for the path-by-path narrative and PR #85 for the prefix cache.
• genie-claw#410 — the transformer vs Mamba vs hybrid benchmark (all measurements on-device).
• genie-claw — the on-device home assistant the runtime serves.
• Falcon-H1 (TII) — the attention + Mamba-2 hybrid family.
• NVIDIA Jetson Orin Nano — the target hardware.
• llama.cpp — the baseline runtime, and the engine behind the architecture benchmark.