DEV Community: xbill

Comparing Qwen and Gemma on the Low End: Qwen Is Faster Without Drafting, Gemma Has Better Quality

xbill — Sun, 18 Oct 2026 20:20:20 +0000

This article provides a step by step deployment guide for Qwen3.5-2B to a local, desktop and laptop hosted GPU enabled system, served by llama.cpp beside Gemma 4 E2B on the same cards. A suite of Python MCP tools is built to simplify management of the llama.cpp hosted deployment.

https://github.com/xbill9/qwen-dev/tree/main/local-llamacpp-3050oem-qwen352b-q4_0

Qwen3.5-2B decodes 1.12x as fast as Gemma 4 E2B on an RTX 3050 and prefills 1.24–1.31x as fast. On three tasks, with thinking off for both models, Gemma answers better on every one: its 4-bit build leads Qwen's 4-bit build by 15.5 to 26.0 points and Qwen's unquantized bf16 by 10.0 to 19.5. With each model's drafter turned on, the speed lead disappears. The gaps survive scoring every cut-off or tool-less Qwen answer as right.

What is this project trying to Do?

The rigs in this repository serve Gemma 4 E2B from local hardware: an RTX 3050 OEM in a desktop and a GTX 1650 Ti in a laptop. Both run Google's quantization-aware-trained (QAT) 4-bit build, and both have measured scores for it on the same three tasks.

The question here: what does a Qwen of the same size class give on the same cards? Qwen3.5-2B counts 1.943B parameters in its text GGUF. Everything below runs both models with the same flags and the same task code, and on one llama.cpp binary except where the Scope paragraph says otherwise.

At This Point You Should Have…

An NVIDIA GPU with a CUDA driver — here an RTX 3050 OEM (8 GB, Ampere, driver 615.78.08) and a GTX 1650 Ti (4 GB, Turing)
llama.cpp built with CUDA — commit 7182ef71e here
Python 3 for the converter and the MCP server — the system python3, no virtualenv
The repository cloned, and local-llamacpp-3050oem-qwen352b-q4_0/ as your working directory
Claude Code, or any MCP client that speaks stdio

The Two Models

	Qwen3.5-2B	Gemma 4 E2B
Layers	24: three of every four Gated DeltaNet (linear attention), every fourth full attention	35
Parameters, as llama-bench counts them	1.943B	4.629B, most of it per-layer embeddings
Vocabulary	248,320, tied input and output	262,144
4-bit release	none; no QAT checkpoint and no GGUF published	Google's QAT Q4_0
Drafter	an MTP layer inside the checkpoint	a separate 4-layer assistant model

Only Qwen's 6 full-attention layers hold a KV cache. The 18 DeltaNet layers keep a fixed-size state per request, so Qwen's prefill cost grows more slowly with prompt length.

Step 1 — Build the Qwen Q4_0 File

Qwen publishes bf16 weights only, so the rig builds its own 4-bit file. make quantize converts the Hugging Face checkout to a bf16 GGUF and runs the standard llama-quantize Q4_0, with no importance matrix:

python3 ~/llama.cpp/convert_hf_to_gguf.py ~/models/Qwen3.5-2B --outtype bf16 --outfile Qwen3.5-2B-bf16.gguf
~/llama.cpp/build/bin/llama-quantize Qwen3.5-2B-bf16.gguf Qwen3.5-2B-Q4_0.gguf Q4_0

The result, from the rig's tpu.env:

#   convert_hf_to_gguf.py --outtype bf16   -> 3,897,388,000 B, 335 tensors
#   llama-quantize ... Q4_0 (no imatrix)   -> 1,236,741,088 B
# llama-quantize's standard Q4_0 mix: 194 matrices Q4_0 (806.4 MB), token_embd
# Q6_K (417.2 MB, which is also the output head since it is tied), 140 F32 norms

This file is plain rounding from bf16. Gemma's Q4_0 file holds weights Google trained to sit on the 4-bit grid. Step 7 measures what that difference is worth.

Step 2 — Serve It

make serve runs llama-server with the Gemma rig's flags, so the two serve the same 131,072-token window:

llama-server -m Qwen3.5-2B-Q4_0.gguf --host 0.0.0.0 --port 8080 -ngl 99 -c 131072 \
  -ctk f16 -ctv f16 -fa 1 -t 6 -tb 12 --parallel 1 --reasoning off --metrics

Qwen3.5 thinks by default. --reasoning off puts answers in content, as on the Gemma rig; a request can still turn thinking on with "chat_template_kwargs": {"enable_thinking": true}.

At the full context it uses 2995 MiB of the 3050's 8192, and 2951 MiB of the 1650 Ti's 4096. The server log shows the MTP layer loaded and left idle:

0.01.266.683 W model has unused tensor blk.24.attn_norm.weight (size = 8192 bytes) -- ignoring
0.01.266.690 W model has unused tensor blk.24.post_attention_norm.weight (size = 8192 bytes) -- ignoring
0.01.266.699 W model has unused tensor blk.24.attn_q.weight (size = 4718592 bytes) -- ignoring

Step 3 — Check That llama.cpp Runs Qwen Correctly

A low score could come from the runtime. Three checks cover the prompt going in, the probabilities coming out and the tool calls Qwen writes.

Prompts. Every prompt was rendered through llama-server's /apply-template and through transformers' apply_chat_template. All 400 tool-calling prompts, with their tools, and 100 math prompts came out byte-identical.

Next-token probabilities. The reference is transformers in float32 on the CPU over Qwen's bf16 weights, on 65 prompts from the quality suite:

Against transformers fp32	Same chosen answer	Answer-probability difference, median
transformers bf16, GPU	61 of 65	0.0128
llama.cpp bf16 GGUF	62 of 65 🥇	0.0027 🥇
llama.cpp Q4_0 GGUF	43 of 65	0.1136

llama.cpp's bf16 file sits closer to the float32 reference than transformers' own bf16 path does. The drift comes from the 4-bit rounding.

Tool calls. Every BFCL reply that made no tool call was searched for text that looks like a call the parser could have missed (tool_call, <function, a JSON object): 0 of 79 at Q4_0 and 0 of 48 at bf16. Those replies solve the problem in prose, such as working the Pythagorean theorem step by step.

Step 4 — Measure Speed on One Binary

llama-bench ran both models on the same build, card and hour, five cells alternating between models, each started with the GPU at 40 °C or below:

llama-bench -p 512,2048 -n 128 -r 5 -ngl 99 -fa 1 -t 6

The summary script's output on the RTX 3050:

pp512: qwen 5193.44  gemma 4203.68  qwen/gemma 1.235x  qwen fa1/fa0 1.042x  qwen rep/first 0.993x  gemma rep/first 0.997x
pp2048: qwen 5316.13  gemma 4050.45  qwen/gemma 1.312x  qwen fa1/fa0 1.078x  qwen rep/first 0.996x  gemma rep/first 0.999x
tg128: qwen 132.84  gemma 118.30  qwen/gemma 1.123x  qwen fa1/fa0 1.012x  qwen rep/first 1.001x  gemma rep/first 1.002x

	RTX 3050	GTX 1650 Ti
Qwen decode, tok/s	132.84	85.26
Gemma decode, tok/s	118.30	81.17
Qwen / Gemma, decode	1.12x	1.05x
Qwen / Gemma, prefill at 512 / 2,048 tokens	1.24x / 1.31x	1.43x / 1.51x

Qwen's prefill rises with prompt length (5193 to 5316 t/s) while Gemma's falls (4204 to 4050). That fits the DeltaNet layers, whose cost does not grow with the prompt.

Step 5 — Measure Answer Quality

Three tasks, the same task code and settings as the Gemma runs, 8 requests at a time, greedy:

Suite: 13 classification sets, 3,880 questions, scored from the next-token probabilities
GSM8K: 1,319 math word problems, 2,048-token limit
BFCL simple: 400 tool-calling requests

Qwen's suite prompts end in Qwen's own chat template with thinking closed, and its answers are read from Qwen's own tokens. Every pair of runs is compared question by question, with a 95% range from resampling.

RTX 3050	Gemma E2B QAT Q4	Qwen Q4_0	Qwen − Gemma	95% range
Suite	0.683 🥇	0.422	−0.260	−0.281 to −0.240
GSM8K	0.898 🥇	0.743	−0.155	−0.177 to −0.132
BFCL simple	0.925 🥇	0.710	−0.215	−0.265 to −0.165

Every range is far from zero. On the suite, Qwen is ahead only on boolq and squad2.

Where Qwen loses points:

Option letters. Asked for A, B or C, it often writes the option's name instead. On multinli it answers neutral on 298 of 299 questions.
One-sided answers. It calls 248 of 300 comments toxic, where the gold labels mark 32.
Long math answers. 100 GSM8K replies hit the 2,048-token limit, against 5 for Gemma. With thinking off it still deliberates inside the answer.
Prose instead of tool calls. 79 of its 116 wrong BFCL answers made no tool call, mostly by solving the problem in text. Gemma skipped the tool 6 times.

Step 6 — Score Qwen Unquantized

Qwen's 4-bit file is plain rounding, so part of the gap may be the rounding. The bf16 GGUF from Step 1 ran through the same three tasks on the same binary, the same day:

RTX 3050	Qwen Q4_0	Qwen bf16	Gemma E2B QAT Q4
Suite	0.422	0.488	0.683 🥇
GSM8K	0.743	0.798	0.898 🥇
BFCL simple	0.710	0.765	0.925 🥇

Rounding to Q4_0 costs Qwen 5.5 to 6.5 points per task, every range clear of zero.
Gemma's 4-bit build still leads unquantized Qwen by 19.5, 10.0 and 16.0 points. The rounding accounts for a quarter to a third of the gap; the rest is the model.
The same habits show at bf16: most comments called toxic, 48 BFCL replies with no tool call, 52 GSM8K replies at the length limit.

Step 7 — Measure What QAT Is Worth to Gemma

Gemma's 4-bit file comes from quantization-aware training; Qwen's comes from rounding. To compare like with like, Gemma E2B's ordinary bf16 release was rounded to Q4_0 with the exact recipe used for Qwen, and scored on the GTX 1650 Ti beside the other three:

GTX 1650 Ti	Gemma QAT Q4_0	Gemma rounded Q4_0	Qwen bf16	Qwen Q4_0
Suite	0.681 🥇	0.632 🥈	0.488	0.430
GSM8K	0.901 🥇	0.881 🥈	0.807	0.734
BFCL simple	0.917 🥇	0.895 🥈	0.762	0.688

Share of Qwen Q4_0's gap to Gemma QAT	Suite	GSM8K	BFCL
Whole gap	−0.251	−0.167	−0.230
QAT's part	−0.048	−0.020	−0.022
QAT's share	19%	12%	10%

QAT is worth 2 to 5 points to Gemma at 4 bits. Rounded Gemma still beats unquantized Qwen on all three tasks, by 14.4, 7.4 and 13.3 points.

Is the Comparison Fair to Qwen?

Four objections a Qwen user will raise, each checked against the per-question results.

Gemma is the bigger model. llama-bench counts 4.629B parameters for Gemma E2B against 1.943B for Qwen, most of Gemma's in per-layer embeddings. By file size, Qwen's bf16 GGUF is 3.90 GB and Gemma's QAT Q4_0 is 2.61 GB, and the larger Qwen file loses every task to the smaller Gemma one.

The token limit cut Qwen's math answers off. Scoring every GSM8K reply that hit the 2,048-token limit as right:

RTX 3050, GSM8K	Scored	Every cut-off reply right	Gemma E2B QAT Q4
Qwen Q4_0	0.743	0.812	0.898 🥇
Qwen bf16	0.798	0.836	0.898 🥇

The tool-calling task punishes prose answers. Scoring every BFCL reply with no tool call as right:

RTX 3050, BFCL simple	Scored	Every tool-less reply right	Gemma E2B QAT Q4
Qwen Q4_0	0.710	0.908	0.925 🥇
Qwen bf16	0.765	0.885	0.925 🥇

The suite scores answer format. Partly right. On three subsets Qwen gives one answer almost every time or writes the option name in place of the letter: civil_comments, summeval-consistency and multinli. With those removed, the bf16 gap to Gemma on the GTX 1650 Ti shrinks from 19.3 points to 11.2. The suite is the most format-sensitive of the three results; Gemma still leads on it.

Q4_K_M and importance-matrix builds of Qwen were not scored. The bf16 file is the most any 4-bit build can recover, and Gemma leads it on every task.

Step 8 — Run the Bf16 File on a 4 GB Card

Qwen's bf16 GGUF holds 3.62 GiB of weights, and the 1650 Ti has 4 GiB. With every layer on the card, the 501.52 MiB compute buffer fails to allocate. Moving the last 6 of 24 blocks to the CPU is the fewest that fits:

llama-server -m Qwen3.5-2B-bf16.gguf -ngl 99 -ot 'blk\.(18|19|20|21|22|23)\.=CPU' \
  -c 8192 -ctk f16 -ctv f16 -fa 1 -t 6 -tb 12 --parallel 1 --reasoning off --metrics

GTX 1650 Ti	Decode, tok/s	Prefill at 512, t/s	GPU memory
Q4_0, all on the card	85.11 🥇	511.43 🥇	2951 MiB at 131k context
bf16, 6 blocks on the CPU	20.22	490.87	3207 MiB at 8k context

Decode drops to 0.24x and prefill holds at 0.96x. The split costs no score: bf16 on this card scored 0.488, 0.807 and 0.762, level with the 3050.

Step 9 — Draft With Qwen's MTP Head

The Q4_0 file already carries Qwen's multi-token-prediction layer. llama.cpp drafts from it with --spec-type draft-mtp and no extra file:

llama-server -m Qwen3.5-2B-Q4_0.gguf -ngl 99 -c 8192 -fa 1 \
  --spec-type draft-mtp --spec-draft-n-max 1

The log confirms it and reports acceptance per request:

0.02.045.088 I common_speculative_init_result: creating MTP draft context against the target model '/home/xbill/models/Qwen3.5-2B-gguf/Qwen3.5-2B-Q4_0.gguf'
0.04.830.818 I slot print_timing: id  0 | task 8 | draft acceptance = 0.91729 (  122 accepted /   133 generated), mean len =  1.92

Six prompts (code, chat, story, explain, facts, summary), one request at a time, geometric mean against no drafting on the RTX 3050:

Drafted tokens	1	2	3	4	6
Greedy	1.149x 🥇	1.145x	0.998x	0.852x	0.647x
Sampled, temperature 0.8	1.086x 🥇	1.047x	0.920x	0.782x	0.596x

1 drafted token is the setting to serve. Code gains most (1.25x greedy); story least (1.09x).
Greedy replies stay byte-identical to no drafting up to 3 drafted tokens on the 3050.
On the same card, Gemma's separate drafter gave 1.39x at 2 tokens.

Step 10 — Race the Two Drafters on One Card

On the GTX 1650 Ti, both models on one binary, one request at a time, each with its best drafter setting of 1 token:

GTX 1650 Ti, greedy tok/s	No drafter	1 drafted token
Qwen Q4_0, own MTP head	82.2 🥇	91.0 (1.107x)
Gemma E2B, QAT drafter	77.9	92.1 (1.181x) 🥇
Qwen / Gemma	1.054	0.988

Sampled at temperature 0.8 the result is the same: 86.1 against 90.7 tok/s. Gemma's drafter gains more, so the two tie.

Drafting leaves Qwen's answers level: GSM8K 0.748 and BFCL 0.698 with the MTP head, against 0.734 and 0.688 without. At 8 requests at once, each task took 1.08–1.09x as long with drafting.

🔎 Tip: Pass Qwen's Tokenizer to Probability-Scored Tasks

A task that scores answers from next-token probabilities has to look up the answer labels in the model's own vocabulary. Qwen's differs from Gemma's, so give the scorer --tokenizer Qwen/Qwen3.5-2B; the default Gemma ids would read the wrong tokens.

🔎 Tip: Two Hosts, Two Different Hashes

The 3050 host and the 1650 Ti host built Qwen's Q4_0 from the same revision on the same llama.cpp commit, and got files of the same size with different SHA-256s (1c27608b… and bd99d5ae…). The bf16 files differ too, so the change enters at conversion. They pick the same answer on 62 of 65 prompts and score within 2.2 points on every task. Pin a hash per host, and keep one file if you need byte-identical results across machines.

Compare and Contrast

	Qwen3.5-2B Q4_0	Gemma 4 E2B QAT Q4_0
Decode, RTX 3050	132.84 tok/s 🥇	118.30 tok/s
Decode with drafter, GTX 1650 Ti	91.0 tok/s	92.1 tok/s 🥇
Suite / GSM8K / BFCL, RTX 3050	0.422 / 0.743 / 0.710	0.683 / 0.898 / 0.925 🥇
4-bit build	rounded locally	QAT, published by Google 🥇
Drafter	built in, no extra file	separate 59 MB file

So, Which One?

Gemma 4 E2B. Qwen3.5-2B decodes faster without a drafter and prefills faster on prompts up to 2,048 tokens. Gemma answers better on all three tasks measured, at 4 bits against Qwen's bf16, and with both drafters on the speed difference is gone.

Qwen's built-in MTP head is a convenience: drafting with no extra file, and greedy text unchanged at 1 drafted token. Choose Qwen when prefill speed matters more than the answers; how its lead grows past 2,048-token prompts is unmeasured here.

Summary

The goal of this article was to serve Qwen3.5-2B with llama.cpp beside Gemma 4 E2B on two budget GPUs and compare them on speed and answer quality. The key to the solution was running both models on one llama.cpp binary with the same flags and the same task code, and scoring Qwen at bf16 and Gemma rounded the way Qwen's file is built. The results were:

🟢 llama.cpp runs Qwen3.5-2B correctly: identical prompts to Qwen's template, and its bf16 probabilities sit closer to float32 than transformers' bf16 does.
🟢 Qwen is faster without drafting: 1.12x decode and 1.24–1.31x prefill on the RTX 3050, 1.05x and 1.43–1.51x on the GTX 1650 Ti.
❌ Qwen scores lower on all three tasks: 15.5 to 26.0 points behind Gemma at 4 bits, 10.0 to 19.5 behind at bf16.
🟢 The gaps hold in Qwen's favour: scoring every cut-off GSM8K reply and every tool-less BFCL reply as right still leaves Qwen behind.
⚠️ Rounding is a quarter to a third of Qwen's gap, and QAT is a tenth to a fifth of Gemma's lead.
🟢 Qwen's MTP head drafts with no extra file: 1.15x greedy at 1 token, greedy text unchanged.
⚠️ With both drafters on, the two tie on speed: 91.0 against 92.1 tok/s on the GTX 1650 Ti.

Scope: one RTX 3050 OEM (8 GB) and one GTX 1650 Ti (4 GB), llama.cpp 7182ef71e for every Qwen run and for Gemma on the 1650 Ti, measured 2026-10-16 and 2026-10-17. The Gemma quality figures on the 3050 come from 2026-10-13 on llama.cpp 77c4540ce; Gemma's runs on the 1650 Ti on 7182ef71e scored level with them. Speed cells ran 3 to 5 repeats; each quality task ran once per file. Both models ran with thinking off and greedy decoding throughout; Qwen thinks by default, and thinking on was not run for either. The three tasks are classification, grade-school math and simple tool calling: no coding, Chinese, instruction-following or long-context task. Prefill was measured up to 2,048 tokens and every Qwen run served one request at a time, except the 8-request quality runs. No importance-matrix, Q4_K_M or Q8_0 build of either model was scored. The scoring code and the two cards were set up for this repository's Gemma work.

The strategy for using MCP for serving and comparing Qwen3.5-2B and Gemma 4 E2B on local GPUs was validated with an incremental step by step approach.

References

Can a QAT 12B Fit an 8 GB GPU? Gemma 4 12B on an RTX 3050

xbill — Sun, 18 Oct 2026 14:10:29 +0000

This article provides a step by step guide to serving Gemma 4 12B, Google's quantization-aware trained (QAT) build, on an 8 GB NVIDIA RTX 3050 with llama.cpp, with every layer on the GPU and Google's multi-token prediction (MTP) drafter for speculative decoding. A suite of Python MCP tools is built to simplify management of the llama.cpp hosted deployment.

This is the third part of a short series. Part 1, Gemma 4's MTP Drafter, Explained, covers how the drafter works and how the exact Q4_0 repack is built. Part 2, Do Drafters Work on Small Models?, measures E2B and E4B on this card and a GTX 1650 Ti. This part moves up to the 12B.

https://github.com/xbill9/gemma4-dev

Can a QAT 12B fit an 8 GB GPU? Yes, with every layer on the card and room for a working context:

RTX 3050 OEM, 8 GB	Google's 12B Q4_0 GGUF	exact Q4_0 rebuild
File size	6.98 GB	6.72 GB
Most context that loads	28,672 tokens	40,960 tokens
Generation	27.7 tok/s	28.6 tok/s
With the 12B drafter	—	56.4 tok/s
Most context with the drafter, 4-bit cache	—	81,920 tokens

The rebuild stores every weight on the 4-bit grid QAT trained, including the embedding table that Google's file keeps at 6 bits. That one table frees 247.5 MiB of the card, which buys 12,288 more tokens of context. Google's 12B drafter, rebuilt the same way, doubles generation, and a 4-bit context cache gives it 81,920 tokens of context with no change in answers.

What Is Being Measured?

Gemma 4 12B is a dense model with 48 layers and 11.9 billion parameters. Google publishes it already trained for 4-bit weights, as safetensors (google/gemma-4-12B-it-qat-q4_0-unquantized) and as a llama.cpp GGUF (google/gemma-4-12B-it-qat-q4_0-gguf). At 4 bits its weights take about 6.7 GB, which leaves an 8 GB card very little for anything else.

The questions are whether the whole model fits on the card, how much conversation fits beside it, how fast it generates, and what Google's 12B drafter adds. The card is the same RTX 3050 OEM that Part 2 measured with E2B and E4B.

At This Point You Should Have…

An NVIDIA GPU with 8 GB and a CUDA driver — here an RTX 3050 OEM, driver 615.78.08
The CUDA toolkit, 13.4 here, and cmake
llama.cpp at a commit with Gemma 4 and --spec-type draft-mtp — every number here is from 7182ef71e
Python 3.10 or newer, the system python3, no virtualenv, with numpy
About 40 GB of free disk for the source checkpoint and the build
The repository cloned, with the rig local-llamacpp-3050oem-12b-q4_0/
Claude Code, or any MCP client that speaks stdio

How the 12B Differs From E2B and E4B

	E4B	12B
Layers	42	48
Hidden size	2,560	3,840
Per-layer embedding table	yes, read from disk	none
Layers that own a KV cache	24 of 42	all 48
Sliding window	512	1,024
Generation on the 3050, no drafter	61.2 tok/s	28.6 tok/s

E2B and E4B carry a large per-layer embedding table that llama.cpp reads a few rows at a time from the memory-mapped file, so much of their file never reaches the card. The 12B has no such table: almost the whole file goes to the GPU. Its embedding table doubles as the output layer, so it sits on the card as well.

The 12B generates at 0.47x E4B's speed on the same card and llama.cpp build.

Step 1 — Size the Context Cache Before Downloading

The context cache (KV cache) decides how much conversation fits. Gemma 4 12B mixes two kinds of attention layers:

8 full-attention layers, each with one key/value head of width 512. At 16-bit precision that is 8 × 512 × 2 (key and value) × 2 bytes = 16 KiB per token.
40 sliding-window layers, each with 8 key/value heads of width 256, which only look back 1,024 tokens. llama.cpp holds a fixed 1,536 positions for them (the window plus one 512-token batch): 480 MiB, whatever the context.

So 40,960 tokens of context cost 640 + 480 MiB. Those are the figures llama.cpp's allocation log prints:

I load_tensors:        CUDA0 model buffer size =  6390.19 MiB
I llama_kv_cache:      CUDA0 KV buffer size =   640.00 MiB
I llama_kv_cache:      CUDA0 KV buffer size =   480.00 MiB
I sched_reserve:      CUDA0 compute buffer size =   161.52 MiB

With 6.4 GB of weights on an 8 GB card, every 247 MiB saved on the weights is about 15,800 tokens of context (arithmetic, at 16 KiB per token).

Step 2 — Rebuild the GGUF on Its Trained Grid

Google's 12B GGUF is 6.131 GB of Q4_0 and 0.826 GB of Q6_K. The Q6_K is a single tensor, token_embd, the 262,144 × 3,840 embedding table, at 787.5 MiB.

QAT trained that table on a 4-bit grid like everything else: in each block of 32 values, every value is a block step times an integer from −8 to 7. Q4_0 stores one fp16 step and a 4-bit level per value, the same shape. gguf_exact.py reads each block's trained step off its own values (the largest magnitude divided by 8, 7, … 1, whichever puts every value on an integer level, refined by least squares) and writes every weight matrix as Q4_0. It copies Google's metadata, tokenizer and chat template byte for byte.

hf download google/gemma-4-12B-it-qat-q4_0-unquantized --revision b6ed862 --local-dir src
hf download google/gemma-4-12B-it-qat-q4_0-gguf gemma-4-12b-it-qat-q4_0.gguf --revision 29d0977 --local-dir gguf
python3 gguf_exact.py src gguf/gemma-4-12b-it-qat-q4_0.gguf gemma-4-12B-it-q4_0-exact.gguf

It is the same script that built the E2B and E4B files, unchanged: the 12B's tensor names already match. It rebuilt 329 tensors in 17 minutes, and every block landed on the grid. Per tensor kind, 96.77–96.88% of values match Google's bf16 weights bit for bit once rounded back to bf16; the rest differ by the fp16 rounding of the step.

	Google's GGUF	exact rebuild
File	6,975,879,296 B	6,716,356,736 B
`token_embd`	Q6_K, 787.5 MiB	Q4_0, 540.0 MiB
Weights on the card	6,637.69 MiB	6,390.19 MiB

🔎 Tip: Google's 12B GGUF Is Already Text Only

The vision encoder ships as a separate file, mmproj-gemma-4-12b-it-qat-q4_0.gguf (167 MiB). Leave it out and the model is text only, as the rebuild is. On an 8 GB card those 167 MiB are another 10,700 tokens of context (arithmetic).

Step 3 — Fetch the Published Files

Both rebuilds are on Hugging Face:

xbill9/gemma-4-12B-it-qat-q4_0-exact-gguf, file gemma-4-12B-it-q4_0-exact.gguf, SHA-256 038337ee0375f602…
xbill9/gemma-4-12B-it-qat-assistant-q4_0-exact-gguf, the drafter, file gemma-4-12B-it-qat-assistant-q4_0-exact.gguf, SHA-256 f22f6366ca76102a…

The rig's make download and make draft fetch them and check the hashes in tpu.env:

make download
make draft
/home/xbill/models/gemma-4-12B-it-qat-assistant/gemma-4-12B-it-qat-assistant-q4_0-exact.gguf: OK

Step 4 — Find the Most Context That Fits

fit.sh starts llama-server -ngl 99 -fa 1 at rising context sizes and records the largest that loads, with the card otherwise idle:

RTX 3050, 8,192 MiB	16-bit cache	8-bit cache (`q8_0`)	4-bit cache (`q4_0`)
Google's GGUF	28,672 tokens	not tested	not tested
exact rebuild	40,960 tokens	65,536 tokens	131,072 tokens
exact rebuild + drafter	24,576 tokens	49,152 tokens	81,920 tokens

Google's file fails at 32,768 tokens; the rebuild loads there with 556 MiB to spare and fails at 49,152. The 8-bit cache halves every cache figure, sliding layers included, and the 4-bit cache halves it again. Step 10 measures what that costs.

The model's own limit is 131,072 tokens. Without the drafter, the 4-bit cache reaches it on this card.

Step 5 — Measure Speed Against Google's File

bench.sh runs llama-bench, five repetitions, in four passes that alternate which file goes first, with the card cooled to 40 °C before each:

pass	file	prompt, 512 tokens	generation, 128 tokens
1	Google	981.57 tok/s	27.68 tok/s
2	exact	987.49 tok/s	28.64 tok/s
3	exact	986.97 tok/s	28.64 tok/s
4	Google	977.83 tok/s	27.68 tok/s

The rebuild generates 1.035x faster, the same in both passes of each file, and reads prompts at the same speed. Every generated token reads the output layer in full, and the rebuild's output layer is 247.5 MiB smaller. The gain is smaller than E4B's 1.06x and E2B's 1.11x because the output layer is a smaller share of what a 12B reads per token (inferred, not measured).

Generation slows as the conversation grows:

tokens already in context	0	8,192	16,384
prompt, 512 tokens	981.09 tok/s	801.58 tok/s	692.32 tok/s
generation	28.63 tok/s	26.81 tok/s	26.33 tok/s

Step 6 — Check How Close It Stays to bf16

run_kld.sh scores the same text with Google's bf16 weights and with each 4-bit file, and compares the predicted next-token distributions (llama-perplexity --kl-divergence, wikitext-2, 16 chunks of 512 tokens, on the CPU):

against bf16	Google's GGUF	exact rebuild
Mean KL divergence	0.543	0.125
99th percentile KL divergence	6.217	1.915
Same top token	71.77%	89.14%
Perplexity, relative to bf16	1.386	1.257

The rebuild is 4.4x closer to bf16 by mean divergence, and picks the same top token 89% of the time against 72%.

Compare the two columns with each other only. The 12B instruction-tuned model reads raw Wikipedia text poorly, at a bf16 perplexity of 312 against E4B's 40 on the same test, and every divergence is larger in that regime. E4B's rebuild scores 0.0010 on this test.

Step 7 — Rebuild the 12B Drafter

Google's drafter for this model is google/gemma-4-12B-it-qat-q4_0-unquantized-assistant: 4 layers of width 1,024 with queries only, reading the 12B's own KV cache, and a tied 262,144-token embedding. It ships as bf16 on a 4-bit grid.

llama.cpp's converter writes a Q8_0 file, which the drafter version of gguf_exact.py uses for metadata:

python3 convert_hf_to_gguf.py hf --outtype q8_0 --outfile gemma-4-12B-it-qat-assistant-q8_0.gguf
python3 gguf_exact.py hf gemma-4-12B-it-qat-assistant-q8_0.gguf gemma-4-12B-it-qat-assistant-q4_0-exact.gguf

Every block landed on the grid. The rebuild is 253,708,960 bytes against 465,128,096 for Q8_0, and its weights take 226.90 MiB on the card against 428.53 MiB. This drafter is 4.7x the size of the E2B and E4B drafters.

Step 8 — Measure the Drafter

The rig serves with four extra flags for the drafter, and a 4-bit context cache (Step 10):

llama-server -m gemma-4-12B-it-q4_0-exact.gguf -ngl 99 -c 81920 -ctk q4_0 -ctv q4_0 -fa 1 --reasoning off \
  -md gemma-4-12B-it-qat-assistant-q4_0-exact.gguf --spec-type draft-mtp -ngld all --spec-draft-n-max 2

spec_ab.py runs six prompts (code, advice, a story, an explanation, a list, a summary), 256 tokens each, three greedy repeats plus one sampled pass, one server per setting at -c 8192, the card cooled before each. Serving without a drafter ran first and last, at 28.1 tok/s both times.

greedy, geometric mean of 6 prompts	tok/s	gain	drafted tokens kept	memory
no drafter	28.1	—	—	7,255 MiB
Q8_0 drafter, 2 tokens	53.4	1.90x	0.661	7,771 MiB
exact Q4_0 drafter, 2 tokens	56.4	2.01x	0.658	7,569 MiB
exact Q4_0 drafter, 1 token	45.6	1.62x	0.759	7,569 MiB
exact Q4_0 drafter, 3 tokens	56.7	2.02x	0.561	7,569 MiB

Per prompt at 2 tokens: code 2.40x, a list of facts 2.30x, advice 1.94x, an explanation 1.93x, a summary 1.82x, a story 1.75x. Sampled at temperature 0.8: 1.90x.

The drafter doubles the 12B's generation speed. The gain is larger than E4B's 1.72x on this card in Part 2: the 12B keeps 66% of drafted tokens against E4B's 48%.

🔎 Tip: For the 12B, the 4-Bit Drafter Is Faster Than the 8-Bit One

For E2B and E4B the Q4_0 and Q8_0 drafters ran level. For the 12B the Q4_0 drafter is 1.06x faster (1.04–1.07x on each prompt), at the same share of tokens kept, and uses 202 MiB less memory. This drafter is 4.7x larger, so its own weight reads are likely a larger share of each drafted step (inferred, not measured). On an 8 GB card the 202 MiB is also about 12,900 tokens of context (arithmetic).

🔎 Tip: Greedy Text Changes With the Drafter

With the drafter on, greedy replies repeat exactly from run to run, but differ from the replies without a drafter in all 18 at 2 drafted tokens and 9 of 18 at 1. Checking several tokens in one pass runs different GPU kernels, which can change a near-tie word. Step 9 checks whether the answers change.

Step 9 — Check the Answers

The rig's own settings, one request at a time: GSM8K (1,319 grade-school math questions, up to 2,048 tokens per reply) and BFCL simple (400 single function calls), with the drafter and without. E4B's scores are from Part 2's run of the same tests on the same card.

one request at a time	12B, drafter	12B, no drafter	E4B, drafter	E4B, no drafter
GSM8K	0.960	0.963	0.930	0.931
BFCL simple	0.950	0.948	0.913	0.913

The 12B answers more questions than E4B on both tests. Paired question by question, it is 3.0 points ahead on GSM8K (95% range 1.8 to 4.3) and 3.8 points ahead on BFCL (1.0 to 6.5), both with the drafter.

The drafter changes no score. On GSM8K the runs with and without it differ by 0.3 points (95% range −0.8 to +0.2), on BFCL by one question of 400. The greedy text differs, as in Step 8: 332 of 1,319 GSM8K replies match word for word. Without the drafter GSM8K took 2.34x as long, 4 h 16 min against 1 h 49 min.

Step 10 — Choose the Context Cache

The rig serves one long conversation at a time, so the context it holds and its speed with that context matter most. kv_sweep.py loads each cache type at its largest context, with and without the drafter, puts 0 to 98,304 tokens of a long document ahead of three questions, and measures generation (256 greedy tokens, geometric mean of the three):

tokens already in context	0	16,384	32,768	65,536	98,304
16-bit cache, drafter (24,576)	57.5 tok/s	49.8	—	—	—
8-bit cache, drafter (49,152)	55.1	43.3	39.3	—	—
4-bit cache, drafter (81,920)	55.1	42.4	37.8	30.8	—
4-bit cache, no drafter (131,072)	27.5	24.1	22.2	19.2	16.8

The drafter keeps paying with a long conversation held: 2.00x at 0 tokens, 1.70x at 32,768, 1.60x at 65,536. A smaller cache costs the drafter more than it costs plain generation: at 16,384 tokens the 4-bit cache runs at 0.85x the 16-bit one with the drafter and 0.94x without. Prompt reading is level across cache types, about 720 tok/s at 16,384 tokens and 408 at 65,536.

The same GSM8K and BFCL runs as Step 9, at each cache's largest context, paired question by question against the 16-bit cache:

one request at a time	GSM8K	BFCL simple
16-bit cache, drafter	0.960	0.950
8-bit cache, drafter	0.960	0.948
4-bit cache, drafter	0.961	0.950
4-bit cache, no drafter	0.962	0.950

The cache type changes no score: every paired range includes zero. A recall test plants eight six-digit codes through the document and asks for all of them: the 16-bit, 8-bit and 4-bit caches return the same share, about 85% at 16,384 tokens and 70% at 32,768, and the drafter changes neither.

3 drafted tokens ran level with 2 at every depth and 1 ran well behind, so 2 stays. Larger prompt batches (-ub 1024) do not load beside the drafter.

Compare and Contrast

RTX 3050 OEM, 8 GB	🥇 Gemma 4 12B	🥈 Gemma 4 E4B
File	6.72 GB	4.22 GB
Generation, no drafter	28.6 tok/s	61.2 tok/s
Generation, 2-token drafter	56.4 tok/s	100.0 tok/s (Part 2)
Drafter gain	2.01x	1.72x
Context with the drafter, 16-bit cache	24,576 tokens	131,072 tokens
Context with the drafter, 4-bit cache	81,920 tokens	not measured
GSM8K, with the drafter	0.960	0.930
BFCL simple, with the drafter	0.950	0.913

E4B's numbers are from llama.cpp 77c4540ce (Part 2) except the no-drafter generation, measured beside the 12B on 7182ef71e.

So, Which One?

For answers, the 12B. It scores 3 to 4 points higher than E4B on math and on function calls, and with the drafter it generates at 56.4 tok/s, near E4B's 61.2 without one.

For one long conversation, the 12B with a 4-bit cache and the drafter: 81,920 tokens at 55 tok/s with nothing held and 31 tok/s at 65,536, with the same answers as the 16-bit cache. Past 81,920 tokens, drop the drafter for the full 131,072 at about half the speed.

For speed, E4B. It generates at 100 tok/s with its drafter and holds the full 131,072-token context with a 16-bit cache.

Either way, serve with the drafter: it doubles the 12B's speed and costs no answers.

What Comes Next

Google publishes the same QAT drafters for 26B-A4B and 31B, both too large for this card. The next measurements are the 12B under several requests at once, which share the 81,920 tokens, and the 12B on a 16 GB card with its full context and the drafter.

Summary

The goal of this article was to find out whether Gemma 4 12B, in Google's QAT build, runs on an 8 GB RTX 3050 with every layer on the GPU. It does. The key to the solution was rebuilding Google's GGUF with every weight on the 4-bit grid QAT trained, which frees the memory a working context needs, and rebuilding the 12B drafter the same way. The results were:

🟢 The exact Q4_0 rebuild fits with every layer on the card at up to 40,960 tokens of context, against 28,672 for Google's file
🟢 It generates 1.035x faster than Google's file and stays 4.4x closer to bf16 by mean divergence
🟢 Google's 12B drafter, rebuilt as exact Q4_0, doubles generation to 56.4 tok/s, and runs 1.06x faster than the Q8_0 conversion
🟢 The 12B answers more questions than E4B on the same card: GSM8K 0.960 against 0.930, BFCL 0.950 against 0.913, and the drafter changes neither score
🟢 A 4-bit context cache holds 81,920 tokens with the drafter, 3.3x the 16-bit cache, with no change in GSM8K, BFCL or long-document recall
⚠️ With the drafter, the 4-bit cache generates at 0.85x the 16-bit cache once 16,384 tokens are held
⚠️ The 12B generates at 0.47x E4B's speed on the same card
⚠️ The full 131,072-token context fits only without the drafter, at about half the speed

Scope: Gemma 4 12B-it QAT only, text only, one RTX 3050 OEM (8 GB) in an i7-5820K desktop, llama.cpp 7182ef71e, Google's 12B Q4_0 GGUF against the exact Q4_0 rebuild, Google's 12B QAT drafter as exact Q4_0 and Q8_0, six short prompts for speed, divergence on 8,192 tokens of wikitext-2 on the CPU, GSM8K and BFCL simple one request at a time, 16-bit, 8-bit and 4-bit context caches with 0 to 98,304 tokens of one document in context, one request at a time throughout. E4B comparisons from Part 2 used llama.cpp 77c4540ce.

The strategy for using MCP for local accelerator comparison was validated with an incremental step by step approach.

References

Part 1, Gemma 4's MTP Drafter, Explained: https://dev.to/gde/gemma-4s-mtp-drafter-explained-speculative-decoding-and-a-4-bit-drafter-at-60-of-the-size-1oc0
Part 2, Do Drafters Work on Small Models?: https://dev.to/gde/do-drafters-work-on-small-models-gemma-4-e2b-and-e4b-on-two-small-gpus-1b4g
Repository: https://github.com/xbill9/gemma4-dev
12B rig: https://github.com/xbill9/gemma4-dev/tree/main/local-llamacpp-3050oem-12b-q4_0
The rebuild and divergence: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-12b-q4_0/benchmarks/runs/2026-10-16-exact-gguf-12b-3050oem/REPORT.md
Memory and speed: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-12b-q4_0/benchmarks/runs/2026-10-16-first-light-3050oem/REPORT.md
The drafter: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-12b-q4_0/benchmarks/runs/2026-10-16-drafter-q4exact-12b-3050oem/REPORT.md
Answers: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-12b-q4_0/benchmarks/runs/2026-10-16-accuracy-12b-3050oem/REPORT.md
Context cache: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-12b-q4_0/benchmarks/runs/2026-10-17-kvcache-12b-3050oem/REPORT.md
12B GGUF: https://huggingface.co/xbill9/gemma-4-12B-it-qat-q4_0-exact-gguf
12B drafter: https://huggingface.co/xbill9/gemma-4-12B-it-qat-assistant-q4_0-exact-gguf
Google's 12B QAT GGUF: https://huggingface.co/google/gemma-4-12B-it-qat-q4_0-gguf
llama.cpp: https://github.com/ggml-org/llama.cpp

Do Drafters Work on Small Models? Gemma 4 E2B and E4B on Two Small GPUs

xbill — Fri, 16 Oct 2026 17:30:10 +0000

This article provides a step by step guide to serving Gemma 4 E2B and E4B with Google's QAT assistant drafters through llama.cpp's speculative decoding, and measuring what the drafter is worth on two small NVIDIA cards: a desktop RTX 3050 OEM and a laptop GTX 1650 Ti. A suite of Python MCP tools is built to simplify management of the llama.cpp hosted deployment.

This is the second part of a short series. Part 1, Gemma 4's MTP Drafter, Explained, covers how speculative decoding and Gemma 4's multi-token prediction (MTP) drafter work and how the exact Q4_0 repack is built. This part has the measurements.

https://github.com/xbill9/gemma4-dev

A drafter is the guesser in speculative decoding: it proposes the next few tokens, and the main model checks them all in one pass.

Do drafters work on small models? Yes. Gemma 4 E2B and E4B are already small, and their drafter is a 4-layer model of 59 MB. Served one request at a time, it makes both models faster on both cards, and GSM8K and BFCL scores stay the same:

one request at a time	RTX 3050	GTX 1650 Ti
E4B	1.72x	1.49x
E2B	1.40x	1.18x
Answer quality	unchanged	unchanged

The smaller model gains less. E2B is already fast per token, so each drafted step costs a larger share of what it saves (inferred from the draft-length results below, not measured directly). Both cards keep the same share of the drafted tokens, within 0.03, so the gap between the cards comes from the cards.

The gain has limits. With several requests at once it shrinks on both cards, and on the 1650 Ti it turns into a loss from 4 requests (E4B) or 8 (E2B). On the 1650 Ti, E2B's gain is gone by about 32k tokens of conversation. Too many drafted tokens also cost speed on the slower card.

The drafter here is an exact Q4_0 repack of Google's weights. It matches the converter's Q8_0 file on speed and answers at 60% of its size, and beats Google's bf16 weights on speed on every card and model.

What Is Being Measured?

A language model writes one token per forward pass, and every pass reads every weight. For one user on a small GPU, each pass spends most of its time moving weights, which leaves compute to spare. Speculative decoding uses it: a small drafter guesses the next few tokens, and the main model checks all of them in one pass. It keeps the guesses up to the first one it disagrees with and adds one token of its own, so the output is the main model's either way. When most guesses are kept, one pass yields several tokens instead of one.

Gemma 4's drafter uses multi-token prediction (MTP). It is a 4-layer model that reads the main model's hidden state and reuses its KV cache, so it keeps no cache of its own and costs little more than its weights.

MTP is one of several ways to do speculative decoding. llama.cpp 77c4540ce offers ten through --spec-type: a separate small draft model (draft-simple), draft-eagle3, draft-mtp, draft-dflash, draft-dspark, and five n-gram methods. MTP is the one Google builds drafters for: E2B, E4B, 12B, 26B-A4B and 31B each have a published assistant checkpoint (E2B's config names it Gemma4AssistantForCausalLM), which llama.cpp runs as draft-mtp and vLLM as "method":"mtp". This article measures MTP only. Part 1 explains this in full.

Google publishes an assistant drafter for each Gemma 4 QAT checkpoint. Both cards here serve the exact Q4_0 GGUFs of gemma-4-E2B-it and gemma-4-E4B-it with the drafters stored back as Q4_0 on their trained 4-bit grid, 59 MB each. The question is what the drafter is worth on each card, and why the two cards differ.

At This Point You Should Have…

An NVIDIA GPU with a CUDA driver — here an RTX 3050 OEM (driver 615.78.08) and a GTX 1650 Ti with Max-Q Design (driver 615.71.09)
The CUDA toolkit, 13.4 here, and cmake
llama.cpp at a commit with --spec-type draft-mtp — both cards ran 77c4540ce for every comparison below
Python 3.10 or newer, the system python3, no virtualenv
The repository cloned, with the four rigs local-llamacpp-{3050oem,1650ti}-{2b,4b}-q4_0/
Claude Code, or any MCP client that speaks stdio

The Two Cards

	RTX 3050 OEM	GTX 1650 Ti Max-Q
Host	i7-5820K desktop, 46 GiB	i7-10750H laptop (Lenovo Yoga 9), 15 GiB
Memory	8192 MiB	4096 MiB
Architecture	Ampere, compute capability 8.6, tensor cores	Turing TU117, compute capability 7.5, no tensor cores
Power limit	120 W	about 40 W
E2B, no drafter	108.2 tok/s	76.5 tok/s
E4B, no drafter	58.3 tok/s	38.0 tok/s

Without a drafter the 3050 generates 1.41x faster on E2B and 1.53x on E4B. Different hosts, so treat ratios between the cards as indicative.

Step 1 — Repack the Drafter on Its Trained Grid

Google publishes the QAT assistant drafters as google/gemma-4-E2B-it-qat-q4_0-unquantized-assistant and google/gemma-4-E4B-it-qat-q4_0-unquantized-assistant. Each is a 4-layer model that reads the main model's KV cache, stored as bf16 safetensors.

Quantization-aware training already put the weights on a 4-bit grid: in every group of 32 values there are at most 16 distinct levels, evenly spaced. bf16 spends 16 bits storing each of those levels.

llama.cpp's convert_hf_to_gguf.py has no way to know that, so it writes Q8_0: 8.5 bits per weight, re-rounded onto an 8-bit grid of its own. The file shrinks, and the values move. Only 11.7% of the drafter's weights come back bit-identical.

Q4_0 is 4.5 bits per weight: one fp16 step per group of 32, and a 4-bit level for each value. That is enough to hold the trained grid, provided each group's step is read off its own values. gguf_exact.py does this for each group:

amax = np.abs(g).max(-1, keepdims=True)
for m in range(1, 9):                      # try steps amax/1 .. amax/8
    cand = amax / m
    q = np.clip(np.rint(g / cand), -8, 7)
    good = np.abs(g - q * cand).max(-1, keepdims=True) <= amax * 2**-7

The first step that puts every value in the group on an integer level wins. A least-squares fit then refines it, and it is rounded to fp16, Q4_0's step type. A group that fits no step stops the build, since nothing is re-rounded. On both drafters no group was off the grid.

The script takes Google's safetensors and the converter's Q8_0 file, copies the Q8_0 file's metadata byte for byte (tokenizer, chat template, settings), and writes every weight matrix as Q4_0. That covers 23 tensors per drafter, including the token embedding, which holds 86–87% of the values rewritten:

python3 gguf_exact.py <google-snapshot-dir> gemma-4-E2B-it-qat-assistant-q8_0.gguf gemma-4-E2B-it-qat-assistant-q4_0-exact.gguf
python3 gguf_exact.py <google-snapshot-dir> gemma-4-E2B-it-qat-assistant-q8_0.gguf gemma-4-E2B-it-qat-assistant-bf16.gguf --bf16

--bf16 writes Google's weights unchanged into the same layout, so all three formats can be compared with identical metadata.

E2B drafter	Q8_0 (converter)	exact Q4_0 (repack)	bf16 (Google's weights)
File size	97.8 MB	59.2 MB	170.2 MB
Bits per weight	8.5	4.5	16
Error against Google's weights	0.00494	0.00138	0
Values bit-identical, read back in fp32	11.7%	29.4%	100%
Values bit-identical, rounded to bf16	—	96.5–97.0% per tensor	100%

The repack is 60% of the Q8_0 file and its error is 3.6x lower. Rounded back to bf16, the format Google stored, 97% of its values match Google's bit for bit. The E4B drafter gives the same picture: 59.7 MB, error 0.00138, 96.4–97.1% per tensor.

🔎 Tip: The Converter Drops the Masked Embedding

llama.cpp's converter leaves out the checkpoint's masked_embedding tensors, and the repack keeps its tensor list. All three drafter files compared here lack them, so the comparison is fair. The files work with --spec-type draft-mtp as they are.

Step 2 — Fetch the Published Drafters

Both repacks are on Hugging Face:

xbill9/gemma-4-E2B-it-qat-assistant-q4_0-exact-gguf, file gemma-4-E2B-it-qat-assistant-q4_0-exact.gguf, SHA-256 060f5a0158f8120d…
xbill9/gemma-4-E4B-it-qat-assistant-q4_0-exact-gguf, file gemma-4-E4B-it-qat-assistant-q4_0-exact.gguf, SHA-256 dc965cb963e56b86…

They pair with the main models' exact Q4_0 GGUFs, built the same way: xbill9/gemma-4-E2B-it-qat-q4_0-exact-gguf and xbill9/gemma-4-E4B-it-qat-q4_0-exact-gguf.

Each rig's make draft downloads its drafter at a pinned revision and checks the hash from tpu.env:

curl -fL -o $(SPEC_DRAFT_MODEL).part https://huggingface.co/$(SPEC_DRAFT_HF)/resolve/$(SPEC_DRAFT_HF_REVISION)/$(notdir $(SPEC_DRAFT_MODEL))
echo "$(SPEC_DRAFT_SHA256)  $(SPEC_DRAFT_MODEL)" | sha256sum -c -

Both cards loaded the same two files, checked by hash.

How the Repack Compares on Both Cards

Each format was served on both cards at 2 drafted tokens, six prompts, greedy. Gain against no drafter:

2 drafted tokens	3050 E2B	1650 Ti E2B	3050 E4B	1650 Ti E4B
Q8_0	1.39x	1.09x	1.72x	1.40x
exact Q4_0	1.40x	1.11x	1.72x	1.42x
bf16	1.23x	0.96x	1.59x	1.30x

Memory with each drafter loaded, at -c 8192:

MiB	no drafter	Q8_0	exact Q4_0	bf16
3050 E2B	2,795	2,915	2,875	2,983
3050 E4B	4,415	4,537	4,499	4,607
1650 Ti E2B	1,485	1,595	1,557	1,663
1650 Ti E4B	2,853	2,965	2,927	3,035

The repack is level with Q8_0 on speed on both cards and uses 38–40 MiB less memory.
bf16 is slowest everywhere. On the 1650 Ti it makes E2B slower than serving without a drafter.
The share of drafted tokens kept is the same for all three formats, within 0.012 on both cards. The drafter's precision changes what each drafted step costs, never which tokens get kept.
Greedy text depends on drafting, never on the drafter's format. At a given draft length, Q8_0, the repack and bf16 give the same 18 greedy replies, on the 3050 for both models and on the 1650 Ti for E2B (the 1650 Ti's E4B text does not repeat on f95b0d9, so it cannot be compared there).

Answer quality on the 3050, GSM8K (1,319 problems) and BFCL simple (400 function calls), 8 requests at a time:

	no drafter	Q8_0	exact Q4_0
E2B GSM8K	0.8984	0.8992	0.9022
E2B BFCL	0.9250	0.9225	0.9225
E4B GSM8K	0.9340	0.9310	0.9325
E4B BFCL	0.9100	0.9150	0.9150

Every drafted score is level with the same day's no-drafter run, and the repack is level with Q8_0: E2B GSM8K +0.0030 (95% range −0.0061 to +0.0129), identical BFCL answers on E2B.

Google also publishes non-QAT drafters, google/gemma-4-E2B-it-assistant and google/gemma-4-E4B-it-assistant. Their weights are on no 4-bit grid, so they were converted to Q8_0. Against the QAT targets they keep fewer drafted tokens on E2B and are level or behind at every draft length on both models: E2B 1.34x against the repack's 1.40x on the 3050, and 1.16x against 1.18x on the 1650 Ti.

The 1650 Ti's format run used llama.cpp f95b0d9, the build its rigs serve; every other comparison here uses 77c4540ce on both cards.

Is the Repack Worth It?

Yes, as the drafter file to ship. Against the two files a user would otherwise serve:

exact Q4_0 repack against	Q8_0 (converter)	bf16 (Google's weights)
Speed	level on both cards	faster on every card and model
Share of drafted tokens kept	level, within 0.012	level, within 0.012
GSM8K and BFCL	level	not scored
File size	60%	35%
Card memory	38–40 MiB less	106–108 MiB less
Match to Google's weights	3.6x lower error	97% bit-identical in bf16

🟢 Against bf16 it is the clear choice. It is faster everywhere, and on the 1650 Ti bf16 makes E2B slower than serving without a drafter.
🟢 Against Q8_0 it is a free improvement: smaller, closer to the trained weights, and level on speed and answers.
⚠️ The memory it saves over Q8_0 is under 1% of either card. It counts only on a 4 GB card running near its limit.
⚠️ Speed stays level with Q8_0 because the drafter's precision leaves the kept tokens unchanged. Every format keeps the same share, so a more exact drafter guesses as well as a rounded one and gains its speed from size alone.

Step 3 — Serve With the Drafter

The drafter adds four flags to the rig's normal command line:

llama-server -m gemma-4-E4B-it-q4_0-exact.gguf --host 127.0.0.1 --port 8080 \
  -ngl 99 -c 8192 -ctk f16 -ctv f16 -fa 1 --reasoning off --metrics \
  -md gemma-4-E4B-it-qat-assistant-q4_0-exact.gguf --spec-type draft-mtp -ngld all --spec-draft-n-max 2

--spec-draft-n-max is the number of tokens drafted per step. Each rig's make serve builds this line from tpu.env, and an empty SPEC_DRAFT_MODEL serves without the drafter. The drafter costs 72–84 MiB of memory on either card.

Step 4 — Measure One Request at a Time

speed.py runs six prompts (code, advice, a story, an explanation, a list, a summary), 256 tokens each, three greedy repeats plus one sampled pass at temperature 0.8, with the prompt cache off. Each draft length gets its own server, the card cools before each one, and the first drafted setting runs again at the end to check for drift. The same script, prompts and llama.cpp build ran on both cards.

python3 speed.py <rig> <out>

The Result

Greedy, geometric mean over the six prompts. "Kept" is the share of drafted tokens the main model accepted.

	3050 tok/s	3050 gain	3050 kept	1650 Ti tok/s	1650 Ti gain	1650 Ti kept
E2B, no drafter	108.2	—	—	76.5	—	—
E2B, 1 token	148.9	1.38x	0.609	90.3	1.18x	0.602
E2B, 2 tokens	151.2	1.40x	0.475	83.6	1.09x	0.469
E2B, 3 tokens	143.0	1.32x	0.384	77.3	1.01x	0.378
E4B, no drafter	58.3	—	—	38.0	—	—
E4B, 1 token	89.8	1.54x	0.624	56.5	1.49x	0.635
E4B, 2 tokens	100.0	1.72x	0.481	54.5	1.44x	0.506
E4B, 3 tokens	99.3	1.70x	0.398	50.9	1.34x	0.414

The share of tokens kept matches across the cards at every length. What differs is how much each extra drafted token costs: on the 3050 E4B gains from 1 token to 2 and barely loses at 3, while on the 1650 Ti every extra token makes it slower. The likely reason is that each drafted step costs a larger share of a main-model step on the older card (inferred, not measured).

The repeated setting landed within 0.3% (3050) and 2.0% (1650 Ti) of its first run.

🔎 Tip: Pick the Draft Length Per Card

The best draft length is a property of the card as much as the model. The 3050 rigs serve 2 tokens for both models. On the 1650 Ti, 1 token leads for both models on this build, and E2B at 3 tokens is no faster than serving without a drafter. A setting copied from a faster card loses most of the gain on a slower one.

🔎 Tip: The Same Pattern on a Tesla T4

A Tesla T4 (Turing, 15,360 MiB) on Google Cloud shows the same answer under a different engine. With vLLM 0.29, the W4A16 builds of E2B and E4B and Google's bf16 drafters, the same six prompts one request at a time speed up E2B 1.29x and E4B 1.57x at 3 drafted tokens, and 1.31x and 1.62x at 4. There the gain keeps rising to 4 drafted tokens, against 1 on the 1650 Ti, another case of the draft length belonging to the card. vLLM cannot load the GGUF repack, and engine, target files and drafter format all differ, so the T4 numbers stay out of the tables here (T4 drafter run).

Step 5 — Measure With Long Text Ahead

longctx.py puts the first 0 to 120,000 tokens of the NVIDIA driver README ahead of a story, a code request or a summary of the text, and measures generation speed alone. Same document, by hash, on both cards. The 1650 Ti's E4B stops at 32,768 tokens, the most that fits in 4 GiB with a drafter loaded.

Gain at 1 drafted token, mean over the three requests:

tokens ahead	3050 E2B	1650 Ti E2B	3050 E4B	1650 Ti E4B
0	1.41x	1.27x	1.61x	1.52x
8,192	1.38x	1.18x	1.55x	1.53x
32,768	1.33x	1.03x	1.48x	1.45x
65,536	1.27x	0.97x	1.46x
120,000	1.23x	0.89x	1.41x

On the 3050 the drafter pays at every depth for both models; at 2 tokens it does better still, E4B 1.86x with nothing ahead and 1.55x at 120k. On the 1650 Ti E2B's gain is gone by 32k tokens and turns into a loss beyond. The share of tokens kept holds steady as the text grows on both cards, so the drop comes from each checking step getting dearer with a long context.

Step 6 — Measure Several Requests at Once

sweep.py sends 1 to 16 requests at once, each with a unique prompt, in two shapes: 512 tokens in and 128 out, and 128 in and 512 out. Both cards ran it with one shared KV cache (-kvu); the 3050 at -c 65536 --parallel 64, the 1650 Ti at -c 16384 --parallel 16.

Total output with the drafter divided by without:

requests at once	1	2	4	8	16
3050 E2B, 512 / 128	1.56	1.17	1.28	1.28	1.05
1650 Ti E2B, 512 / 128	1.15	1.04	1.03	0.95	0.97
3050 E2B, 128 / 512	1.79	1.28	1.43	1.54	1.18
1650 Ti E2B, 128 / 512	1.33	1.12	1.09	0.91	0.96
3050 E4B, 512 / 128	1.86	1.38	1.32	1.38	1.14
1650 Ti E4B, 512 / 128	1.28	1.07	0.90	0.93	0.92
3050 E4B, 128 / 512	2.05	1.44	1.41	1.60	1.26
1650 Ti E4B, 128 / 512	1.65	1.18	0.79	0.83	0.83

sweep.py's prompt repeats one paragraph, which the drafter guesses unusually well, so these ratios run higher than on everyday chat. On mixed chat prompts the 3050 measured E4B 1.08–1.24x and E2B 0.92–1.17x at 2–8 requests, and a loss at 32–64 for both. Read the table for the trend: the 3050 keeps a gain through 16 requests here, and the 1650 Ti loses from 4 or 8.

With long prompts the drafter pays further. longmixed.py puts a different 1,000–3,500-token stretch of the driver README ahead of each of six tasks (a summary, a bash script, a quiz and so on), 256 tokens out, greedy, at 1 to 32 requests on the 3050 (-c 131072 --parallel 32), in two passes run in opposite order. The drafter keeps 0.56–0.62 of its tokens with a document in context, against 0.43–0.50 on short chat. Total output with the drafter divided by without, both passes:

requests at once	1	2–16	32
3050 E2B, per-slot cache	1.34	1.02–1.14	1.02–1.04
3050 E2B, shared cache	1.26	1.04–1.13	0.96–0.97
3050 E4B, per-slot cache	1.54	1.08–1.16	1.03–1.05
3050 E4B, shared cache	1.61	1.08–1.18	0.99

The drafter pays at every level up to 16 requests and breaks even at 32. This run covers the 3050 only.

🔎 Tip: The Shared KV Cache Has a Price on Both Cards

llama-server gives each slot its own KV cache by default. On llama.cpp 77c4540ce that layout halved generation speed at random with several requests in flight, on both cards: per-slot repeats spread 28–36% at 2 requests on the 1650 Ti. -kvu gives all slots one shared cache and the spread drops to about 1%. It also costs output as the total text held grows. At 16 requests in sweep.py the shared cache delivered 0.86–0.94x of per-slot output on the 3050 and 0.59–0.81x on the 1650 Ti. The drafter's ratio is nearly the same under either layout. With 1,000–3,500-token prompts on the 3050 the per-slot cache is the faster layout and showed no random slowdown: it beat the shared cache in 204 of 232 paired batches, and the shared cache gave 0.68–0.82x of its output at 16 requests and 0.49–0.66x at 32. All four rigs serve one request at a time, where the two layouts are identical.

Step 7 — Check the Answers

GSM8K (1,319 problems) and BFCL simple (400 function calls), one request at a time, with and without the drafter:

	3050 none	3050 drafter	1650 Ti none	1650 Ti drafter
E2B GSM8K	0.9014	0.8969	0.8939	0.8969
E2B BFCL	0.9225	0.9250	0.9250	0.9225
E4B GSM8K	0.9310	0.9295	0.9363	0.9310
E4B BFCL	0.9125	0.9125	0.9125	0.9125

Every difference has a 95% range that includes zero. The drafter finishes GSM8K 1.76x (E2B) and 2.23x (E4B) faster on the 3050 and 1.34x and 1.78x faster on the 1650 Ti. The 3050 ran at -c 131072 with 2 drafted tokens for both models; the 1650 Ti at -c 8192 with each rig's setting, 1 token for E2B and 2 for E4B.

🔎 Tip: Greedy Text Can Change With the Drafter

At temperature 0 a drafted reply can differ from the undrafted one at a near-tie word, because checking several tokens at once runs different GPU kernels. On both cards E2B's drafted text differs from its undrafted text on every prompt at 2 tokens. E4B at 1 token matches on all 18 short replies on both cards. On the 3050 E4B also matches at 2 tokens for short prompts and starts to differ from about 2,000 tokens of prompt. The prompt cache changes greedy text by itself as well. A test that needs identical text must serve without the drafter and send cache_prompt: false. The scores above show that the changed words do not change the answers.

Compare and Contrast

	🥇 RTX 3050 OEM	🥈 GTX 1650 Ti Max-Q
Best gain, E2B, one request	1.40x at 2 tokens	1.18x at 1 token
Best gain, E4B, one request	1.72x at 2 tokens	1.49x at 1 token
Fastest drafted E4B	100.0 tok/s	56.5 tok/s
Share of tokens kept, 1 token	0.61–0.62	0.60–0.64
E2B gain with 120k tokens ahead	1.23x	0.89x
Gain at 8 requests at once (`sweep.py`)	1.28–1.60x	0.83–0.95x
Answer quality	unchanged	unchanged
GSM8K faster with the drafter	1.76x / 2.23x	1.34x / 1.78x

So, Which One?

Drafters work on small models. Turn the drafter on for both cards when they serve one user. It costs no answers and 72–84 MiB of memory, and E4B gains the most on both: 1.72x on the 3050 and 1.49x on the 1650 Ti. Set the draft length per card, 2 tokens on the 3050 and 1 on the 1650 Ti.

On the 1650 Ti, leave E2B's drafter off for long conversations past about 32k tokens and for several users at once. On the 3050, the drafter keeps paying with long text ahead, and with long prompts it pays up to 16 users at once and breaks even at 32.

What Comes Next

Google publishes the same two drafters, QAT and non-QAT, for the larger Gemma 4 models: 12B, 26B-A4B and 31B. None of them has been measured here yet, and whether they fit beside a drafter on these two cards is unchecked. The next measurements are the larger models with their drafters, and the NVIDIA L4 (Ada, compute capability 8.9, 24 GB), a generation newer than the 3050.

Summary

The goal of this article was to find out whether Google's Gemma 4 assistant drafters speed up small models on small NVIDIA cards served by llama.cpp. They do. The key to the solution was running the same llama.cpp build, scripts, prompts and drafter files on both cards, one setting per server, and checking answer quality alongside speed. The results were:

🟢 One request at a time the drafter speeds up E4B 1.72x on the RTX 3050 and 1.49x on the GTX 1650 Ti, and E2B 1.40x and 1.18x
🟢 Both cards keep the same share of drafted tokens, within 0.03 at every draft length
🟢 GSM8K and BFCL scores are unchanged with the drafter on both cards
🟢 The exact Q4_0 repack of the drafter is 60% of the Q8_0 file and level with it on speed on both cards; Google's bf16 weights are slowest on both
⚠️ The best draft length differs: 2 tokens on the 3050, 1 on the 1650 Ti
⚠️ With long text ahead, E2B's gain on the 1650 Ti is gone by 32k tokens; on the 3050 it holds at 1.23x at 120k
🟢 With 1,000–3,500-token prompts on the 3050 the drafter pays at every level up to 16 requests at once
❌ With 4–8 or more requests at once the drafter costs output on the 1650 Ti

Scope: Gemma 4 E2B and E4B only, MTP speculative decoding only, two cards on two hosts, llama.cpp 77c4540ce on both, the exact Q4_0 E2B and E4B GGUFs with Google's QAT assistant drafters as exact Q4_0 (plus Q8_0 and bf16 for the format comparison), six short prompts for speed, one 120k-token document for long text, sweep.py at 1–16 requests with a shared KV cache, six long-prompt tasks at 1–32 requests on the 3050 with both cache layouts, GSM8K and BFCL simple one request at a time, each setting run once with a repeat check. The T4 tip comes from a separate vLLM run.

The strategy for using MCP for local accelerator comparison was validated with an incremental step by step approach.

References

Part 1, Gemma 4's MTP Drafter, Explained: https://dev.to/gde/gemma-4s-mtp-drafter-explained-speculative-decoding-and-a-4-bit-drafter-at-60-of-the-size-1oc0
Repository: https://github.com/xbill9/gemma4-dev
RTX 3050 rigs: https://github.com/xbill9/gemma4-dev/tree/main/local-llamacpp-3050oem-2b-q4_0 and https://github.com/xbill9/gemma4-dev/tree/main/local-llamacpp-3050oem-4b-q4_0
GTX 1650 Ti rigs: https://github.com/xbill9/gemma4-dev/tree/main/local-llamacpp-1650ti-2b-q4_0 and https://github.com/xbill9/gemma4-dev/tree/main/local-llamacpp-1650ti-4b-q4_0
Drafter repack, 3050: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-2b-q4_0/benchmarks/runs/2026-10-14-drafter-q4exact-3050oem/REPORT.md
Drafter formats, 1650 Ti: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-1650ti-2b-q4_0/benchmarks/runs/2026-10-14-drafter-formats-1650ti/REPORT.md
3050 draft length and long text: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-2b-q4_0/benchmarks/runs/2026-10-14-drafter-gaps-3050oem/REPORT.md
1650 Ti draft length and long text: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-1650ti-2b-q4_0/benchmarks/runs/2026-10-14-drafter-gaps-1650ti/REPORT.md
3050 several requests at once: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-2b-q4_0/benchmarks/runs/2026-10-15-concurrency-drafter-kvu-3050oem/REPORT.md
3050 several requests at once, long prompts: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-2b-q4_0/benchmarks/runs/2026-10-15-concurrency-longmixed-3050oem/REPORT.md
1650 Ti several requests at once and answers: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-1650ti-2b-q4_0/benchmarks/runs/2026-10-15-drafter-load-1650ti/REPORT.md
Tesla T4 drafters under vLLM: https://github.com/xbill9/gemma4-dev/blob/main/gpu-vllm-t4-2b-w4a16/benchmarks/runs/2026-10-15-drafter-sweep-t4/REPORT.md
3050 answers: https://github.com/xbill9/gemma4-dev/blob/main/local-llamacpp-3050oem-2b-q4_0/benchmarks/runs/2026-10-14-accuracy-fullctx-3050oem/REPORT.md
Drafters: https://huggingface.co/xbill9/gemma-4-E2B-it-qat-assistant-q4_0-exact-gguf and https://huggingface.co/xbill9/gemma-4-E4B-it-qat-assistant-q4_0-exact-gguf
llama.cpp: https://github.com/ggml-org/llama.cpp

Renting Out One T4 to Three Customers: Gemma 4, vLLM Multi-LoRA and a Rust Gateway

xbill — Fri, 16 Oct 2026 16:48:42 +0000

This is a submission for the Hacktoberfest Open-Source AI Challenge Week 2: Side Hustle

This article provides a step by step guide to renting out one NVIDIA T4 on a Google Cloud VM to several small customers at once, with Gemma 4 E2B served by vLLM and a small Rust gateway in front. Every log, script and number is committed.

https://github.com/xbill9/t4-timeshare

One customer sending 4,000-token prompts pushes everyone else's first token from 0.28 s to 29.90 s on a shared T4. vLLM's priority scheduling brings that back to 6.52 s. The gateway adds what turns a shared GPU into a business: keys, a fine-tune per customer that no one else can call, a prefix cache per customer, and a ledger that matched vLLM's own token counters exactly. Billed by GPU time, the long-prompt customer pays 27.2% of the VM for 22.0% of the tokens.

What I Built

A small AI job costs too much either way. A GPU of your own sits idle most of the day. A per-token API bill grows with every customer you add, and every prompt leaves your building.

The side hustle is the middle: one person rents one T4, and sells shares of it to a few small businesses at a flat monthly price. Each customer gets:

an API key that works with any OpenAI client
its own fine-tuned adapter on the shared base model, callable by that customer only
a prefix cache that other customers cannot read through response timing
a bill computed from metered tokens, weighted by what each token costs the GPU

The operator gets one config file and a ledger.

The three demo tenants are picked so each one loads the T4 differently:

Tenant	Job	Shape
A	Summarise long daily logs	Few requests, 4,096-token prompts
B	Triage short support emails	One 512-token request every 2 s
C	Tag product rows, with its own adapter	One 512-token request every 2 s

Demo

The gateway listens on port 8080 and speaks the OpenAI chat API. day3/demo.sh, run with bash -v:

curl -s localhost:8080/v1/chat/completions -H "Authorization: Bearer demo-key-B" \
  -d '{"messages":[{"role":"user","content":"Label this email: my parcel arrived broken"}],"max_tokens":20}' \
  | jq -c '{model, usage: .usage.total_tokens}'
{"model":"xbill9/gemma-4-E2B-it-qat-q4_0-w4a16-ct-text-emb4","usage":37}
curl -s localhost:8080/v1/chat/completions -H "Authorization: Bearer demo-key-C" \
  -d '{"messages":[{"role":"user","content":"Label this email: my parcel arrived broken"}],"max_tokens":20}' \
  | jq -c '{model, usage: .usage.total_tokens}'
{"model":"rand","usage":26}
curl -s localhost:8080/v1/chat/completions -H "Authorization: Bearer demo-key-B" \
  -d '{"model":"rand","messages":[{"role":"user","content":"hi"}]}' -w ' %{http_code}\n'
{"error":"adapter belongs to another tenant"} 403
curl -s localhost:8080/v1/chat/completions -d '{"messages":[{"role":"user","content":"hi"}]}' -w ' %{http_code}\n'
{"error":"missing or unknown key"} 401
tail -4 ledger.jsonl
{"completion_tokens":20,"model":"xbill9/gemma-4-E2B-it-qat-q4_0-w4a16-ct-text-emb4","prompt_tokens":17,"queue_ms":0.0006799999999999999,"status":200,"tenant":"B","total_ms":355.14126200000004,"ts":1792165075.2638927}
{"completion_tokens":9,"model":"rand","prompt_tokens":17,"queue_ms":0.000778,"status":200,"tenant":"C","total_ms":498.805413,"ts":1792165075.7809048}
{"refused":"adapter belongs to another tenant","status":403,"tenant":"B","ts":1792165075.7979941}
{"refused":"missing or unknown key","status":401,"tenant":null,"ts":1792165075.8125255}

The same request from B and C reaches two different models: B gets the base model, C gets its adapter without naming it. B asking for C's adapter is refused, and so is a request with no key. Every one of the four lands in the ledger.

Code

https://github.com/xbill9/t4-timeshare

Path	What
`gateway/src/main.rs`	The gateway, 245 lines of Rust (axum, reqwest, tokio)
`gateway/tenants.toml`	Keys, adapters, priorities and in-flight caps
`day1/`	Checks on the T4 before building: LoRA, cache salt, streamed usage, priority, prefill cost
`day2/`	The noisy-neighbour test straight to vLLM
`day3/`	The same test through the gateway, isolation checks, the bill and VM pricing

How I Built It

At This Point You Should Have…

A Google Cloud VM with one NVIDIA T4. This one is an n1-standard-2 in us-west2-b
vLLM 0.29.0 serving xbill9/gemma-4-E2B-it-qat-q4_0-w4a16-ct-text-emb4, Gemma 4 E2B with Google's quantization-aware 4-bit weights, fp16, at most 8 sequences at once
Rust, for the gateway
A clone: git clone https://github.com/xbill9/t4-timeshare

Step 1 — Check the Engine Can Do Multi-Tenant Work

Four features decide whether one vLLM can serve several paying customers. Each was checked on the T4 before any gateway code was written (day1/RESULTS.md):

Check	Result
LoRA adapters on the 4-bit weights, Turing GPU	🟢 A zero adapter returns the base model's output, 40 of 40 tokens. A random adapter differs from token 1. 6 of 6 requests in one mixed batch match their solo output.
Per-request prefix-cache salt	🟢 Same salt, repeat: 1,216 of 1,217 tokens from cache. Different salt: 0.
Token usage on streamed responses	🟢 Present with `stream_options.include_usage`, absent without it. The gateway sets it.
Priority scheduling	🟢 With `--scheduling-policy priority`, a high-priority request queued last finished first.

The 6 of 6 is what makes per-customer fine-tunes cheap: requests for different adapters share one batch on one GPU, and each still gets its own model's answer.

Step 2 — Measure What Hurts on a T4

A customer's request costs the GPU in two parts: reading the prompt (prefill) and writing the answer (decode). day1/prefill_cost.py times prefill alone, one request at a time, 5 per length, with the prefix cache unable to help:

prompt ~  256: median    106.9 ms  min    103.4  max    108.7
prompt ~  512: median    262.0 ms  min    258.9  max    267.9
prompt ~ 1024: median    730.8 ms  min    717.0  max    737.4
prompt ~ 2048: median   2180.8 ms  min   2149.6  max   2223.6
prompt ~ 4096: median   7122.4 ms  min   7050.4  max   7363.7
prompt ~ 8192: median  25364.0 ms  min  25261.6  max  25697.8
prefill fit over 30 requests: -12.8 ms + 374.32 us x n + 336.9989 ns x n^2; worst error 2.1%
decode: median 9.34 ms per output token over 5 pairs (256 tokens minus 1 token, same length prompt)

Prefill grows with the square of the prompt. One 4,096-token prompt costs as much GPU time as 767.8 output tokens. On a T4 the customer who hurts the others is the one with long prompts, more than the one with many requests.

Step 3 — Watch One Customer Stall the Others

day2/m1_direct.py runs the three tenants straight against vLLM. B and C each send one 512-token request every 2 seconds and stream 64 tokens back. For 180 seconds, A keeps four 4,096-token summaries in flight. Every prompt starts with a random tag, so nothing comes from cache.

First come, first served, every request in the same queue:

[quiet] B: 90/90 ok  first token median    0.29 s  95th    0.29 s  max    0.34 s | total median    0.98 s  95th    0.99 s
[flood] A: 16/16 ok  first token median   30.82 s  95th   45.37 s  max   46.12 s | total median   50.08 s  95th   54.03 s
[flood] B: 90/90 ok  first token median   29.78 s  95th   44.47 s  max   47.81 s | total median   35.07 s  95th   48.91 s

A support-email triage that answered in a third of a second now waits half a minute for its first word. B pays for A's prompts.

Step 4 — Put the Gateway in Front

gateway/src/main.rs does six things per request, in order:

Key to tenant. Authorization: Bearer <key> looks the tenant up in tenants.toml. No key or an unknown one: 401.
Model check. The tenant gets its own adapter by default, or the base model on request. Another tenant's adapter: 403.
Priority. The tenant's priority goes into the request body, overwriting whatever the client sent.
Cache salt. cache_salt is set to the tenant's name, so cached prompts are shared within a tenant and never across tenants.
In-flight cap. The request waits for one of the tenant's slots, a tokio semaphore per tenant.
Ledger. When the response ends, or the client disconnects, one JSON line is written with the tenant, the model, the status, vLLM's own prompt and output token counts, the time spent waiting for a slot and the total time.

The demo tenants:

[[tenant]]
name = "A"
key = "demo-key-A"
max_inflight = 1
priority = 10

[[tenant]]
name = "C"
key = "demo-key-C"
adapter = "rand"
max_inflight = 8
priority = 0

A has one slot and runs last. B and C have 8 slots each, the most vLLM runs at once. C's rand adapter is a random rank-8 LoRA from day1/make_adapters.py, standing in for a customer's fine-tune.

Streaming responses pass straight through. The gateway watches the stream for the one chunk that carries usage and records it when the stream ends. A request that dies halfway is still in the ledger, and its slot is released at that moment.

Step 5 — Check Isolation

python3 day3/isolation.py

PASS  no key: 401
PASS  unknown key: 401
PASS  B asks for C's adapter: 403
PASS  B asks for a model that is not served: 404
PASS  B gets the base model: xbill9/gemma-4-E2B-it-qat-q4_0-w4a16-ct-text-emb4
PASS  C gets its own adapter by default: rand
PASS  C's answer differs from the base model's: 'Here are a few examples of how a spare G' vs 'Here are three common uses for a spare G'
PASS  B's repeat is served from B's cache: 2400 hit tokens (first send 0)
PASS  A sending B's prompt gets none of B's cache: 0 hit tokens
9 of 9 pass

The last two lines are the privacy check. When B sends a long prompt twice, the second send is served from cache. A sends the identical prompt and gets nothing from B's cache, so A cannot tell from response time whether B ever sent it. The hit counts come from vLLM's /metrics, not from timing.

Step 6 — Rerun the Noisy Neighbour Through the Gateway

day3/run_m1.sh
python3 day3/compare.py

run                                   B+C first token median    95th     max  B+C done  A done
quiet, no A (direct)                                  0.28 s  0.29 s  0.37 s       180       0
direct, first come first served                      29.90 s 44.49 s 47.81 s       180      16
direct, A at low priority                             6.52 s 14.71 s 16.31 s       180      16
gateway: priority + A capped at 1                     7.06 s 12.11 s 13.55 s       180       7

Priority does most of the work: B and C's median wait falls from 29.90 s to 6.52 s. Capping A at one request shortens the worst waits, from 14.71 s to 12.11 s at the 95th percentile. A pays for it, finishing 7 summaries in the window against 16.

B and C still wait about 7 s, close to one 4,096-token prefill (7.1 s in Step 2). A request that arrives while A's prompt is being read waits for that read to finish.

🔎 Tip: size the cap above what the tenant needs. B sends a request every 2 seconds and each one spends about 12 seconds on the GPU under load, so B needs about 6 slots. With B capped at 4, the same run gave B a median time inside vLLM of 11.61 s and a median wait for a gateway slot of 37.77 s (day3/cap4/). A cap is for the customer you want to slow down.

Step 7 — Check the Ledger Against vLLM

A bill is only as good as its meter. day3/bill.py adds up the ledger and compares it with vLLM's own prompt_tokens_total and generation_tokens_total counters, read before and after the run:

ledger vs vLLM counters: prompt 122,442 vs 122,442, output 12,416 vs 12,416; match: True

Every token vLLM processed in the run is attributed to a tenant, and no tenant is charged for a token vLLM did not process.

Step 8 — Send the Bill

The VM's on-demand price comes from the Cloud Billing Catalog API, fetched during the build (day3/pricing.us-west2.txt):

T4 0.410000 + 2 vCPU x 0.037970 + 7.5 GiB x 0.005089 = 0.5241 USD/hour
x 730 hours = 382.60 USD/month

bill.py prices each request in GPU seconds, using the prefill fit from Step 2 plus 9.34 ms per output token, and splits the month by each tenant's share. Over the gateway run:

tenant   requests refused  prompt tok  output tok token share     GPU s GPU share  bill USD
A               7       0      28,718         896       22.0%      58.7     27.2%    104.22
B              90       0      46,803       5,760       39.0%      78.4     36.4%    139.11
C              90       0      46,921       5,760       39.1%      78.5     36.4%    139.26

A sent 22.0% of the tokens and used 27.2% of the GPU. A flat price per token would undercharge the long-prompt customer and pass the difference on to the other two. Weighting prompt tokens by their measured prefill cost puts it back.

Compare and Contrast

	Straight to vLLM	vLLM with priority	Through the gateway
B and C, median first token under load	29.90 s	🥈 6.52 s	🥇 7.06 s
B and C, 95th percentile	44.49 s	🥈 14.71 s	🥇 12.11 s
Keys per customer	❌ no	❌ no	🟢 yes
Adapter only its owner can call	❌ no	❌ no	🟢 yes
Prefix cache per customer	❌ no	❌ no	🟢 yes
Per-customer bill, checked against vLLM	❌ no	❌ no	🟢 yes

So, Which One?

Turn on --scheduling-policy priority for any shared vLLM: on these runs it is the single biggest change, and it costs one flag. Selling the GPU takes the gateway, because priority alone still lets any client claim any priority, call any adapter and read any cache. The gateway sets all three from the key.

Prior Work

The T4 rig: the VM, vLLM 0.29.0 install, Turing patch and the vllm-t4 script come from my earlier gemma4-dev T4 serving work.
gemma-rust: reqwest client patterns for the OpenAI-compatible endpoint.
dev-hacktoberfest (Week 1): every result checked by code that did not produce it.
Open source used: vLLM (Apache 2.0), Gemma 4 E2B (Gemma terms), axum, reqwest, tokio.

Why Does Open Innovation Matter?

Every part of this business needs the weights.

The model runs on hardware the operator controls. That is possible only with downloadable weights, and it is what lets the operator promise a customer that its emails never reach a model vendor.
One base model, one small adapter per customer. Per-customer fine-tunes that share one GPU and one batch need the weights and an engine that serves adapters side by side. vLLM's multi-LoRA does that on a 2018-era T4.
The cost is flat. The operator pays 382.60 USD a month whatever the customers send. A heavy customer costs the others waiting time, which the scheduler controls, and costs the operator nothing extra.
Everything can be checked. Engine, gateway and model are open, so the cache isolation and the bill are tested against vLLM's own counters, and a customer can read the code that bills them.

Prize Categories

Gemma: Gemma 4 E2B, served locally on a T4 with vLLM, with per-customer LoRA adapters on Google's quantization-aware 4-bit weights.

Summary

The goal of this project was to sell shares of one T4 to several small customers without one customer ruining the service for the others. The key to the solution was vLLM's priority scheduling, adapters and cache salt, controlled per customer by a gateway that holds the keys. The results were:

🟢 A long-prompt customer pushed other customers' first token from 0.28 s to 29.90 s. Priority brought it to 6.52 s; the gateway's cap on that customer brought the 95th percentile from 14.71 s to 12.11 s.
🟢 9 of 9 isolation checks pass: keys, adapter ownership and a prefix cache per customer.
🟢 The ledger matched vLLM's token counters exactly: 122,442 prompt and 12,416 output tokens.
🟢 A per-adapter customer runs in the same batch as base-model customers, at about the same speed.
⚠️ B and C still wait about one long prefill, around 7 s, while A is active. Chunking long prompts further, or giving A its own hours, are the next steps.
❌ A's throughput falls from 16 summaries to 7 in the window. That is the price A pays for a cheaper share.

Scope: one T4 on one Google Cloud VM, vLLM 0.29.0, one 180-second run per configuration. The two direct runs and the gateway run used separate server starts with the same flags; in the direct runs C used the base model, through the gateway it used its adapter. The token bucket, the per-customer fair queue and a tenant-count sweep from SPEC.md are not built. The VM price excludes disk and network.

The strategy for renting out one T4 to several customers with Gemma 4 and vLLM was validated with an incremental step by step approach.

References

Don't Assume Your Cloud T4 Runs at Spec: Gemma 4 Decode Falls 7% as the Card Warms

xbill — Fri, 16 Oct 2026 15:14:10 +0000

This article provides a step by step guide to measuring how a cloud GPU's temperature changes LLM decode speed, on the NVIDIA Tesla T4 of a Google Cloud VM serving Gemma 4 E2B with vLLM. Every log, script and number is committed.

https://github.com/xbill9/t4-thermal-drift

The same T4, on the same VM, serving the same model the same prompt, decoded at 115.3 tokens per second at 45 °C and 106.8 at its steady 77 °C: 7.4% slower, for as long as the load lasts. The card runs at its 70 W power limit the whole time. A hotter chip draws more power at the same clock, so the limit pulls the clock down from about 1,457 MHz to about 1,267. The thermal throttle never fires.

What is this project trying to Do?

A rented GPU is sold by name and by the hour. A T4 is a T4, and the spec sheet gives it one number: 8.1 TFLOPS of FP32, 65 TFLOPS of FP16. Benchmarks on it are run for a minute or two and quoted as the card's speed.

This project asks whether that number holds when the card is hot. It logs the card's temperature, clock, power and throttle state every 100 ms while it decodes, and lines every request up with the card's state at the moment its tokens were generated.

The answer shapes three things a customer cares about: what speed a T4 delivers under sustained load, how much a short benchmark overstates it, and whether an A/B test run back to back is fair.

At This Point You Should Have…

A Google Cloud VM with one NVIDIA T4 and root access, so nvidia-smi -lgc can lock the clock
vLLM serving a model on the card. This one runs vLLM 0.29.0 with xbill9/gemma-4-E2B-it-qat-q4_0-w4a16-ct-text-emb4, Gemma 4 E2B in 4-bit weights, fp16
A clone: git clone https://github.com/xbill9/t4-thermal-drift

The Hardware

name, driver_version, memory.total [MiB], power.limit [W], persistence_mode
Tesla T4, 615.71.09, 15360 MiB, 70.00 W, Enabled

    Max Clocks
        Graphics                                       : 1590 MHz
        SM                                             : 1590 MHz

        Current Power Limit                            : 70.00 W
        Default Power Limit                            : 70.00 W
        Max Power Limit                                : 70.00 W


VM	`n1-standard-2`, `us-west2-b`
GPU	Tesla T4, passive cooling, 70 W limit, 1,590 MHz maximum clock
Server	vLLM 0.29.0, one request at a time, greedy, 256 output tokens

The T4 has no fan. It depends on the host's airflow, which the cloud provider controls. Google runs the card at the highest power limit it allows.

Step 1 — Log the Card Every 100 ms

thermal.py runs nvidia-smi in the background for the whole of each run, writing eight fields every 100 ms:

timestamp,temperature.gpu,clocks.sm,clocks.mem,power.draw,clocks_event_reasons.active,utilization.gpu,pstate

Each request records the wall time of every streamed token. analyze.py takes the samples that fall between a request's first and last token, so each decode rate is matched to the temperature, clock and power while it was generated.

clocks_event_reasons.active is the field that says why the clock is below maximum. The two values that matter here are 0x4, the software power cap, and 0x20 / 0x40, the software and hardware thermal slowdowns.

Step 2 — Run a Heat Soak From a Cool Card

The soak waits until the card is below 50 °C, starts the server, and sends one prompt back to back, greedy, 256 tokens each:

python3 thermal.py soak runs/2026-10-15-first-t4/soak-unlocked-5min-1 --minutes 5

soak-start {'temp_c': 41}
10 114.0 (54, 1515)
20 112.3 (63, 1545)
30 110.7 (68, 1470)
40 108.6 (74, 1440)
50 105.7 (79, 1500)
60 106.8 (76, 1425)
70 107.4 (76, 1455)
80 106.4 (77, 1500)
90 106.7 (77, 1500)
100 107.1 (76, 1440)
110 106.9 (76, 1305)
120 106.9 (77, 1455)
soak-end {'requests': 125, 'temp_c': 77}

Every tenth request prints its decode rate, then the card's temperature and SM clock read once, after the request. The speed falls as the temperature climbs, and both settle within about a minute. A single clock reading jumps by 200 MHz from one moment to the next, so Step 3 uses the 100 ms samples taken while each request decoded.

A 20-minute soak, run the day before, shows the same shape from a 57 °C start.

Step 3 — Read Decode Speed by Temperature

python3 bands.py runs/2026-10-15-first-t4

The 5-minute soak, every request grouped by the card's mean temperature while it decoded:

Card temperature	Requests	Decode tok/s	SM clock MHz	Power W
45–49 °C	4	115.3	1,457	69.3
55–59 °C	5	113.5	1,404	68.9
65–69 °C	10	111.1	1,350	68.2
70–74 °C	8	109.4	1,316	67.8
75–79 °C	81	106.9	1,267	66.9

The 20-minute soak:

Card temperature	Requests	Decode tok/s	SM clock MHz
55–59 °C	1	112.7	1,390
65–69 °C	6	110.8	1,346
70–74 °C	7	108.0	1,296
75–79 °C	462	105.6	1,241

== soak-unlocked-1: requests 486, below 95 tok/s 1
after 2 min: n 437, median 105.6 tok/s, trend +0.03 tok/s per 10 min, temp 76-78 C, SM median 1241 MHz
r(decode, temperature) -0.86  r(decode, SM clock) 0.80
== soak-unlocked-5min-1: requests 125, below 95 tok/s 0
after 2 min: n 73, median 106.8 tok/s, trend +1.01 tok/s per 10 min, temp 76-80 C, SM median 1266 MHz
r(decode, temperature) -0.97  r(decode, SM clock) 0.99

One request in the 20-minute soak ran at 54.1 tokens per second, 36 seconds in, and is left out of the bands and correlations.

Cool to hot costs 7.4% in one soak and 6.3% in the other, which started warmer. Once the card holds at 76–78 °C the rate stays flat for 18 minutes, so the loss follows temperature, and time on its own adds nothing. Per request, decode tracks the SM clock with a correlation of 0.99.

Step 4 — Find What Pulls the Clock Down

== clock event reasons seen, all runs (0x1 idle, 0x4 SW power cap, 0x8 HW slowdown, 0x20 SW thermal, 0x40 HW thermal)
soak-locked1200-1                             samples  13352 reasons ['0x1', '0x4'] max temp 81 C
soak-unlocked-1                               samples  13384 reasons ['0x1', '0x4'] max temp 81 C
soak-unlocked-5min-1                          samples   3509 reasons ['0x1', '0x4'] max temp 80 C

Across every run, 154,815 samples between them, the card reported two reasons only: idle, and the software power cap. The thermal slowdowns never fired, and the hottest reading was 83 °C. During the unlocked soaks the power cap was active in every temperature band, from 45 °C up.

So the card is at its 70 W limit from the first request. At 45 °C those 70 W buy 1,457 MHz. At 77 °C they buy 1,267.

Step 5 — Measure the Leakage at a Held Clock

Hold the clock fixed and the only thing left to change the power is the chip's temperature:

python3 thermal.py soak runs/2026-10-15-first-t4/soak-locked1200-1 --lock 1200

power at a held 1200 MHz, not capped (6485 samples of 100 ms):
  66 C n=   16 median 62.0 W
  68 C n=   19 median 63.4 W
  70 C n=   22 median 64.0 W
  72 C n=   26 median 64.3 W
  74 C n=   23 median 65.1 W
  76 C n=  261 median 65.7 W
  78 C n=  271 median 66.0 W
  slope 0.26 W per C

Same clock, same work, 4.0 W more at 78 °C than at 66 °C. A hotter chip leaks more current. NVIDIA's TensorRT documentation names the effect and names this card: "higher GPU temperature also leads to more leakage current in the circuits, which increases the power consumed by the GPU at a specific clock frequency. Therefore, for GPUs more likely to be power throttled like NVIDIA T4, poor cooling can lead to lower stabilized clock frequency with power throttling and, thus, worse performance."

Step 6 — Lock the Clock and See What It Costs

The same soak with the SM clock locked at 1,200 MHz:

Card temperature	Requests	Decode tok/s	SM clock MHz	Power-capped
65–69 °C	5	104.8	1,200	0%
70–74 °C	7	104.5	1,200	1%
75–79 °C	460	104.1	1,198	29%

Locking shrinks the drift from 7% to 0.7%. It does so by giving up speed: 104.1 tokens per second hot and locked, against 106.8 hot and unlocked and 115.3 cool. Even at 1,200 MHz the card reaches its power limit in 29% of samples once it is hot.

NVIDIA's own guidance says the same of locked clocks: "the average performance will not be as good as when the clock is floating."

Step 7 — Check What a Back-to-Back A/B Test Measures

The project began with a speculative-decoding comparison: Gemma 4 E2B alone against Gemma 4 E2B with its 4-bit drafter, three draft tokens, each on its own server, six prompts, three requests each. Run order and cooling were varied:

python3 order.py runs/2026-10-15-first-t4

order-cool50-nonefirst                     none>qat-n3  none: 111.1 tok/s n=18 56->71 C  qat-n3: 144.3 tok/s n=18 56->70 C  speed-up 1.298x
order-cool50-qatfirst                      qat-n3>none  qat-n3: 144.6 tok/s n=18 56->71 C  none: 111.5 tok/s n=18 55->71 C  speed-up 1.297x
order-nocool-nonefirst                     none>qat-n3  none: 111.3 tok/s n=18 56->71 C  qat-n3: 142.3 tok/s n=18 64->76 C  speed-up 1.279x
order-nocool-qatfirst                      qat-n3>none  qat-n3: 145.1 tok/s n=18 56->65 C  none: 108.5 tok/s n=18 63->77 C  speed-up 1.337x

With the card cooled to 50 °C before each server, the speed-up is 1.297x or 1.298x whichever runs first. Without the cool-down, the second server starts on a warm card and runs 1.9% to 2.5% slower than the same server run first, and the same speed-up reads anywhere from 1.279x to 1.337x.

What Do Google and NVIDIA Promise?

NVIDIA's T4 datasheet lists, in its specifications table:

Single-Precision       8.1 TFLOPS
Mixed-Precision        65 TFLOPS
(FP16/FP32)
INT8                   130 TOPS
Thermal Solution       Passive

No clock, no condition and no footnote is attached. Google's Compute Engine GPU page repeats 65, 130 and 260 TOPS for the T4, also without conditions. The Compute Engine SLA defines downtime as "loss of external connectivity or persistent disk access": it covers uptime and says nothing on speed.

The published figure is a maximum-clock figure. FP32 throughput is 2,560 CUDA cores times two operations per clock times the clock, so 8.1 TFLOPS needs 1,582 MHz, the card's 1,590 MHz maximum. The same arithmetic at the clocks measured here:

SM clock	FP32, by arithmetic	Share of the spec clock
1,590 MHz, the maximum	8.14 TFLOPS	100%
1,430 MHz, cool, 44 °C	7.32 TFLOPS	89.9%
1,241 MHz, hot, 77 °C	6.35 TFLOPS	78.1%
1,200 MHz, locked	6.14 TFLOPS	75.5%

These rows are arithmetic from the clock, not measured TFLOPS. Decode fell less than the clock did, 7.4% against 13.0%, because decoding one token also reads every weight from memory, and the memory clock holds.

🔎 Tip: Log the Clock Event Reasons With Every Benchmark

nvidia-smi --query-gpu=clocks_event_reasons.active,clocks.sm,temperature.gpu,power.draw -lms 100 costs nothing and turns "it seemed slower" into a reason code. A benchmark that quotes tokens per second without the clock it ran at leaves out the variable that moved the number by 7% here.

🔎 Tip: Cool to the Same Temperature Before Every Server

Stopping the server lets a T4 in this host fall from 71 °C to 50 °C in under four minutes. Loaded with a model and idle, the same card sat at 72–76 °C for 30 minutes and never got there. Stop the server, wait for the temperature, then start the next one.

Compare and Contrast

	Cool, unlocked	Hot, unlocked	Hot, locked 1,200 MHz
Card temperature	45–49 °C	75–79 °C	75–79 °C
Decode	🥇 115.3 tok/s	🥈 106.8 tok/s	🥉 104.1 tok/s
SM clock	1,457 MHz	1,267 MHz	1,198 MHz
How long it lasts	about a minute	as long as the load	as long as the load
Drift cool to hot	—	7.4%	0.7%

So, Which One?

Don't assume. Under sustained load this T4 delivers the hot number, 106.8 tokens per second on Gemma 4 E2B, and the cool number exists for the first minute. Quote the hot number, log the clock beside it, and cool every arm of an A/B test to the same temperature before it starts. Lock the clock when consistency matters more than speed, and count what it costs: 2.5% below hot and unlocked here.

The customer pays the same hourly rate for 115 or 107 tokens per second, and the cooling that decides between them belongs to the provider.

Summary

The goal of this article was to find whether a cloud T4's temperature changes LLM decode speed, and why. The key to the solution was logging the card's clock, power, temperature and throttle reason every 100 ms and matching each request to the samples taken while it decoded. The results were:

❌ Decode fell 7.4%, from 115.3 to 106.8 tokens per second, as the card warmed from 45 °C to its steady 77 °C, and 6.3% in a second soak that started warmer
⚠️ The card sits at its 70 W limit from the first request. At a held 1,200 MHz it drew 4.0 W more at 78 °C than at 66 °C, so the limit pulls the clock down as it heats
🟢 No thermal throttle fired in 154,815 samples, and the hottest reading was 83 °C
⚠️ Locking the clock at 1,200 MHz cut the drift to 0.7% and cost 2.5% against hot and unlocked
❌ Without a cool-down between servers, the same speculative-decoding speed-up read 1.279x to 1.337x depending on order. With one it read 1.297x and 1.298x
⚠️ Google and NVIDIA publish the T4's throughput with no condition attached, and it corresponds to the 1,590 MHz maximum clock

Scope: one Tesla T4 on one n1-standard-2 VM in us-west2-b, driver 615.71.09, vLLM 0.29.0, Gemma 4 E2B in 4-bit weights, one request at a time, measured 2026-10-15 and 2026-10-16. Two unlocked soaks (20 and 5 minutes) and one locked soak at 1,200 MHz; each order condition ran once. Other T4 hosts, batched serving and other models were not run. Parts of the analysis and writing were done with AI assistance (Claude); every figure comes from the committed output files.

The strategy for measuring thermal drift on a cloud T4 was validated with an incremental step by step approach.

References

Code, logs and this article's evidence: https://github.com/xbill9/t4-thermal-drift
NVIDIA T4 datasheet: https://www.nvidia.com/content/dam/en-zz/Solutions/Data-Center/tesla-t4/t4-tensor-core-datasheet-951643.pdf
NVIDIA TensorRT, Hardware/Software Environment for Performance Measurements: https://docs.nvidia.com/deeplearning/tensorrt/10.x.x/performance/hw-sw-environment.html
Google Compute Engine GPUs: https://docs.cloud.google.com/compute/docs/gpus
Compute Engine SLA: https://cloud.google.com/compute/sla
nvidia-smi documentation: https://docs.nvidia.com/deploy/nvidia-smi/index.html

EmbeddingGemma 2 on Cloud Run: Minimum Specs, and When an NVIDIA L4 Pays for Itself

xbill — Thu, 15 Oct 2026 22:50:48 +0000

This article provides a step by step sizing guide for EmbeddingGemma 2 text embeddings on Google Cloud Run. It measures the smallest CPU instance the model fits in, tests an int4 repack of the weights, and prices a burst of work on a 1 vCPU service against the same burst on an NVIDIA L4 GPU.

https://github.com/xbill9/embedding-cloudrun

What is this project trying to Do?

The first article in this series deployed EmbeddingGemma 2 to Cloud Run three ways and priced each per million documents: https://dev.to/gde/embeddinggemma-2-on-cloud-run-an-int8-cpu-service-against-an-nvidia-l4-15ac

EmbeddingGemma 2 is a compact model. With the vision and audio encoders dropped and the Linears stored as int8, the text model is a 203.4 MiB file. This follow-up asks the minimum-spec question: is a basic 1 vCPU instance good enough, and when does the L4 cost less?

Two numbers decide it. How much memory the CPU service needs, and how much an L4 bills for a short burst of work, idle time included.

The Short Answer

A basic 1 vCPU CPU service is good enough, and it is the cheapest option for most of what an embedding service does. Add the L4 only for big batch jobs.

Minimum spec: 1 vCPU, 2 GiB, int8 weights. Queries take a median 133 ms, a million documents cost USD 27.67, and an idle service costs nothing. 1 GiB is too small. Documents near 512 tokens need a batch size of 8 or less, or 4 GiB.
The L4 bills its idle time. Each isolated burst was billed about eight minutes of L4 time, USD 0.14, even when the work took seconds. The L4 wins once a single burst passes about 5,300 documents: a million documents cost USD 1.83 on it against USD 27.67 on the CPU.
An L4 kept warm costs USD 764 a month. The CPU service embeds 27.6 million documents for the same money.
The int4 repack is slower on CPU. The file is 28% smaller, but queries take 1.52x as long and it needs 3 GiB, because this server runs it in float32. Retrieval quality is about the same: SciFact nDCG@10 of 0.834 against 0.836 for int8.

In practice: serve queries and small indexing batches from the 1 vCPU int8 service. For a job of tens of thousands of documents, deploy the L4, run the job, and let it scale to zero.

Where do I start?

The strategy is an incremental step by step approach, smallest instance first.

The int8 repack is measured for memory on a laptop at one thread, the same thread count Cloud Run gives it. An int4 repack is built and deployed beside it. Then the L4 is swept for throughput, cold started repeatedly, and its billed seconds are read back from Cloud Monitoring. One script does the cost arithmetic from those files.

At This Point You Should Have…

The repository cloned and models/ staged with scripts/prepare_models.sh, as in the first article
The emb-cpu-repack service in us-central1 and emb-l4-original in us-east4, from scripts/deploy.sh
The gcloud CLI signed in, for the Cloud Monitoring query in Step 6

Step 1 — How Little Memory Does the CPU Service Need?

scripts/cpu_memory.py loads the int8 repack at one thread, embeds 32 copies of a document at a given batch size, and prints the process's peak memory. A fresh process each time, documents of 135 tokens and of 512 tokens (the model's limit):

for n in short long; do for b in 1 4 8 16 32; do python3 scripts/cpu_memory.py $b $n; done; done

batch  1 doc 135 tokens: peak RSS 1382 MiB
batch  4 doc 135 tokens: peak RSS 1369 MiB
batch  8 doc 135 tokens: peak RSS 1345 MiB
batch 16 doc 135 tokens: peak RSS 1468 MiB
batch 32 doc 135 tokens: peak RSS 1871 MiB
batch  1 doc 512 tokens: peak RSS 1385 MiB
batch  4 doc 512 tokens: peak RSS 1572 MiB
batch  8 doc 512 tokens: peak RSS 1893 MiB
batch 16 doc 512 tokens: peak RSS 2356 MiB
batch 32 doc 512 tokens: peak RSS 3551 MiB

The loaded model and PyTorch take about 1.35 GiB before any batch, so 1 GiB is too small at any setting. The service from the first article runs at 2 GiB with a batch of 32, which holds for 135-token documents at 1,871 MiB. A batch of 32 documents at 512 tokens peaks at 3,551 MiB.

So the 2 GiB service needs batches of 8 or fewer for long documents, or 4 GiB at a batch of 32.

🔎 Tip: Make Batch Size a Request Field

The server now takes batch_size per request, 1 to 256, default 32:

class EmbedRequest(BaseModel):
    texts: list[str] = Field(min_length=1, max_length=256)
    prompt_name: str | None = None
    dim: int = 768
    batch_size: int = Field(32, ge=1, le=256)  # texts per forward pass

A client sending long documents to a 2 GiB instance asks for "batch_size": 8 and the instance stays inside its memory limit without a redeploy. /health also reports rss_mib and peak_rss_mib, and on a GPU the GPU's total, allocated and peak memory, so the numbers in this article can be read from a live service.

Step 2 — Does an int4 Repack Shrink It Further?

repack/repack_text_w4a16_emb4.py stores every Linear in the text model as int4, in groups of 32 with a bf16 scale, in the compressed-tensors pack-quantized layout. Activations stay in float. The file is 146.1 MiB against the int8 repack's 203.4 MiB.

Quality on SciFact (300 queries over 5,183 abstracts, 768 dimensions) with repack/eval_scifact.py:

weights	nDCG@10	recall@10	top 10 shared with float32
🥇 float32 original	0.854	0.944	100%
🥈 int8 repack	0.836	0.926	85.5%
🥉 int4 repack	0.834	0.934	80.7%

Both repacks keep about 98% of the original's nDCG@10: 97.9% for int8 and 97.6% for int4.

Deployed as cpu-w4a16 on the same 1 vCPU shape, with three runs of each variant alternating on the same day:

python3 scripts/cost.py results/runs/cpu-repack-2026-10-15*.json results/runs/cpu-w4a16-*.json

variant        runs          query ms            docs/s  USD / 1M queries       USD / 1M docs
cpu-repack        3     133 (132-135)     1.0 (1.0-1.1)              4.25 27.67 (26.16-28.55)
cpu-w4a16         3     201 (190-216)     0.8 (0.8-0.8)              6.74 38.09 (37.20-39.43)

The server unpacks the int4 weights to float32 once at load, so every matrix multiply runs in float32, while the int8 weights run on PyTorch's int8 CPU kernels. Queries take 1.52x as long on int4 and documents cost 38% more. The int4 service also needs 3 GiB: loading plus one 32-document batch peaks at 2,127 MiB locally, against 1,856 MiB for int8.

The smaller file buys a smaller image, 580 MB compressed against 636 MB, and nothing else on a CPU. The int8 repack stays the minimum spec.

Step 3 — How Fast Is the L4 From a Client?

The first article timed the L4 inside the server: 331 documents per second, median of four runs. scripts/sweep.py measures what a client receives, across request sizes and concurrent requests, with the request size also used as the batch size:

python3 scripts/sweep.py l4-original $L4_URL --sizes 32,64,128,256 --concurrency 1,2,4 --requests 8

Documents per second as the client saw them, network and JSON included:

documents per request	1 at a time	2 at a time	4 at a time
32	81.4	148.1	162.4
64	102.6	🥇 171.8	167.4
128	111.9	166.1	152.6
256	115.9	155.8	151.9

The best cell is 171.8 documents per second, at 64 documents per request and two requests at a time. The server's own timer on one request at a time gives 328.3 documents per second at a batch of 32 and falls to 192.5 at 256, so smaller batches suit this GPU.

Peak GPU memory reached 14,956 MiB reserved by PyTorch at four requests of 256, out of 22,563 MiB on the L4. The model itself holds 516 MiB.

Every response carries 768 numbers per document as JSON. That transfer is the likely limit on the client rate; this was not tested.

Step 4 — Cold Starts on the L4

scripts/coldstart_series.sh waits 20 minutes with no traffic, sends one request, then a second, and repeats:

scripts/coldstart_series.sh $L4_URL 8 1200 results/coldstart-l4-series/run
python3 scripts/coldstart_summary.py results/coldstart-l4-series/run-*.txt

file                                   variant       first request s  model load s  the rest s  warm s
run-1.txt                              original                 38.9           6.4        32.4    0.09
run-2.txt                              original                 17.1           4.8        12.3    0.11
run-3.txt                              original                 27.4           5.9        21.5    0.11
run-4.txt                              original                 17.3           4.7        12.7    0.10
run-5.txt                              original                 27.9           6.3        21.5    0.10

Five samples completed. The first request took 17.1 to 38.9 seconds, of which the model load was 4.7 to 6.4. A warm request took about 0.1 seconds.

The 1 vCPU services took longer: 58.4 and 83.5 seconds in the first article, and 134.8 seconds for the int8 repack and 164.9 for int4 on 2026-10-15.

🔎 Tip: A Cold Start Series Needs a Quiet Service

Each sample is cold only if nothing else reaches the service during the wait. A health check, a test, or a sweep in that window starts an instance, and the next sample measures a warm request. Run the series alone, and leave the wait longer than the time Cloud Run takes to scale an idle instance to zero.

Step 5 — What Does One Burst Cost on the L4?

A GPU service on Cloud Run is billed for its instance: every second from start until it scales to zero, idle seconds included. Cloud Monitoring records that time as run.googleapis.com/container/billable_instance_time. scripts/billed_time.py reads it per minute for one service:

python3 scripts/billed_time.py emb-l4-original us-east4 2026-10-15T18:15:00Z 2026-10-15T20:10:00Z

minute ending (UTC)     billed s  instances
2026-10-15T18:32:00         20.8  {'active': 1.0, 'idle': 0.0}
2026-10-15T18:33:00         60.0  {'active': 1.0, 'idle': 0.0}
...
2026-10-15T18:42:00         60.0  {'active': 1.0, 'idle': 0.0}
2026-10-15T18:53:00         24.7  {'active': 1.0, 'idle': 0.0}
...
total billed instance seconds 2731.5

Each cold cycle in Step 4 served two health requests and was billed 265 to 621 seconds, a mean of 474. At USD 1.0465 per hour, one isolated burst costs about USD 0.14 before it embeds anything.

The 1 vCPU service uses request-based billing: it is charged only while a request is in flight, and an idle instance costs nothing.

Step 6 — How Big Must a Burst Be?

scripts/tradeoff.py takes the billed seconds per burst from Step 5, the client rate from Step 3 and the int8 service's cost per document from Step 2, and does the arithmetic:

python3 scripts/tradeoff.py

L4 USD/hour 1.0465; billed s per isolated burst 265-621, mean 474
  one isolated burst costs USD 0.1379 (mean), the price of 4,982 documents on the CPU
L4 best end-to-end 171.8 docs/s -> USD 1.69 per 1M docs while busy; CPU USD 27.67

burst of N documents (one cold L4 cycle vs CPU)
         N   CPU USD    L4 USD
       100    0.0028    0.1380
     1,000    0.0277    0.1396
    10,000    0.2767    0.1548
   100,000    2.7670    0.3071
 1,000,000   27.6700    1.8302
break-even burst size 5,307 documents

L4 kept warm all month (traffic at least every ~8 min): USD 764/month, equal to 27.6M documents/month on the CPU

Below about 5,300 documents in a burst, the 1 vCPU service is cheaper. Above it, the L4 is cheaper, and at a million documents it costs USD 1.83 against USD 27.67.

An L4 that receives traffic at least every few minutes never scales to zero and costs USD 764 a month. The 1 vCPU service embeds 27.6 million documents a month for the same money.

Compare and Contrast

	1 vCPU int8	1 vCPU int4	NVIDIA L4
Weights	🥈 203.4 MiB	🥇 146.1 MiB	🥉 float32 original, run in bf16
Memory needed	🥇 2 GiB, or 4 GiB for long documents at batch 32	🥈 3 GiB	🥉 16 GiB and a GPU
Query, median	🥈 133 ms	🥉 201 ms	🥇 57 ms
USD per 1M documents, busy	🥈 27.67	🥉 38.09	🥇 1.69 from a client
One isolated burst of 1,000 documents	🥇 USD 0.03	🥈 USD 0.04	🥉 USD 0.14
Cold start	🥈 58–135 s	🥉 165 s	🥇 17–39 s
SciFact nDCG@10	🥈 0.836	🥉 0.834	🥇 0.854 for the float32 weights

So, Which One?

For queries and for indexing in batches under about 5,300 documents, the 1 vCPU int8 service is good enough and is the cheapest. It answers a query in a median 133 ms, costs nothing between requests and needs no GPU quota. Give it 2 GiB, with batch_size at 8 for documents near 512 tokens, or 4 GiB.

For a batch job of tens of thousands of documents or more, deploy the L4, run the job, and let it scale to zero. One burst bills about eight minutes of L4 time, and past 5,307 documents that is cheaper than the CPU.

The int4 repack is the smallest file and the slowest CPU service here, because this server runs it in float32. A runtime with int4 kernels may change that; none was tested.

Step 7 — Teardown

scripts/down.sh

It deletes the services in both regions. All of them were deployed with --min-instances 0, so an idle service costs nothing while it stays deployed.

Summary

The goal of this article was to find the minimum Cloud Run spec for EmbeddingGemma 2 and the point where an NVIDIA L4 costs less. The key to the solution was measuring the L4's billed instance time per burst from Cloud Monitoring, idle time included, and pricing a burst on each service from that. The sizing results were:

🟢 The int8 repack runs on 1 vCPU and 2 GiB, at a median 133 ms per query and USD 27.67 per million documents
🟢 The L4 costs USD 1.69 per million documents at its best client rate of 171.8 documents per second
🟢 The L4 is cheaper for bursts above 5,307 documents; the 1 vCPU service is cheaper below
🟢 L4 cold starts took 17.1 to 38.9 seconds over five samples, with the model load under 7 seconds
⚠️ Each isolated L4 burst was billed 265 to 621 seconds, a mean of 474, about USD 0.14
⚠️ 512-token documents at a batch of 32 need 3,551 MiB; at 2 GiB they need a batch of 8 or fewer
⚠️ An L4 kept warm all month costs USD 764, the 1 vCPU price of 27.6 million documents
❌ The int4 repack is 28% smaller and 1.52x slower per query on CPU, and needs 3 GiB
❌ 1 GiB is too small for any batch size

Scope: one Google Cloud project, CPU services in us-central1 and the L4 in us-east4, Cloud Run prices from the Billing Catalog on 2026-10-10, measured 2026-10-15. Memory was measured locally on a laptop at one thread, not inside Cloud Run. The int8 and int4 services ran three benchmark runs each, alternating. The L4 sweep ran once, with eight requests per cell, from one client outside Google Cloud. The cold start series stopped after five of eight samples, each cycle a single request after 20 minutes idle; the billed time per burst comes from those five cycles. The break-even assumes each burst runs at the sweep's best client rate.

The strategy for sizing EmbeddingGemma 2 on Cloud Run was validated with an incremental step by step approach.

References

Repository: https://github.com/xbill9/embedding-cloudrun
First article in this series: https://dev.to/gde/embeddinggemma-2-on-cloud-run-an-int8-cpu-service-against-an-nvidia-l4-15ac
EmbeddingGemma 2 on Hugging Face: https://huggingface.co/google/embeddinggemma-2
int8 repack: https://huggingface.co/xbill9/embeddinggemma-2-w8a8-ct-text-emb4
int4 repack: https://huggingface.co/xbill9/embeddinggemma-2-w4a16-ct-text-emb4
Cloud Run GPU support: https://cloud.google.com/run/docs/configuring/services/gpu
Cloud Run billing settings: https://cloud.google.com/run/docs/configuring/billing-settings
Cloud Run pricing: https://cloud.google.com/run/pricing
Cloud Run metrics: https://cloud.google.com/monitoring/api/metrics_gcp_p_z#gcp-run

Gemma 4's MTP Drafter, Explained: Speculative Decoding and a 4-Bit Drafter at 60% of the Size

xbill — Thu, 15 Oct 2026 01:18:39 +0000

This article provides a step by step introduction to speculative decoding with Gemma 4's multi-token prediction (MTP) drafter, served by llama.cpp on a local, desktop hosted GPU enabled system. It is the first of a short series: this part explains what the drafter is and why you would use it, and the next part has the measurements.

https://github.com/xbill9/gemma4-dev/tree/main/local-llamacpp-3050oem-2b-q4_0

Gemma 4 ships a small companion model for every size, a drafter that guesses the next few tokens so the main model can check them in one pass. Google's QAT drafters are trained onto a 4-bit grid; stored that way they shrink from 97.8 MB to 59.2 MB and are published on Hugging Face.

What is this project trying to Do?

The rigs in this repository serve Gemma 4 E2B and E4B from local hardware: an RTX 3050 OEM in a desktop and a GTX 1650 Ti in a laptop. Both already run the exact 4-bit QAT builds of the main models. Google also publishes a drafter for each Gemma 4 model, and llama.cpp can now run it.

This part covers three things: how speculative decoding works, how Gemma 4's drafter differs from the usual draft model, and the exact Q4_0 repacks of Google's QAT drafters that these rigs serve.

Why Is Generating One Token at a Time Slow?

A language model writes one token per forward pass, and every pass reads every weight. For one user on a small GPU, the arithmetic is light and the memory traffic is heavy: the card spends most of each step moving weights, not computing with them.

That leaves spare compute on every step. Checking several tokens in one pass reads the same weights once and does a little more arithmetic, so it takes little more time than producing one. Speculative decoding exists to use that spare capacity.

How Does Speculative Decoding Work?

Each step has two parts.

Draft. A small, fast model guesses the next few tokens.
Verify. The main model, called the target, runs one forward pass over all the guesses at once and checks each one against what it would have produced.

The target accepts guesses from the front until the first one it disagrees with. It discards the rest and adds one token of its own at that point. A step therefore always produces at least one token, and as many as the number drafted plus one when every guess is right.

The output is the target's: every token either matches what the target would have chosen or is the target's own replacement. The drafter only changes how many tokens each expensive pass produces.

Google's model card puts it this way:

When used in a Speculative Decoding pipeline, the draft model predicts several tokens ahead, which the target model then verifies in parallel.

What Makes Gemma 4's Drafter Different?

A classic draft model is a smaller copy of the same family, with its own layers, its own KV cache and its own pass over the whole conversation. Gemma 4's drafter is an add-on to the target. Google calls the approach multi-token prediction and describes the drafter as a lightweight 4-layer model that uses the target model's activations and KV cache.

The E2B drafter's config.json on Hugging Face shows how small it is:

"architectures": ["Gemma4AssistantForCausalLM"],
"backbone_hidden_size": 1536,
"hidden_size": 256,
"num_hidden_layers": 4,
"num_kv_shared_layers": 4,
"layer_types": ["sliding_attention", "sliding_attention", "sliding_attention", "full_attention"],
"vocab_size": 262144,
"tie_word_embeddings": true

Property	What it means
`backbone_hidden_size: 1536`	the drafter reads the E2B target's 1536-wide hidden state (2560 for E4B)
`hidden_size: 256`, 4 layers	a tiny transformer of its own
`num_kv_shared_layers: 4`	all four layers reuse the target's KV cache; the drafter computes queries only
`vocab_size: 262144`	the same tokenizer as the target, so its guesses are the target's tokens

Two consequences follow. The drafter keeps no separate cache of the conversation, so its memory cost stays its weights and a small working buffer. And it starts each guess from what the target already worked out for the current token, which is why so small a model can guess well.

The checkpoint also carries a clustered shortcut for its output layer (num_centroids: 2048). llama.cpp's converter leaves that out and scores the full vocabulary.

Why Would You Use a Drafter?

🟢 Faster answers for one user. On a GPU waiting on memory, each pass that verifies two or three tokens does the work of two or three passes. Google's card claims up to 3x.
🟢 No retraining and no change to the main model. The target file, its settings and its context stay as they are; the drafter is one more file and two server flags.
🟢 Small memory cost. The drafter shares the target's KV cache, so it costs its weights and a small working buffer.
⚠️ The gain depends on the text. Predictable output such as code is guessed well; free-form stories are guessed less often, and each rejected guess is wasted work.
⚠️ The gain depends on load. With many users at once the GPU is already busy, and spare capacity for checking guesses shrinks. Part 2 measures where it stops paying.
⚠️ Greedy text can change at near-ties. Where two words score almost the same, checking several tokens in one batched pass can tip the choice the other way; batched arithmetic rounding differently is the likely cause, untested here. The answers stay the target's; part 2 shows the scores.

Which Drafter Goes With Which Model?

Google publishes two drafters per Gemma 4 size: -it-assistant for the bf16 model and -it-qat-q4_0-unquantized-assistant for the QAT builds. Google's QAT model card is explicit about pairing them:

When using multi-token prediction (speculative decoding) with an assistant model alongside a QAT target model, the assistant model must also be a QAT checkpoint with the same precision to ensure compatibility.

These rigs serve the exact Q4_0 QAT builds of E2B and E4B, so they take the QAT drafters.

Step 1 — Repack the QAT Drafter Exactly

Google ships the QAT drafter as bf16, but the values already lie on a 4-bit grid: in every weight matrix, each block of 32 holds at most 16 distinct values. That is Q4_0's own layout, so the weights can be stored at 4 bits without rounding them again.

llama.cpp's converter writes the drafter as Q8_0. gguf_exact.py, the same builder behind the exact main-model GGUFs, finds each block's trained step and stores the drafter as Q4_0 instead. No block was off the grid.

E2B drafter	File	Error vs Google's bf16
bf16, Google's weights	170.2 MB	0
Q8_0, llama.cpp converter	97.8 MB	0.00494
Exact Q4_0	59.2 MB 🥇	0.00138 🥇

The exact Q4_0 file is 39% smaller than Q8_0 and 3.6x closer to Google's weights by relative error (arithmetic from the build reports). The E4B drafter comes out the same: 59.7 MB against 98.7 MB.

Step 2 — Download It From Hugging Face

Both drafters are published, with their build script and evidence:

xbill9/gemma-4-E2B-it-qat-assistant-q4_0-exact-gguf
xbill9/gemma-4-E4B-it-qat-assistant-q4_0-exact-gguf

In the rig, make draft downloads the pinned revision and checks its hash:

060f5a0158f8120d25137866b90e04e21cd0d2d23de1db7a9cc3b7f00e60bd11  gemma-4-E2B-it-qat-assistant-q4_0-exact.gguf
dc965cb963e56b86fbc04db8efec83163886085e88514574650a57fa55a2ee24  gemma-4-E4B-it-qat-assistant-q4_0-exact.gguf

Step 3 — Serve the Target With the Drafter

llama.cpp runs Gemma 4's drafter as its own speculative type, draft-mtp. Two flags added to the usual server command:

llama-server -m gemma-4-E2B-it-q4_0-exact.gguf -ngl 99 -fa 1 \
  -md gemma-4-E2B-it-qat-assistant-q4_0-exact.gguf \
  --spec-type draft-mtp -ngld all --spec-draft-n-max 2

Flag	Meaning
`-md`	the drafter file
`--spec-type draft-mtp`	run it as a Gemma 4 MTP drafter on the target's state
`-ngld all`	put the drafter on the GPU with the target
`--spec-draft-n-max 2`	guess at most two tokens per step

The server log confirms the drafter loaded and reports, per request, how many guesses the target accepted:

I common_speculative_init_result: loading draft model '/home/xbill/models/gemma-4-E2B-it-qat-assistant/gemma-4-E2B-it-qat-assistant-q4_0-exact.gguf'
I slot print_timing: id  0 | task 8 | draft acceptance = 0.84656 (  160 accepted /   189 generated), mean len =  2.68

That request was a coding prompt: 160 of 189 guesses accepted, and each target pass produced 2.68 tokens on average.

How Much Memory Does the Drafter Cost?

Measured on the RTX 3050 at 8,192 tokens of context, with the E2B target alone at 2,795 MiB:

Drafter added	GPU memory	Cost of the drafter
none	2,795 MiB	—
bf16	2,983 MiB	+188 MiB
Q8_0	2,915 MiB	+120 MiB
Exact Q4_0	2,875 MiB 🥇	+80 MiB 🥇

The exact drafter costs a third less memory than the Q8_0 conversion. With E4B the pattern holds: +84 MiB against +122. On a 4 GB laptop card every one of those megabytes is context or headroom.

The smaller file also leaves the guesses unchanged: Q4_0, Q8_0 and Google's bf16 drafter had the target accept the same share of drafted tokens. Part 2 covers what that means for speed.

🔎 Tip: Where the Drafter Runs

llama.cpp's draft-mtp is the route measured here. Two others exist:

Hugging Face Transformers: pass assistant_model= to generate(), which is the route Google documents.
vLLM: --speculative-config '{"method":"mtp", ...}', measured in this repository on a Cloud TPU v5e.

🔎 Tip: Start With Two Drafted Tokens

Every extra guess must be right for the one after it to count. Two drafted tokens per step was the best setting on the RTX 3050 for both models, and one on the GTX 1650 Ti laptop card. Raise it only for highly predictable output such as code.

Compare and Contrast

	Classic draft model	Gemma 4 MTP drafter
What it reads	the conversation, from scratch	the target's hidden state and KV cache
Own KV cache	yes, grows with context	none
Size	a smaller model of the family	4 layers, 256 wide
Tokenizer	must match	same 262,144-token vocabulary

E2B QAT drafter	File	GPU memory added
bf16	170.2 MB	+188 MiB
Q8_0	97.8 MB	+120 MiB
Exact Q4_0	59.2 MB 🥇	+80 MiB 🥇

So, Which One?

For a Gemma 4 QAT target served by llama.cpp, the exact Q4_0 QAT drafter. It matches the target's precision as Google requires, it is the smallest file, the closest to Google's weights and the cheapest in memory, and the target accepts its guesses as often as it accepts the larger files'.

Turn it on for one user, or a few, on a GPU that waits on memory. The next part measures how much faster that is, on which kinds of text, and where adding users makes the drafter stop paying.

Summary

The goal of this article was to explain speculative decoding with Gemma 4's MTP drafter and to introduce exact 4-bit repacks of Google's QAT drafters. The key to the solution was Google's drafter design, a 4-layer model that reuses the target's state and KV cache, trained onto the same 4-bit grid as the QAT target. The results were:

🟢 A drafter guesses, the target verifies several tokens in one pass, and every token is still the target's.
🟢 Gemma 4's drafter keeps no KV cache of its own: 4 layers of width 256 that read the target's hidden state.
🟢 Exact Q4_0 QAT drafters on Hugging Face: 59.2 MB for E2B and 59.7 MB for E4B, 39–40% smaller than llama.cpp's Q8_0 conversion and 3.6x closer to Google's weights.
🟢 +80 MiB of GPU memory for the E2B drafter, against +120 MiB for Q8_0 and +188 MiB for bf16.
⚠️ The gain depends on the text and on the number of users; part 2 has the measurements.
⚠️ Greedy text can change at near-ties when verification runs in batches.

Scope: Google's QAT drafters for Gemma 4 E2B and E4B, repacked with gguf_exact.py and served by llama.cpp 77c4540ce on one RTX 3050 OEM (8 GB) at 8,192 tokens of context on 2026-10-14. The descriptions of the drafter come from Google's model cards, its MTP documentation and the checkpoints' config.json.

The strategy for adding Gemma 4's MTP drafter to a local llama.cpp deployment was validated with an incremental step by step approach.

References

Gemma 4 on a Used 8 GB RTX 3050: 116 Tokens/s for One User, 618 for 64

xbill — Thu, 15 Oct 2026 00:46:07 +0000

This article provides a step by step deployment guide for Gemma 4 E2B and E4B to a local, desktop hosted GPU enabled system: an RTX 3050 OEM with 8 GB, in a 2014-era Haswell-E desktop. A suite of Python MCP tools is built to simplify management of the llama.cpp hosted deployment.

https://github.com/xbill9/gemma4-dev/tree/main/local-llamacpp-3050oem-2b-q4_0

A used RTX 3050 OEM, listed at $179.97, serves Gemma 4 E2B at 116 tokens per second for one user and 618 tokens per second across 64 users, with the model's full 131,072-token context in 3.6 GB of its 8 GB. Against an AWS T4 on demand, the card costs as much as 342 hours of rental.

What is this project trying to Do?

Earlier rigs in this repository served Gemma 4 from rented TPUs and GPUs, and from hardware already on the desk: a 2021 laptop with a 4 GB GTX 1650 Ti, its CPU, and two Snapdragon ARM laptops. This rig asks what the cheapest dedicated card adds. The RTX 3050 OEM is an 8 GB Ampere card with tensor cores, and used ones were listed on eBay at about $145 to $210 on 2026-10-14.

The rig serves the same model file as the 1650 Ti rig, byte for byte, with the same server flags. The questions are single-user speed, how much context fits, how many users one card carries, whether the bigger E4B model is worth its cost, and what all of it costs against the alternatives.

Speculative decoding with Google's drafter models is a separate article. Every figure here was measured without it.

At This Point You Should Have…

An Nvidia GPU with 8 GB of memory. This one is an RTX 3050 OEM: GA106, compute capability 8.6, tensor cores, 120 W power limit.
The CUDA toolkit (13.4 here) and cmake.
About 7 GB of disk for both model files.
A clone of xbill9/gemma4-dev.

The Hardware

power.limit [W], clocks.max.sm [MHz], clocks.max.memory [MHz]
120.00 W, 2100 MHz, 7001 MHz


GPU	NVIDIA GeForce RTX 3050 OEM, PCI `10de:2508` (GA106), compute capability 8.6, 8192 MiB
Driver	615.78.08
CPU	Intel i7-5820K, 6 cores / 12 threads, AVX2
Board / RAM	ASUS X99-A, 46 GiB
OS	Debian sid, gcc 16.2.0, CUDA 13.4

The OEM 3050 came in prebuilt desktops and is a different configuration from the retail RTX 3050 8 GB. The host is a 2014-era X99 board; with every layer on the card, the CPU barely matters.

Step 1 — Build llama.cpp for the Card

export PATH=/usr/local/cuda/bin:$PATH
make build

The target runs cmake with GGML_CUDA=ON, CMAKE_CUDA_ARCHITECTURES=86 and GGML_NATIVE=ON, then builds llama-server and llama-bench. It printed the commit, 77c4540ce, which goes into tpu.env as LLAMA_CPP_COMMIT. The build took 10 minutes 24 seconds at -j10.

Step 2 — Download the Exact QAT GGUF

make download

The target downloads the file and checks it against MODEL_SHA256 in tpu.env. The file is xbill9/gemma-4-E2B-it-qat-q4_0-exact-gguf, SHA-256 419db9a6bf3bc15d770c85ccf9216827aa88c64fe8f3d72fbe5674f48efe2dc8. It is Google's quantization-aware-trained (QAT) Gemma 4 E2B with every trained 4-bit value packed into Q4_0 exactly. On the 1650 Ti it measured 32x closer to bf16 by KL divergence than Google's own GGUF, and faster.

Step 3 — Configure the Server

tpu.env is the rig's configuration, and every value in it carries the measurement that set it:

N_GPU_LAYERS=99
CONTEXT_SIZE=131072
KV_CACHE_TYPE=f16
FLASH_ATTENTION=1
THREADS=6
PARALLEL_SLOTS=1
REASONING=off
OVERRIDE_TENSORS=per_layer_token_embd\.weight=CUDA0

Setting	Why
`-ngl 99`	every layer on the card
`-c 131072`	the model's trained maximum; Step 8
`-fa 1`	flash attention: 1.133x prefill, 1.071x decode
`-ctk f16 -ctv f16`	`q8_0` costs 7.4% of decode
`-ot per_layer_token_embd…=CUDA0`	the per-layer embedding table on the card; Step 7
`--reasoning off`	answers start at once; a request can still ask for thinking

Step 4 — Start It and Check What Is Serving

make serve
nvidia-smi --query-compute-apps=pid,process_name,used_memory --format=csv

pid, process_name, used_gpu_memory [MiB]
47321, /home/xbill/llama.cpp/build/bin/llama-server, 1510 MiB

That is the first-light configuration: 8,192 tokens of context and the per-layer table left in host memory. The MCP server's attestation reads the device from the running process, since a CPU-only build would answer on the same port:

device=gpu · pid=47321 · -ngl 99 · mapped: ggml-cuda, libcublas, libcuda, libcudart · /home/xbill/llama.cpp/build/bin/llama-server
mismatch: None

One chat request at temperature 0, "In one sentence, what is a TPU?":

A TPU, or Tensor Processing Unit, is a specialized hardware accelerator designed to significantly speed up the training and inference of machine learning models, particularly those involving deep learning.

The server's own timings for that request: 34 tokens at 102.56 tokens per second.

Step 5 — Measure Single-User Speed

llama-bench, a 512-token prompt and 128 generated tokens, five repetitions per cell, the card cooled to 40 °C before each cell:

fa1-f16-t6 pp512=3859.96±111.73 tg128=116.14±0.25 77c4540ce
fa0-f16-t6 pp512=3408.19±110.65 tg128=108.46±0.22 77c4540ce
fa1-q8-t6 pp512=3708.48±118.09 tg128=107.57±0.23 77c4540ce
fa1-f16-t4 pp512=3838.78±111.16 tg128=115.87±0.22 77c4540ce
fa1-f16-t6-rep pp512=3835.54±102.85 tg128=115.90±0.22 77c4540ce

116.14 tokens per second generated, 3,860 tokens per second of prompt read. The repeated last cell lands within 0.6% of the first, and the card held 1890 to 1905 MHz without passing 44 °C.

On the 1650 Ti the same file generates at about 81.7 tokens per second and reads prompts at about 345. That makes this card about 1.42x faster at generating and 11x faster at reading. Generating one token reads every weight once, so it follows memory speed, where the two cards are closer. Reading a prompt is matrix work, which this card runs on tensor cores and the 1650 Ti has none of.

Step 6 — Keep the Defaults That Measure Faster

Setting	Prompt	Generate	Kept
flash attention off	0.883x	0.934x	❌ on
`q8_0` KV cache	0.961x	0.926x	❌ f16
4 CPU threads in place of 6	0.995x	0.998x	⚠️ either

The KV cache for E2B is small, so quantizing it saves memory the card has plenty of and costs speed. Thread count makes no difference because the CPU does no model work.

Step 7 — Move the Per-Layer Embedding Table onto the Card

The E in E2B is a per-layer embedding table, per_layer_token_embd, 1.32 GB and 51% of the file. llama.cpp leaves it memory-mapped in host RAM by default and reads one row per token. With 6 GB of the card unused, the rig moves it with -ot. Card against host, mean of two passes each, order host, card, card, host:

test	host	card	card / host
pp512	3845.18	3982.67	1.036
pp2048	3759.36	3848.42	1.024
tg128	115.91	115.98	1.001

A 2.4 to 3.6% gain on prompts, nothing on generation, for 1,260 MiB. Both card passes beat both host passes at both prompt lengths. The 1650 Ti measured no gain from the same move: this card reads prompts 11x faster, so fetching rows from host memory is a bigger share of its work.

Step 8 — Serve the Full 131k Context

Gemma 4 E2B was trained to 131,072 tokens. Only its 3 full-attention layers keep a cache that grows with context; the 12 sliding-window layers stay at 1,024 cells.

`-c`	KV, full-attention	KV, sliding	process on the card
8,192	48 MiB	12 MiB	2770 MiB
32,768	192 MiB	12 MiB	2938 MiB
131,072	768 MiB	12 MiB	3610 MiB

The full context costs 840 MiB over 8,192 and leaves 4.5 GB of the card free. Short chats do not notice: 256 generated tokens ran at 108.30 tokens per second at 8,192 and 107.99 at 131,072.

Long conversations do slow down, because every new token attends to everything before it:

tokens already in the conversation	prompt t/s	generate tok/s	vs empty
0	3959.63	116.02	1.000
8,192	2898.48	107.06	0.923
32,768	1690.19	90.68	0.782
65,536	1079.92	75.02	0.647
120,000	669.25	58.14	0.501

Generation stays over 100 tokens per second to about 8k tokens of conversation and halves by 120k.

Step 9 — Serve Many Users at Once

For the concurrency sweep the server ran 64 slots of 1,024 tokens (-c 65536 --parallel 64), 2612 MiB on the card. sweep.py sent 1 to 64 simultaneous requests, each with a unique prompt, five repeats per level, median reported:

c=1      93.28 tok/s aggregate  TTFT    212.7 ms  TPOT   9.34 ms  per-stream 107.03
c=4     196.09 tok/s aggregate  TTFT    669.1 ms  TPOT  15.64 ms  per-stream  63.95
c=16    316.07 tok/s aggregate  TTFT   2702.4 ms  TPOT  30.12 ms  per-stream  33.20
c=64    358.05 tok/s aggregate  TTFT   8516.4 ms  TPOT 113.20 ms  per-stream   8.84

That is 512 tokens in and 128 out. With 128 in and 512 out, chat-shaped traffic:

c=1     104.74 tok/s aggregate  TTFT    134.9 ms  TPOT   9.34 ms  per-stream 107.02
c=4     259.73 tok/s aggregate  TTFT    330.3 ms  TPOT  14.74 ms  per-stream  67.83
c=16    505.42 tok/s aggregate  TTFT   1114.3 ms  TPOT  29.32 ms  per-stream  34.10
c=64    618.12 tok/s aggregate  TTFT   3890.0 ms  TPOT  95.33 ms  per-stream  10.49

3.84x and 5.90x the single-user output at 64 users. Past 16 users the gain flattens while each user's speed keeps falling, so 16 is the practical ceiling: 34 tokens per second each, faster than anyone reads.

llama.cpp still reads prompts one request at a time, so time to first token about doubles with each doubling of users. The 1650 Ti ran into that wall: at 512 in and 128 out its output reached only 1.47x single-user. This card reads prompts fast enough that batched generation shows through anyway.

The card drew at most 110.93 W of its 120 W limit, held 1890 MHz in every busy sample and peaked at 68 °C. The server's prompt cache answered 0 tokens, so every prompt was read in full.

Step 10 — Bring Up E4B on the Same Card

local-llamacpp-3050oem-4b-q4_0 is the E4B sister rig: same card, binary and flags, with the exact Q4_0 E4B file (SHA-256 5462cc10…, 4.20 GB).

at 131,072 context	E2B	E4B
Process on the card	3610 MiB	6430 MiB
KV cache	780 MiB	2,088 MiB
Generate, empty context	116.02 tok/s	60.41 tok/s
Prompt, empty context	3959.63 t/s	2230.50 t/s
Generate, 120k in context	58.14 tok/s	30.04 tok/s

E4B generates at 0.52x E2B at every conversation length and fits with 1.7 GB to spare. The two rigs share port 8080 and do not fit on the card together. On the 1650 Ti, E4B fits at 8,192 context with the per-layer table in host memory and generates at about 39.8 tokens per second; this card runs it 1.52x faster with 16x the context.

Step 11 — Measure Answer Quality

Both files served one at a time with identical flags, no drafter, and three task sets: a 13-subset classification suite (3,880 questions), GSM8K math (1,319 problems) and BFCL simple tool calling (400 records), all greedy.

task	E2B	E4B	E4B − E2B	95% range
Classification suite	0.683	0.730 🥇	+0.048	+0.035 to +0.061
GSM8K	0.898	0.933 🥇	+0.035	+0.020 to +0.050
BFCL simple	0.925	0.910	−0.015	−0.043 to +0.013

E4B is better at reading comprehension and math, and level on tool calling, where the range includes zero. Both 4-bit files score within the error ranges of their bf16 references measured on Cloud TPU: E2B suite 0.683 against bf16's 0.683, E4B suite level with its bf16 run. The quantization costs nothing measurable; the gap between the two is the model.

The whole suite took 6.4 minutes for E2B and 11.2 for E4B on this card.

🔎 Tip: Spend the Spare Memory on Context

On an 8 GB card E2B leaves more than half the memory free even at full context. Spend it on the 131k context and on the per-layer table, both measured gains or free. Quantizing the KV cache to q8_0 saves memory the card does not need and costs 7.4% of generation.

🔎 Tip: Size `--parallel` Against `-c`

llama.cpp splits --ctx-size across --parallel slots. -c 131072 --parallel 16 gives each user 8,192 tokens; -c 16384 --parallel 16 gives each 1,024. Set both together, from what each user needs.

Compare and Contrast: Local Hardware

Gemma 4 E2B, the same exact Q4_0 file on every row except the last two, single user, llama-bench generation speed:

Machine	Device	Generate tok/s	vs 3050
Desktop, RTX 3050 OEM 8 GB	GPU, 120 W limit	116.14 🥇	1.00x
Lenovo Yoga 9, GTX 1650 Ti Max-Q 4 GB	GPU, 40 W limit	81.75 🥈	0.70x
Lenovo Yoga 9, i7-10750H	CPU, 6 threads	18.84 🥉	0.16x
Samsung Galaxy Book Go, Snapdragon 7c Gen 2	CPU, 2 cores	8.33	0.07x
Lenovo Yoga C630, Snapdragon 850	CPU, 4 cores	5.00	0.04x

Capability	RTX 3050 OEM	GTX 1650 Ti
Prompt read, 512 tokens	3,860 t/s 🥇	about 345 t/s
Full 131k context, E2B	🟢 3.6 GB of 8	⚠️ not measured
E4B	🟢 131k context, 60 tok/s	⚠️ 8k context, 40 tok/s
512 in / 128 out, best aggregate	🟢 358 tok/s, 3.84x	❌ 1.47x single user
Power under load	110.93 W peak	40 W cap

Compare and Contrast: Cloud

The two cheapest cloud deployments of Gemma 4 E2B measured in this repository serve bf16 weights through vLLM: an AWS g4dn.xlarge with a T4, and one Google Cloud TPU v5e chip.

Operating point	RTX 3050 OEM	T4, `g4dn.xlarge`	TPU v5e-1
One user, output tok/s	93.28 🥈	42.36 🥉	120 🥇
8 users, 512 in / 128 out	224.18	242.47	—
Best aggregate measured	618.12 🥈	243.93 🥉	1,382 🥇
Price	$179.97 once	$0.526/hour	$0.5779/hour, spot

The one-user row compares like shapes for the 3050 and T4 (512 in, 128 out); the TPU figure is 128 in, 128 out. The TPU is faster at every point and serves 2.2x the 3050's best aggregate. The 3050 generates 2.2x the T4's speed for one user and lands within 8% of it at 8 users.

And Price/Performance?

All of this is arithmetic from the measured speeds, the cloud rates recorded in those runs, a used-card listing captured on 2026-10-14, and the US Energy Information Administration's 2026 residential electricity forecast of 18.2 cents per kWh.

The card pays for itself in two to three weeks of continuous rental.

Against	Rate	Hours to $179.97	Days
TPU v5e-1, spot	$0.5779/h	311	13.0
T4 `g4dn.xlarge`, on-demand	$0.526/h	342	14.3
T4 `g4dn.xlarge`, spot	$0.3559/h	506	21.1

Running it costs about two cents an hour. At 110.93 W, the most the card drew in any sample, electricity is $0.0202 per hour against $0.526 for the T4.

Per million output tokens	Cost
RTX 3050, 64 users, 128 in / 512 out, electricity at peak draw	$0.009 🥇
RTX 3050, 8 users, 512 in / 128 out, electricity at peak draw	$0.025 🥈
RTX 3050, one user, electricity at 80 W	$0.039 🥉
TPU v5e-1, 64 users, spot	$0.116
T4, 8 users, on-demand	$0.603
TPU v5e-1, one user, spot	$1.338
T4, one user, on-demand	$3.449

The local rows leave out the card's purchase and the host it sits in; the cloud rows leave out storage and data transfer. A card already paid for wins by an order of magnitude on running cost. The TPU is the faster machine and the better buy only for a burst of work too big for one 3050, or for anyone without a desktop to put the card in.

So, Which One?

For a single person or a small team running Gemma 4 E2B or E4B all day, the used RTX 3050 OEM. It is the only local card here with tensor cores and 8 GB, and both together are what it takes: E4B at full context fits, prompts read at thousands of tokens per second, and 16 users each still get 34 tokens per second.

The 1650 Ti laptop remains the right answer for one user who already owns it. The ARM laptops serve E2B, slowly. Rent the TPU for bursts of work.

Between the two models, E4B buys 4.8 points on the classification suite and 3.5 on math for half the generation speed. For tool calling, E2B is level and twice as fast.

Summary

The goal of this article was to bring Gemma 4 E2B and E4B up on a used RTX 3050 OEM and to measure what it delivers for the money against the other local and cloud deployments in this repository. The key to the solution was the exact Q4_0 QAT files, which leave enough of the 8 GB card free for the full context and for many users at once. The measured results were:

🟢 116.14 tok/s generated and 3,860 t/s of prompt read for one user, E2B, 1.42x and 11x the 1650 Ti.
🟢 The full 131,072-token context in 3610 MiB for E2B and 6430 MiB for E4B, with no cost to short chats.
🟢 618.12 tok/s aggregate at 64 users (128 in / 512 out) and 358.05 at 512 in / 128 out; 34 tok/s per user at 16.
🟢 4-bit quality matches bf16: E2B suite 0.683 against 0.683; E4B 0.730.
🟢 Pays for itself in 311 to 506 hours against the cheapest cloud options, and runs for about $0.02 an hour in electricity (arithmetic).
⚠️ E4B generates at 0.52x E2B for +4.8 suite and +3.5 GSM8K points.
⚠️ Long conversations halve generation speed by 120k tokens of context.
❌ A q8_0 KV cache and flash attention off are both slower on this card.

Scope: one RTX 3050 OEM in an i7-5820K desktop, llama.cpp 77c4540ce with CUDA 13.4, measured 2026-10-13, without speculative decoding. The 1650 Ti, i7-10750H and Snapdragon figures come from other hosts, llama.cpp commits and days with the same model file; the T4 and TPU figures come from vLLM serving bf16 weights under a different benchmark tool; so every cross-machine ratio is indicative. The card price is one used listing's asking price on 2026-10-14, and the electricity figures use the highest power draw sampled or the single-user benchmark's end-of-run reading.

The strategy for running Gemma 4 locally on a used RTX 3050 OEM was validated with an incremental step by step approach.

References

One MCP Server for Six AWS Inference Backends: Bedrock, SageMaker, GPUs, Inferentia and Trainium

xbill — Tue, 13 Oct 2026 03:04:09 +0000

This article provides a step by step guide to the MCP server that operates a six-backend inference demo on AWS: Amazon Bedrock, a SageMaker endpoint, vLLM on EC2 g6 and g5g, and Gemma 4 on Inferentia2 and Trainium. A suite of eight Python MCP tools lists, launches, health-checks, reads the logs of, asks, and tears down every backend from one chat session.

https://github.com/xbill9/aws-inference-strands

What Is the Problem?

The demo behind this server puts one Strands agent in front of six ways to serve a model on AWS. Running it means five billable resources in up to four regions, two of them on AWS Neuron chips, and none of the instances reachable from the internet.

The repository already has shell scripts for all of it: stage/up.sh, stage/forward.sh, stage/check.py, stage/ask.py and stage/down.sh. Each one prints for a person at a terminal.

The MCP server gives the same operations to an agent. It answers "what is running and what is billing", "is inf2 serving yet", "show me the vLLM log on g6" and "tear it all down" with tools whose arithmetic and checks run in code, and it serves demo 1's counting tool to a Strands agent over the protocol itself.

At This Point You Should Have…

An AWS account with the AWS CLI configured, and Bedrock model access for Amazon Nova Micro in us-east-1
Python 3 (3.14.7 here) with mcp, strands-agents[openai,sagemaker] and boto3 installed
The instance profile and SSM-only security group stage/up.sh uses (it creates the security group on first use)
Claude Code, or any MCP client that can start a stdio server
The companion article for how each backend is built

Step 1 — Register the Server

The server is one file, mcp_server.py, started over stdio. .mcp.json in the repository root registers it for Claude Code:

{
  "mcpServers": {
    "strands-demo": {
      "command": "python3",
      "args": ["mcp_server.py"]
    }
  }
}

Any MCP client can start it the same way. Three lines of JSON-RPC on stdin are enough to see what it offers:

printf '%s\n' \
  '{"jsonrpc":"2.0","id":1,"method":"initialize","params":{"protocolVersion":"2025-11-25","capabilities":{},"clientInfo":{"name":"printf","version":"0"}}}' \
  '{"jsonrpc":"2.0","method":"notifications/initialized"}' \
  '{"jsonrpc":"2.0","id":2,"method":"tools/list"}' \
  | python3 mcp_server.py

{
  "protocolVersion": "2025-11-25",
  "capabilities": {"tools": {"listChanged": false}, "prompts": {"listChanged": false},
                   "resources": {"listChanged": false, "subscribe": false}, "experimental": {}},
  "serverInfo": {"name": "strands-demo", "version": ""}
}

The server runs on the MCP Python SDK 2.1.1 and its MCPServer class, and negotiates protocol version 2025-11-25. tools/list returns eight tools.

Step 2 — One File Over the Demo's Own Code

The server imports the two modules the demos already use and adds AWS calls around them.

           Claude Code  /  Strands MCPClient
                     │  stdio, JSON-RPC 2.0
                     ▼
 ┌──────────────────────── mcp_server.py ────────────────────────┐
 │ list_backends  backend_health  backend_logs   launch_log      │
 │ ask_backend    count_ids       launch_backends  teardown      │
 └──┬───────────────┬───────────────┬───────────────┬────────────┘
    │ boto3         │ backends.py   │ demo1_        │ subprocess
    │ EC2, SSM,     │ chat_once()   │ counting.py   │ stage/up.sh
    │ SageMaker     │               │ count()       │ stage/down.sh
    ▼               ▼               ▼               ▼
    4 US regions,   Bedrock,        the 11 ids,     detached,
    by tag          SageMaker,      in memory       output in
                    SSM forwards                    stage/*.log

Four design rules hold the file together:

Reuse, then wrap. ask_backend calls backends.chat_once, the function stage/ask.py, stage/check.py, the benchmark and demo 2 call. count_ids calls demo1_counting.count, the function the in-process tool calls. A number from the server and a number from the scripts come from the same line of code.
Launch and teardown run the stage scripts. launch_backends starts stage/up.sh and teardown starts stage/down.sh, so the MCP route and the documented route create and destroy the same resources.
No inbound ports. Every instance sits in a security group with no inbound rules. Tools that need the box run a shell command on it through SSM Run Command.
The tag is the inventory. stage/up.sh tags every instance ManagedBy=strands-demo-run, and every lookup in the server filters on that tag across us-east-2, us-east-1, us-west-2 and us-west-1. A backend that fell over to another region on a capacity error is still found.

boto3 is synchronous, so each call runs in a worker thread through asyncio.to_thread. That lets one tool ask four regions at once with asyncio.gather:

async def _instances() -> list[dict]:
    per_region = await asyncio.gather(*[_region_instances(r) for r in REGIONS])
    return [i for found in per_region for i in found]

Step 3 — Read-Only, Write and Destructive

Every tool carries MCP tool annotations, so a client can tell an inventory call from one that bills or one that deletes before it calls anything.

READ_ONLY = ToolAnnotations(readOnlyHint=True, idempotentHint=True)
WRITE = ToolAnnotations(destructiveHint=False)
DESTRUCTIVE = ToolAnnotations(destructiveHint=True)

What tools/list returns, read back through the MCP client:

Tool	Annotation	Arguments	What it touches
`list_backends`	read-only	none	EC2, SSM and SageMaker describe calls in 4 regions
`backend_health`	read-only	`backend`	one chat request, on the box via SSM for EC2
`backend_logs`	read-only	`backend`, `tail`	the serving log, via SSM
`ask_backend`	read-only	`backend`, `prompt`, `max_tokens`	one chat request through the forwards
`count_ids`	read-only	`op`, `value`	demo 1's ids, in memory
`launch_log`	read-only	`which`, `tail`	`stage/up.log` or `stage/down.log`
`launch_backends`	write	`names`	`stage/up.sh`; starts billing
`teardown`	destructive	none	`stage/down.sh`; terminates everything

The rest of this article takes the tools in the order a demo day uses them.

Step 4 — `list_backends`: What Is Running and What Is Billing

list_backends searches the four regions for tagged instances, asks SSM whether each running one has checked in, asks SageMaker for the endpoint's status, and adds up what is billing. The counts are computed in the tool, so the agent quotes them.

With nothing launched:

{
  "backends": [
    {"backend": "bedrock", "where": "us-east-1 (managed)", "state": "always on", "ssm": "-"},
    {"backend": "sagemaker", "where": "us-east-2 endpoint gemma-4-e2b-strands", "state": "not found", "ssm": "-"},
    {"backend": "g6", "where": "-", "state": "not launched", "ssm": "-"},
    {"backend": "g5g", "where": "-", "state": "not launched", "ssm": "-"},
    {"backend": "inf2", "where": "-", "state": "not launched", "ssm": "-"},
    {"backend": "trn1", "where": "-", "state": "not launched", "ssm": "-"}
  ],
  "billing_now": 0,
  "ec2_running": 0,
  "sagemaker": "not found",
  "ask_backend_configured": ["bedrock", "g6", "g5g", "sagemaker", "inf2", "trn1"]
}

One call took 0.78 seconds for five AWS describe calls: one per region for EC2 and one for SageMaker. Each running instance adds an SSM call.

ask_backend_configured lists all six backends while five are down. It names every backend with a URL in demo.env, the file stage/forward.sh writes into the stage directory, and that file outlives the forwards. The tool's description says so, and Step 8 shows what a call to one of those backends returns.

Step 5 — `launch_backends` and `launch_log`: Start Billing, Then Read the Launch

launch_backends takes any of g6, inf2, trn1, g5g and sm, checks the names, and starts stage/up.sh with them in the background:

launch_backends {"names": ["g6", "inf2"]}
started stage/up.sh g6 inf2 (pid 2362067); output in stage/up.log, read it with launch_log

The call returns at once. The script picks an AMI, tries each zone in turn, then the next region, and writes what it did to stage/up.log. launch_log reads the log:

strands-demo-g6: launched i-005ebaee08f191e7a   us-east-2a
strands-demo-inf2: no capacity in us-east-2 (subnet-0b0f0d29473b0a624), trying the next zone
strands-demo-inf2: launched i-0bf43d9b77b321ec3 us-east-2b

Next: stage/forward.sh (once the instances are up), then: source stage/demo.env && python3 stage/check.py

inf2 had no capacity in the first zone and launched in the second. Half a minute later list_backends showed both instances running and billing, before either had registered with SSM:

{"backend": "inf2", "where": "us-east-2 i-0bf43d9b77b321ec3 inf2.xlarge", "state": "running", "ssm": "not registered"},
{"backend": "g6",   "where": "us-east-2 i-005ebaee08f191e7a g6.xlarge",   "state": "running", "ssm": "not registered"},
...
"billing_now": 2,
"ec2_running": 2

Step 6 — `backend_health`: Ask on the Box

For the four EC2 backends, backend_health finds the instance by tag and sends one shell command through SSM Run Command. The command runs curl against the server's port on the instance and prints the HTTP code, the time and the body:

cmd = (f"curl -s -m 90 -o /tmp/strands-health.json -w '%{{http_code}} %{{time_total}}' "
       f"-H 'content-type: application/json' -d {shlex.quote(body)} "
       f"http://localhost:{box['port']}/v1/chat/completions; echo; cat /tmp/strands-health.json")

The request never leaves the instance, so the check needs no port forward and no inbound rule. Bedrock and SageMaker are called directly. SageMaker is checked for InService first.

Polled once a minute from launch, the replies trace each instance coming up:

02:13:14Z g6   -> ❌ AWS InvalidInstanceId: Instances not in a valid state for account
02:14:21Z g6   -> ❌ RuntimeError: SSM Failed: 000 0.000140
...
02:24:36Z g6   -> ✅ g6: answered 'Paris' in 1.21s, on i-005ebaee08f191e7a (us-east-2) via SSM
02:26:11Z inf2 -> ❌ RuntimeError: SSM Failed: 000 90.000889
...
02:32:44Z inf2 -> ✅ inf2: answered 'Paris' in 0.27s, on i-0bf43d9b77b321ec3 (us-east-2) via SSM

Each failure names a different stage:

Reply	Meaning
`InvalidInstanceId`	the SSM agent has not registered yet; nothing can run on the box
`000` in under a millisecond	the box answers, the serving port is closed
`000` after 90.0 s	the port is open and the request waited out `curl -m 90`
✅ `'Paris'`	the model is serving and answers correctly

The 90-second replies came from inf2 only. Its server opens port 8080 while it still loads the compiled model onto the NeuronCores, and holds requests until it is ready.

g6 first passed 12.2 minutes after the launch call and inf2 20.4 minutes after it, with the checks a minute apart. g6's first answer took 1.21 seconds and the next three took 0.02 to 0.03 seconds.

Step 7 — `backend_logs`: The Serving Log Without a Shell

backend_logs runs the right log command for each backend through SSM: docker logs for the vLLM container on g6 and the Neuron container on inf2 and trn1, journalctl -u vllm.service on g5g. Until the container exists it falls back to the cloud-init log, which shows the image pull:

g6 i-005ebaee08f191e7a (us-east-2)
Error response from daemon: No such container: vllm
--- serving log not available yet; cloud-init:
...
469e56fbf4cf: Pull complete

Two minutes after g6's first health check passed, the same call returned vLLM's own log, with the health check's request in it:

(APIServer pid=1) INFO:     Application startup complete.
(EngineCore pid=108) WARNING 10-13 02:24:34 [jit_monitor.py:141] Triton kernel JIT compilation during inference: kernel_unified_attention. This causes a latency spike; consider extending warmup to cover this shape/config.
(APIServer pid=1) INFO:     172.17.0.1:34422 - "POST /v1/chat/completions HTTP/1.1" 200 OK

The JIT warning matches the 1.21-second first answer. On inf2 the log showed the load in progress, then the result:

HTTP listening on :8080 — /health=503 'loading' until ready; queue max=8, gen timeout 120.0s
...
READY in 473.8s — serving on :8080 (SLIM host, peak RSS 14.45 GB)
[req] pt=21 ct=1 5.8tok/s 0.17s finish=stop

The first line explains the 90-second health check replies: the port opened before the model was ready.

Step 8 — `ask_backend`: One Timed Request

ask_backend sends one chat request through backends.chat_once and returns the text, the backend's own completion token count, and the time measured in the server. Tokens per second is that count over that time, so it includes the network and the prompt processing. Self-hosted backends are reached through the SSM port forwards stage/forward.sh opens.

Before the forwards, with demo.env left over from an earlier session:

ask_backend {"backend": "g6", "prompt": "hello"}
❌ URLError: <urlopen error [Errno 111] Connection refused>

After stage/forward.sh rewrote demo.env, a new server process read it at start-up:

{"backend": "g6",   "label": "vLLM on EC2 g6 (L4)",            "text": "Hello! How can I help you today? 😊",
 "completion_tokens": 11, "seconds": 0.14,  "tokens_per_second": 78.6}
{"backend": "inf2", "label": "Inferentia2 (inf2, hand-ported)", "text": "Hello! How can I help you today?",
 "completion_tokens": 9,  "seconds": 0.423, "tokens_per_second": 21.3}

ask_backend {"backend": "trn1", "prompt": "hello"}
❌ trn1 is not configured here; configured: bedrock, g6, inf2. Run stage/forward.sh.

list_backends now reported ask_backend_configured: ["bedrock", "g6", "inf2"] and both instances Online in SSM. Replies this short are mostly prompt processing and network, so these rates describe one request each; the companion article has the 100-token benchmark.

Step 9 — `count_ids`: Demo 1's Engine Over MCP

Demo 1 asks how many of eleven ids are 10 or more. The answer is 8. count_ids takes a predicate and returns the exact count, minimum and maximum, and echoes the predicate it ran:

count_ids {"op": ">=", "value": 10}  ->  {"where": "id >= 10", "count": 8, "min": 10, "max": 23}
count_ids {"op": ">",  "value": 10}  ->  {"where": "id > 10",  "count": 7, "min": 11, "max": 23}
count_ids {"op": "=>", "value": 10}  ->  {"error": "op must be one of ['<', '<=', '==', '>', '>=']"}

The second line returns an exact, sourced 7 for a predicate the question did not ask for. The where field is in the response so the predicate can be checked along with the number.

demo1_counting.py --tools mcp offers this tool to a Strands agent through the protocol. Strands' MCPClient starts the server over stdio and a tool filter exposes count_ids alone:

MCPClient(lambda: stdio_client(StdioServerParameters(command=sys.executable, args=[server])),
          tool_filters={"allowed": ["count_ids"]})

The tool runs in the server's process, so the demo cannot record calls as they happen. It reads them back from the agent's conversation instead, and grades them with the same function the in-process run uses:

def calls_from(messages: list) -> list:
    return [{"tool": c["toolUse"]["name"], **c["toolUse"].get("input", {})}
            for m in messages if m["role"] == "assistant" for c in m["content"] if "toolUse" in c]

Twenty runs each on Bedrock, the in-process tool and then the MCP tool:

BACKENDS=bedrock python3 demo1_counting.py --tools engine --runs 20
BACKENDS=bedrock python3 demo1_counting.py --tools mcp --runs 20

True answer: 8  tools: engine
...
backend     runs  correct  right-filter  errors  median s  filters sent
bedrock       20   20/20          20/20       0      1.07  id >= 10

True answer: 8  tools: mcp
...
backend     runs  correct  right-filter  errors  median s  filters sent
bedrock       20   20/20          20/20       0      1.07  id >= 10

Same answers, same filter, same median time. Starting the server costs once per demo run, outside the timed agent calls.

Step 10 — `teardown`: Stop Everything That Bills

teardown takes no arguments and starts stage/down.sh in the background. It is the one tool marked destructive. The script closes the forwards, terminates every instance tagged ManagedBy=strands-demo-run in all four regions, deletes the SageMaker endpoint when it exists, and then lists anything still running in any of the four regions.

teardown {}
started stage/down.sh  (pid 2367066); output in stage/down.log, read it with launch_log

Twenty-one seconds later, launch_log with which set to down:

us-east-2: terminating i-0bf43d9b77b321ec3  i-005ebaee08f191e7a

Sweep (anything listed here is still billing):
(sweep done)

and list_backends:

"billing_now": 0,
"ec2_running": 0,
"sagemaker": "not found"

The sweep and the inventory both read empty while the two instances were still in shutting-down, because both look only at pending, running, stopping and stopped instances.

🔎 Tip: Return the Error as Text the Agent Can Read

Every tool catches its exceptions and returns one line that starts with ❌ and names the cause: an AWS error code and message, a ValueError for a bad backend name, a RuntimeError naming the four regions searched.

backend_health {"backend": "tpu"}  ->  ❌ ValueError: backend must be one of bedrock, sagemaker, g6, g5g, inf2, trn1
launch_backends {"names": ["p5"]}  ->  ❌ unknown backend(s) ['p5']; use g6 inf2 trn1 g5g sm

An agent reads the cause in the reply. The protocol-level isError flag stays false on all of them, so a client that branches on that flag treats a failed health check as a success. Set it where clients act on it.

🔎 Tip: Pass a Health Check on the Answer

backend_health asks "What is the capital of France? Answer in one word." at temperature 0 with 16 tokens, and passes only when the reply contains "paris". A non-empty reply proves the port is open. A broken build of this model family has answered with repeated text, and that reply is non-empty too.

🔎 Tip: Clamp Every Number the Model Sends

tail is clamped to 1 to 2000 lines and max_tokens to 1 to 1024. backend is checked against a fixed list before any AWS call, and launch_backends rejects an unknown name before stage/up.sh starts. A model that asks for a million log lines gets two thousand.

🔎 Tip: Run Long Jobs Detached and Hand Back a Log

stage/up.sh and stage/down.sh each took under half a minute here, and the backends behind them take up to twenty minutes to serve. Both scripts run with start_new_session=True and their output in stage/up.log or stage/down.log, so the tool call returns at once and the job survives the MCP session closing. launch_log reads the log back.

Compare and Contrast

	stage scripts	MCP server
Caller	a person at a terminal	any MCP client: Claude Code, a Strands agent
Launch and teardown	`stage/up.sh`, `stage/down.sh`	the same scripts, detached, read back with `launch_log`
Inventory	`aws ec2 describe-instances` per region	🥇 one call, four regions in parallel, billing count computed
Health	`stage/check.py`, needs the forwards	🥇 on the box via SSM, no forwards
Logs	`aws ssm start-session`, then `docker logs`	🥇 `backend_logs`, cloud-init fallback
One question	`stage/ask.py`, needs the forwards	`ask_backend`, needs the forwards
Safety signal	the script's name	read-only, write, destructive annotations

So, Which One?

🥇 The MCP server to operate the demo from a chat session: inventory, health, logs and teardown in plain questions, with the billing count and the pass or fail computed in code.

🥈 The stage scripts for the talk itself and anything scripted, and for the port forwards, which ask_backend and the demos need and the server does not open.

Both launch and stop the same resources, because one calls the other.

Summary

The goal of this article was to operate a six-backend AWS inference demo through MCP. The key to the solution was one server over the demo's own code, reaching instances through SSM Run Command and launching through the same scripts a person runs. The results were:

🟢 Eight tools in one file over the demo's own code: the same chat_once and count functions the scripts call, and the same stage/up.sh and stage/down.sh.
🟢 Every tool is marked read-only, write or destructive, and tools/list carries the marks to the client.
🟢 Health checks and logs run on the instance through SSM Run Command, with no inbound rules and no port forwards; g6 passed 12.2 minutes after launch and inf2 20.4 minutes after.
🟢 A launch, health checks, logs, requests and a teardown ran from one session, and list_backends read billing_now: 0 21 seconds after teardown.
🟢 Demo 1 over MCP matched demo 1 in-process on Bedrock: 20 of 20 correct with the right filter both ways, median 1.07 seconds both ways.
⚠️ Failed calls return a ❌ line with isError false, so a client that branches on the protocol flag treats them as successes.
⚠️ ask_backend depends on demo.env, which outlives the forwards; it lists backends that are down until stage/forward.sh runs again.
❌ The server opens no port forwards, so ask_backend and the demos still need stage/forward.sh from a terminal.

Scope: one session on 2026-10-13 (UTC), us-east-2, two on-demand instances (g6.xlarge and inf2.xlarge) from launch to teardown, with Bedrock in us-east-1. trn1, g5g and the SageMaker endpoint were not launched, so their tools are shown only in their not-running replies. Health checks were one minute apart, so each time to ready is accurate to about a minute. Demo 1 ran 20 times per tool set on Bedrock only, with the MCP Python SDK 2.1.1, Strands Agents 1.55.0 and Python 3.14.7 on a laptop.

The strategy for using MCP to operate Gemma 4 inference backends across AWS was validated with an incremental step by step approach.

References

EmbeddingGemma 2 on Cloud Run: an int8 CPU Service Against an NVIDIA L4

xbill — Tue, 13 Oct 2026 01:20:25 +0000

This article provides a step by step deployment guide for EmbeddingGemma 2 text embeddings to Google Cloud Run, on a CPU-only instance and on an NVIDIA L4 GPU. Each service is tested, used in a small semantic search demo, benchmarked and priced per million queries and per million documents.

https://github.com/xbill9/embedding-cloudrun

What is this project trying to Do?

This project serves google/embeddinggemma-2 behind one HTTP endpoint on Cloud Run and answers a cost question: what does it take to embed text with it, on a 1 vCPU instance or on an L4?

Cloud Run runs a container image as an HTTPS service. It starts instances when requests arrive, stops them when traffic stops, and needs no cluster, VM or load balancer. A CPU service can be billed per request: CPU and memory are charged only while a request is in flight. A GPU service is billed per instance: the L4 is charged for every second the instance is up.

An L4 instance costs about ten times as much per hour as a 1 vCPU instance. So the comparison turns on two numbers: how much faster the L4 embeds, and how much of its time it spends busy.

Why Run Your Own Embeddings?

Hosted decision APIs now sell work that embedding models have done for years. TypeSafe's Jev, launched on September 15, 2026, answers typed questions about a piece of text: pick one label from a list, place the text on a scale, or return the probability that a statement is true. It is priced at $0.042 per million input tokens, with access through a waitlist, OpenRouter and Vercel's AI Gateway.

Ranking passages against a query and labelling text by its closest description are what an embedding model is built for. EmbeddingGemma 2 is an open model, so it can run in your own Google Cloud project, at Cloud Run's rates, with the text never leaving the project.

This article prices that option per million input tokens, with the same unit Jev uses, and measures its speed and what it takes to deploy. Accuracy on the same items as Jev is not measured here; that head-to-head is the subject of a follow-up.

What is EmbeddingGemma 2?

EmbeddingGemma 2 turns text into a 768-number vector, so that texts with similar meaning land close together. It also carries vision and audio encoders, which this project drops: only the text towers are loaded.

The vectors are Matryoshka vectors. The first 128, 256 or 512 numbers of the 768 are themselves usable embeddings, so a search index can trade a little quality for a smaller footprint. The model also takes a prompt name: SearchQuery for a question and Document for a passage to be searched.

Three Services, One Dockerfile

service	weights	instance	billing
`emb-cpu-repack`	text only, int8 Linears, int4 token table, 203 MiB	1 vCPU, 2 GiB	per request
`emb-cpu-original`	text only, float32	1 vCPU, 4 GiB	per request
`emb-l4-original`	text only, bf16	4 vCPU, 16 GiB, 1 L4	per instance

The repack is built by repack/repack_text_w8a8_emb4.py in the repository: every Linear in the text model stored as int8 with one scale per output channel, and the 262,144-row token table as int4. The int8 Linears run on PyTorch's fbgemm kernels, which are CPU kernels, so the L4 serves the original weights in bf16.

4 vCPU and 16 GiB is the smallest instance Cloud Run attaches an L4 to.

Where do I start?

The strategy for starting the deployment is an incremental step by step approach.

First, the model files are staged. Then each variant is built into an image, deployed, tested against its own contract and against the original weights, used for a search demo, timed, and priced from the Cloud Billing catalog.

At This Point You Should Have…

A Google Cloud project with billing, and the gcloud CLI signed in
The Cloud Run, Artifact Registry and Cloud Build APIs enabled
Docker, for the CPU images
Python 3.12 or newer with torch, transformers and sentence-transformers for the local check and the repack
For the L4: Cloud Run GPU quota, without zonal redundancy, in the region you deploy to

Setup the Basic Environment

git clone https://github.com/xbill9/embedding-cloudrun
cd embedding-cloudrun
python3 -m pip install torch==2.14.1 transformers==5.19.0 sentence-transformers==6.1.0 huggingface_hub

The scripts that talk to a deployed service (test_service.py, demo.py, bench.py, cost.py) use only the Python standard library.

Step 1 — Stage the Models

scripts/prepare_models.sh

The script downloads a snapshot of google/embeddinggemma-2 into models/original, then builds the repack from it into models/repack:

dropped 473.4M tower parameters; checkpoint 203.4 MiB

1.5G    models/original
235M    models/repack

The repack takes about 8 seconds on a laptop CPU and is deterministic: a second build writes a model.safetensors that is byte for byte the same. models/ is gitignored; the Dockerfile copies one of the two directories into the image.

Step 2 — The Server

app/server.py is one FastAPI file with two routes. MODEL_VARIANT picks the weights, and the model loads before uvicorn binds the port, so Cloud Run's startup time includes it:

@app.post("/embed")
def embed(r: EmbedRequest):
    if r.dim not in (128, 256, 512, 768):
        return {"error": "dim must be 128, 256, 512 or 768"}
    t = time.perf_counter()
    e = model.encode(r.texts, prompt_name=r.prompt_name, truncate_dim=r.dim,
                     convert_to_tensor=True, batch_size=32)
    # Normalize in float32: a bf16 model normalizes in bf16, leaving norms up to 0.25% off 1.
    e = torch.nn.functional.normalize(e.float(), dim=-1)
    return {"embeddings": e.cpu().tolist(), "seconds": time.perf_counter() - t}

Every reply carries seconds, the server's own time spent embedding. The benchmark and the cost model use that figure, so network distance between the client and the region does not enter the price.

One request holds 1 to 256 texts. Vectors come back unit length, so a dot product is the cosine similarity.

🔎 Tip: Cloud Run Reserves Paths Ending in z

The conventional health route is /healthz. On Cloud Run a request for it never reaches the container: paths ending in z are reserved, and the front end answers 404. The health route here is /health.

Step 3 — Build and Deploy a CPU Service

scripts/deploy.sh cpu-repack

The script builds the image locally, pushes it to Artifact Registry and deploys:

gcloud run deploy emb-cpu-repack --image $IMAGE --region us-central1 \
  --allow-unauthenticated --min-instances 0 --max-instances 1 --cpu-boost \
  --cpu 1 --memory 2Gi --concurrency 1 --cpu-throttling \
  --set-env-vars TORCH_THREADS=1

Service [emb-cpu-original] revision [emb-cpu-original-00001-r96] has been deployed and is serving 100 percent of traffic.
Service URL: https://emb-cpu-original-1056842563084.us-central1.run.app

real    6m17.699s

The output above is the float32 service, scripts/deploy.sh cpu-original, from build to serving. The flags that matter:

--cpu-throttling selects request-based billing: CPU is allocated, and charged, only during a request.
--min-instances 0 lets the service scale to zero, so an idle service costs nothing.
--max-instances 1 caps a runaway client at one instance.
--concurrency 1 sends one request at a time to the single vCPU.
TORCH_THREADS=1 matters because the container sees every core on the host while Cloud Run limits it to one vCPU by quota. PyTorch would otherwise start a thread per host core.

Compressed in Artifact Registry, model included, the repack image is 636 MB and the float32 image 1270 MB.

Step 4 — Deploy the L4 Service

The CUDA image is 4092 MB compressed in Artifact Registry, so deploy.sh builds it on Cloud Build instead of on local disk:

SERVICE_REGION=us-east4 scripts/deploy.sh l4-original

gcloud run deploy emb-l4-original --image $IMAGE --region us-east4 \
  --allow-unauthenticated --min-instances 0 --max-instances 1 --cpu-boost \
  --cpu 4 --memory 16Gi --gpu 1 --gpu-type nvidia-l4 --no-gpu-zonal-redundancy \
  --no-cpu-throttling --concurrency 4 --set-env-vars TORCH_THREADS=4

Service [emb-l4-original] revision [emb-l4-original-00001-shk] has been deployed and is serving 100 percent of traffic.
Service URL: https://emb-l4-original-1056842563084.us-east4.run.app

real    2m30.320s

--no-gpu-zonal-redundancy gives up capacity reserved in a second zone for failover, at 0.0001867 USD per L4-second against 0.0002909 with it.

🔎 Tip: GPU Quota Is Per Region

The image lives in us-central1, and the same deploy there was refused:

ERROR: (gcloud.run.deploy) spec.template.metadata.annotations[autoscaling.knative.dev/maxScale]: You do not have quota for using GPUs without zonal redundancy. Learn more about GPU zonal redundancy: g.co/cloudrun/gpu-redundancy-help

The quota is set per region and per redundancy mode. List it before choosing a region:

gcloud beta quotas info describe NvidiaL4GpuAllocNoZonalRedundancyPerProjectRegion \
  --service=run.googleapis.com --format=json

{'region': 'europe-west1'} {'value': '3'}
{'region': 'europe-west4'} {'value': '3'}
{'region': 'us-central1'} {}
{'region': 'us-east4'} {'value': '3'}

A service in one region pulls an image from Artifact Registry in another, so only SERVICE_REGION changes. The Cloud Billing catalog lists the same Cloud Run CPU, memory and L4 rates for us-central1 and us-east4.

Step 5 — Test the Deployment

scripts/test_service.py checks one service against the contract the server promises: health, vector length and unit norm at every Matryoshka size, the 128-number vector being the first 128 of the 768, bad input rejected, and a relevant passage ranked above two unrelated ones.

python3 scripts/test_service.py https://emb-cpu-repack-fgasxpwzoq-uc.a.run.app

✅ health: {"variant": "repack", "device": "cpu", "threads": 1, "load_s": 9.03, "max_seq_length": 512}
✅ dim 128: length 128, norm 1.0000
✅ dim 256: length 256, norm 1.0000
✅ dim 512: length 512, norm 1.0000
✅ dim 768: length 768, norm 1.0000
✅ Matryoshka prefix: cosine(first 128 of 768, dim 128) = 1.00000
✅ rejects dim 100: {"error": "dim must be 128, 256, 512 or 768"}
✅ rejects empty texts: HTTP 422
✅ ranking: scores 0.776, 0.511, 0.642
9/9 checks passed

The ranking check asks "How do plants make food from sunlight?" against a passage on photosynthesis, one on the Treaty of Westphalia and one on sourdough. The float32 and L4 services pass the same nine checks:

✅ health: {"variant": "original", "device": "cuda:0", "threads": 4, "load_s": 6.87, "max_seq_length": 512}
✅ dim 128: length 128, norm 1.0000
...
9/9 checks passed

🔎 Tip: Normalize bf16 Embeddings in float32

Sentence Transformers' normalize_embeddings=True normalizes in the model's own precision. On the L4 the model runs in bf16, and vectors that should be unit length came back with norms of 0.9975 at 128 dimensions and 1.0015 at 512, so a dot product used as a cosine is off by the same amount. Loading the same weights in bf16 on a laptop CPU gives the same spread, 0.9977 to 1.0012.

The server asks for the raw vectors and normalizes them in float32:

e = model.encode(r.texts, prompt_name=r.prompt_name, truncate_dim=r.dim,
                 convert_to_tensor=True, batch_size=32)
e = torch.nn.functional.normalize(e.float(), dim=-1)

With that change every norm prints as 1.0000 on all three services.

Step 6 — Does the Repack Match the Original?

--against embeds the same four texts on a second service and compares the vectors:

python3 scripts/test_service.py $REPACK_URL --against $ORIGINAL_URL

✅ agrees with --against: cosine per text min 0.9888 max 0.9978
10/10 checks passed

The L4 in bf16 and the CPU in float32 run the same original weights, and agree almost exactly:

python3 scripts/test_service.py $L4_URL --against $ORIGINAL_URL

✅ agrees with --against: cosine per text min 0.9999 max 1.0000
10/10 checks passed

The repack against the L4 measured 0.9858 to 0.9981. The repack's vectors shift slightly from one Cloud Run host to the next, because its int8 kernels depend on the host CPU, so the same check gives 0.986 to 0.998 across runs.

A cosine near 0.99 says the vectors point the same way. Retrieval is the use that matters, so the repack was also scored on SciFact (300 queries over 5,183 scientific abstracts, 768 dimensions) with repack/eval_scifact.py:

weights	nDCG@10	recall@10	top 10 shared with float32
🥇 float32 original	0.854	0.944	100%
🥈 int4 token table only	0.850	0.939	92.7%
🥉 repack (int8 + int4)	0.836	0.926	85.5%

The repack keeps 97.9% of the original's nDCG@10.

Step 7 — A Simple Demo: Semantic Search

scripts/demo.py embeds the sixteen passages in docs/corpus.txt as documents and the question as a search query, then prints the closest three:

python3 scripts/demo.py $REPACK_URL "Why does bread rise?"

16 passages, dim 768, server 2.59 s, wall 2.99 s
0.853  Bread dough rises because yeast releases carbon dioxide that the gluten network traps.
0.698  A sourdough starter is a mixture of flour and water fermented by wild yeast and lactic acid bacteria.
0.638  The northern lights appear when charged particles from the sun excite gases in the upper atmosphere.

python3 scripts/demo.py $REPACK_URL "Which old document helped decode Egyptian writing?"

16 passages, dim 768, server 2.63 s, wall 3.02 s
0.745  The Rosetta Stone carries the same decree in hieroglyphic, Demotic and Ancient Greek scripts.
0.631  Matryoshka embeddings can be truncated to a shorter prefix and still be used for retrieval.
0.595  A hash map stores key-value pairs and finds a key in constant time on average.

The second question shares no keywords with its answer: "Rosetta Stone" and "hieroglyphic" appear only in the passage. The L4 returns the same ranking:

python3 scripts/demo.py $L4_URL "Which old document helped decode Egyptian writing?"

16 passages, dim 768, server 0.13 s, wall 0.42 s
0.749  The Rosetta Stone carries the same decree in hieroglyphic, Demotic and Ancient Greek scripts.
0.634  Matryoshka embeddings can be truncated to a shorter prefix and still be used for retrieval.
0.592  A hash map stores key-value pairs and finds a key in constant time on average.

All three services put the same passage first for all three demo questions.

Step 8 — Cold Start

With --min-instances 0, the first request after an idle period starts an instance: Cloud Run places it, pulls the image and starts the container, and the model loads before the port opens. scripts/coldstart.sh times that first request, then a second one, after 25 minutes with no traffic:

scripts/coldstart.sh $L4_URL

2026-10-12T17:46:53Z
{"variant":"original","device":"cuda:0","threads":4,"load_s":4.78,"max_seq_length":512}
http=200 total_s=17.271792
second http=200 total_s=0.098748

scripts/coldstart_summary.py splits each first request into the server's model load and the rest:

file                                   variant    first request s  model load s  the rest s  warm s
coldstart-cpu-original.txt             original              48.0          15.5        32.5    0.13
coldstart-cpu-repack-2.txt             repack                58.4           9.5        48.9    0.11
coldstart-cpu-repack.txt               repack                83.5          13.8        69.7    0.13
coldstart-l4-original.txt              original              17.3           4.8        12.5    0.10

The model load is the smaller part on every service. The CPU services took 48 to 84 seconds to answer a first request, and the L4 took 17, with the largest image of the three. Cloud Run's request log for the float32 cold start shows the order: instance started at 17:08:08, first line from the container at 17:08:49, startup check passed at 17:08:56.

A service that cannot wait that long can keep one instance warm with --min-instances 1; Cloud Run pricing lists how idle instances are billed.

Step 9 — Benchmark

scripts/bench.py sends 50 single-query requests, then 10 requests of 32 documents of 141 tokens each (scripts/count_tokens.py, prompt included), and records the server's own time for each:

python3 scripts/bench.py cpu-repack $REPACK_URL

Each run is kept in results/runs/. A 1 vCPU instance lands on whichever host Cloud Run picks, so each variant was run three times, each on a new revision and so a new instance. cost.py reports the median run and the range:

variant        runs          query ms            docs/s  USD / 1M queries       USD / 1M docs
cpu-original      3     251 (227-251)     0.9 (0.7-0.9)              8.93 39.05 (38.37-47.72)
cpu-repack        3      121 (88-128)     1.1 (1.0-1.6)              3.91 25.87 (18.17-27.69)
l4-original       3        56 (54-59) 333.1 (322.7-339.1)             16.27    0.87 (0.86-0.90)

scripts/ratios.py turns the medians into ratios:

L4 docs/s over cpu-repack docs/s: 297x
cpu-original query time over cpu-repack query time: 2.08x
cpu-repack docs/s over cpu-original docs/s: 1.29x
SciFact nDCG@10, repack over float32: 97.9%

On the same 1 vCPU, the repack answers a query in half the time of float32 and embeds documents 1.29x as fast, in half the memory. The L4 embeds 297 times as many documents per second as the repack.

Step 10 — Cost Analysis

scripts/fetch_prices.py reads the Cloud Run rates from the Cloud Billing Catalog API into prices.json, and scripts/cost.py does all the arithmetic from those rates and the benchmark runs:

python3 scripts/cost.py results/runs/*.json

prices fetched 2026-10-10, us-central1

variant        billing    USD/hour up
cpu-repack     request         0.1044
cpu-original   request         0.1224
l4-original    instance        1.0465

A per-request CPU service is charged its server seconds times its per-second rate, plus 0.40 USD per million requests after the first two million each month. An idle CPU service costs nothing. The L4 is charged for every second its instance is up, so its per-document figures in Step 9 hold only while it is kept busy.

And Price/Performance?

An instance-billed L4 that is busy part of the time costs more per document in proportion. cost.py prints the curve against the repack:

L4 cost per 1M docs when the instance is busy part of the time it is up (vs cpu-repack, median runs):
  busy  L4 USD / 1M docs   cpu-repack USD / 1M docs
  100%              0.87                      25.87
   50%              1.75                      25.87
   25%              3.49                      25.87
   10%              8.73                      25.87
    5%             17.46                      25.87
    1%             87.28                      25.87
break-even: the L4 is cheaper per document while busy more than 3.4% of the time it is up

Per million input tokens, the unit Jev is priced in:

USD per 1M input tokens, documents of 141 tokens (median runs):
cpu-original     0.2770
cpu-repack       0.1835
l4-original      0.0062  at 100% busy
l4-original      0.0619  at 10% busy
jev (hosted)     0.0420  list price
l4-original matches Jev's price per input token while busy 14.7% of the time it is up; at 100% busy Jev costs 6.8x as much
cpu-original costs 6.6x Jev's price per input token
cpu-repack costs 4.4x Jev's price per input token

A busy L4 embeds text for a sixth of Jev's list price per input token. The 1 vCPU repack costs 4.4 times Jev's price per token, and nothing at all while idle.

Compare and Contrast

	cpu-repack	cpu-original	l4-original
Query, median	🥈 121 ms	🥉 251 ms	🥇 56 ms
Documents per second	🥈 1.1	🥉 0.9	🥇 333
USD per 1M queries	🥇 3.91	🥈 8.93	🥉 16.27
USD per 1M documents, busy	🥈 25.87	🥉 39.05	🥇 0.87
Cost while idle	🥇 none	🥇 none	🥉 1.05 USD per hour up
Cold start	🥉 58–84 s	🥈 48 s	🥇 17 s
GPU quota needed	🥇 no	🥇 no	🥉 yes
SciFact nDCG@10	🥈 0.836	🥇 0.854	not run; cosine 0.9999 to float32

So, Which One?

For search queries arriving through the day, use the CPU repack. It is the cheapest per query, costs nothing between requests, needs no GPU quota, and keeps 97.9% of the original's retrieval score.

For indexing a corpus, use the L4. It embeds 297 times as fast and costs 0.87 USD per million documents while busy; it pays for itself over the repack once it is busy more than 3.4% of the time it is up. Deploy it for the job and let it scale to zero afterwards.

The float32 CPU service is the reference that the other two are checked against.

Step 11 — Teardown

scripts/down.sh

It deletes the three services in both regions. The images stay in Artifact Registry until the repository is deleted:

gcloud artifacts repositories delete embedding-cloudrun --location us-central1

Summary

The goal of this article was to deploy EmbeddingGemma 2 text embeddings to Cloud Run and price them against a hosted decision API. The key to the solution was one server and one Dockerfile for three services: an int8 repack on 1 vCPU with per-request billing, the float32 original on 1 vCPU as the reference, and the original in bf16 on an NVIDIA L4. The deployment results were:

🟢 All three services pass the same nine checks, and the demo puts the same passage first on all three
🟢 The repack matches float32 at a cosine of 0.986 to 0.998 per text and keeps 97.9% of its SciFact nDCG@10
🟢 The repack answers a query in a median 121 ms on 1 vCPU, at 3.91 USD per million queries and nothing while idle
🟢 The L4 embeds 333 documents per second, 0.87 USD per million documents while busy, 6.8 times below Jev's list price per input token
⚠️ The L4 bills for every second it is up and is cheaper per document than the repack only while busy more than 3.4% of that time
⚠️ Cold starts took 48 to 84 seconds on CPU and 17 on the L4
⚠️ bf16 vectors need float32 normalization to come back unit length
❌ No L4 quota without zonal redundancy in us-central1; the L4 ran in us-east4

Scope: one Google Cloud project, CPU services in us-central1 and the L4 in us-east4, one instance per service, Cloud Run prices from the Billing Catalog on 2026-10-10, measured 2026-10-11 and 2026-10-12. Three benchmark runs per variant, each on a new instance, with server-side timing, one document length (141 tokens) and one request at a time; the L4 was not tested with concurrent requests. Cold starts are single measurements. Jev's figure is its list price per input token; its accuracy against EmbeddingGemma 2 on the same items is not measured here.

The strategy for deploying EmbeddingGemma 2 on Cloud Run was validated with an incremental step by step approach.

References

Repository: https://github.com/xbill9/embedding-cloudrun
EmbeddingGemma 2 on Hugging Face: https://huggingface.co/google/embeddinggemma-2
Cloud Run GPU support: https://cloud.google.com/run/docs/configuring/services/gpu
Cloud Run billing settings: https://cloud.google.com/run/docs/configuring/billing-settings
Cloud Run pricing: https://cloud.google.com/run/pricing
Sentence Transformers: https://sbert.net/
SciFact in BEIR: https://huggingface.co/datasets/mteb/scifact
TypeSafe Jev: https://typesafe.ai
Jev after eight days of independent tests: https://dev.to/aws-builders/jev-after-eight-days-of-independent-tests-level-with-mid-price-llms-behind-the-frontier-1c60

Debian on a Samsung Galaxy Book Go, Written on the Galaxy Book Go

xbill — Tue, 13 Oct 2026 01:00:40 +0000

This article provides a step by step guide to running Debian sid on a Samsung Galaxy Book Go, the Snapdragon 7c Gen 2 Windows-on-ARM laptop. A small set of POSIX shell scripts and one device tree patch turn a community draft into a machine that runs a daily desktop.

https://github.com/xbill9/linux-arm-adventures

Kernel 7.2.9, GNOME 50.5 on Wayland, GPU acceleration, Wi-Fi, Bluetooth, the speakers, the camera, the battery gauge and charging, the backlight, the keyboard, touchpad, power button, lid and suspend all work. The hardware video decoder works in mpv, and browsers decode in software. This article was written, fact-checked and built on the laptop it describes.

What Is the Problem?

The Galaxy Book Go has no upstream support. Mainline Linux has no device tree for it, and Debian's kernel ships only the Acer Aspire 1 tree for the same SoC (SC7180). The community device tree is a draft with Aspire 1 memory addresses and a keyboard interrupt its own comment calls wrong.

ACPI is no way out. A generic ACPI kernel reboots the machine as soon as the SMMU driver loads, and without the SMMU there is no USB and no keyboard. So the laptop boots with a device tree, and every missing node in that tree is a missing device.

Samsung publishes no driver pack for it either. The only copy of the signed firmware is the factory Windows install.

Drivers on ARM Take Boots

On a PC the kernel finds the hardware. On this laptop the kernel knows only what the device tree tells it, and each device needs three things before its driver binds: a node in the tree, Samsung's firmware, and a board fact read out of the Windows side, such as an interrupt pin or a power pin. A driver setup that runs on the Acer Aspire 1, with the same SoC, is a starting guess here. The draft tree's keyboard interrupt, display bridge interrupt and audio memory region came from the Aspire 1, and so did the first guess at the speaker data line. All four were wrong on this board.

Every guess is tested with a reboot. The installed system logged 26 boots over the two days of this bring-up, 16 on 2026-10-11 and 9 on 2026-10-12, plus UEFI shell boots that leave no journal:

Subsystem	Test boots	What the boots settled
Speakers	4, then 1 as default	DSP memory region, DSP front ends, speaker data line, speaker supply
Battery, camera, backlight	2	the first one put the backlight on the wrong PWM channel
Keyboard interrupt	1, plus a UEFI shell boot	the pin, read from Samsung's ACPI tables
GPU firmware	1	the zap shader from the Windows driver store
Display bridge interrupt	1	the interrupt removed

Each risky change goes in as a one-time boot with grub-reboot, so a black screen costs one power cycle. Plan for several boots per device, and check every file the boot reads before each one.

The Dog Food Test

A port that boots to a login screen proves the boot path. A port that the author uses to write about it proves the keyboard, the desktop, the network and the toolchain all at once, for hours.

This article was drafted in GNOME Terminal on the Galaxy Book Go, by a native arm64 Claude Code running in the local Wayland session, with the repository on the internal UFS disk. The capture below was taken from that session:

$ cat /sys/class/dmi/id/sys_vendor /sys/class/dmi/id/product_name
SAMSUNG ELECTRONICS CO., LTD.
Galaxy Book Go (Wi-Fi)

$ uname -r
7.2.9+deb14-arm64

$ loginctl show-session 2 -p Type -p Remote
Remote=no
Type=wayland

$ file -bL ~/.local/bin/claude | cut -d, -f1-2
ELF 64-bit LSB executable, ARM aarch64

$ systemctl is-system-running
running

The same capture also shows where the laptop strains. It has 3.4 GiB of usable memory, and 1.3 GiB was in swap while this was written, with a browser, GNOME and the assistant open together.

At This Point You Should Have…

A Galaxy Book Go (NP340XLA or NP345XLA), Secure Boot off.
A Debian 13 installer stick for the Book Go that carries the community device tree, sc7180-samsung-galaxy-book-go.dtb.
A USB-C hub with a USB keyboard and a USB Ethernet adapter. USB boot is unreliable straight into the port and reliable through a hub.
The Windows DriverStore\FileRepository from a Book Go, copied off before Windows was wiped.
A second machine with SSH access, for the rescue path.

Step 1 — Put the Device Tree in the Boot Menu

The installer writes a stock GRUB menu. It has no devicetree line and no clk_ignore_unused pd_ignore_unused, so the installed kernel boots through ACPI and the machine resets. The installed system had never booted: its journal was empty.

The fix is a GRUB generator, scripts/09_galaxy-book-go, installed as /etc/grub.d/09_galaxy-book-go so it runs before 10_linux. Per kernel it writes three entries, each with the device tree and the clock options:

Entry id	What it is for
`gbgo-<kver>`	normal boot
`gbgo-safe-<kver>`	no GPU or DSP drivers, known-good initramfs, verbose
`gbgo-debug-<kver>`	`ignore_loglevel`

GRUB_DEFAULT names an entry id, never a position, and the menu waits 10 seconds. Every risky change after this is tried once with grub-reboot <id>, so a hang falls back to the default on the next power cycle.

The first boot from the internal disk reached running and the GDM login.

Step 2 — Make GNOME See the Keyboard

The text console took key presses, and GNOME took none. The internal keyboard (hid-over-i2c 0CF2:9028) reports a stray ABS axis, so udev tags it ID_INPUT_TABLET_PAD, and libinput hands it to GNOME as a tablet pad.

A udev rule clears the tags:

ACTION!="remove", SUBSYSTEM=="input", ATTRS{name}=="hid-over-i2c 0CF2:9028 Keyboard", ENV{ID_INPUT_TABLET}="", ENV{ID_INPUT_TABLET_PAD}="", ENV{ID_INPUT_KEYBOARD}="1", ENV{ID_INPUT_KEY}="1"

It must match the input0 node as well as event0. libinput reads the parent's tags too, so a rule on event* alone leaves the device classified as a tablet pad. With both, libinput list-devices reports Capabilities: keyboard pointer.

🔎 Tip: Do Not Restart GDM With Someone Logged In

A running GNOME keeps the old device classification until it restarts. Restarting gdm3 closes every open session. Log out instead.

Step 3 — Take the Firmware From Windows

The device tree asks for four Samsung-signed files under qcom/sc7180/samsung/galaxy-book-go/, and Wi-Fi needs a fifth. All five are in the Windows driver store. Where several driver versions sat side by side, the newest was taken:

File	Runs	Version
`qcdxkmsuc7180.mbn`	GPU zap shader	27.20.0910
`qcmpss7180_nm.mbn`	modem DSP, which hosts Wi-Fi	1.0.0800.8
`wlanmdsp.mbn`	Wi-Fi firmware, served to the modem	1.0.0970
`qcvss7180.mbn`	video decode and encode	27.20.0910
`qcadsp7180.mbn`	audio DSP	1.0.0750.3

A bad firmware file can reset a Qualcomm SoC on every boot. So nothing went into /usr/lib/firmware before it had run once:

Stage the files in /root/gbgo-fw-staging/.
Point /sys/module/firmware_class/parameters/path at the staging directory. That lasts for the running boot only.
Start each remote processor by hand and read the log.
For the GPU, which retries only at boot, use a one-time GRUB entry with firmware_class.path on the kernel command line.
Install with gbgo-firmware install, which also has promote and rollback.

The modem came up and ath10k_snoc reported WLAN.HL.3.2.2.c10-00767. Venus bound its decoder and encoder. The GPU came up as FD618 and gnome-shell left software rendering.

The audio DSP failed safely: segment outside memory range, Boot failed: -22. Samsung's image loads at 0x90b00000, and the draft tree reserves the Aspire 1's region at 0x8e400000. Step 9 fixes the region.

🔎 Tip: Know What Comes From Windows

Five signed binaries run this laptop's GPU, Wi-Fi, video and audio, and the only source for them is a Windows driver store. Debian's firmware packages ship none of them, and Samsung publishes no driver pack. They run on the modem, audio and video processors, outside what Linux can inspect. This repository keeps their hashes and source paths, and the files themselves stay out of it.

From Windows	Used for
Five firmware files	GPU, Wi-Fi, video decode, audio DSP
ACPI tables	keyboard interrupt, camera power pin, battery and charger addresses
Audio calibration database	speaker port and data line

A Book Go whose Windows was wiped without a copy of its driver store has no source for the firmware. Copy Windows\System32\DriverStore\FileRepository off before installing Linux. The Yoga C630 has the same dependency: its firmware comes from Lenovo's Windows driver releases, mirrored on GitHub.

Step 4 — Win the Wi-Fi Boot Race

On some boots Wi-Fi took about 60 seconds, and on others 17. The modem asks tqftpserv for wlanmdsp.mbn a few milliseconds after the service starts. If it asks first, the WLAN start stalls, and the modem's watchdog crashes it after 40 seconds (DOG detects stalled initialization). The kernel restarts it, and the second try works.

gbgo-wifi-kick runs after rmtfs and tqftpserv. If no Wi-Fi interface exists 10 seconds later, it stops and starts the modem's remote processor:

for r in /sys/class/remoteproc/remoteproc*; do
  [ "$(cat "$r/name")" = "$NAME" ] && R=$r
done
[ "$(cat "$R/state")" = running ] && echo stop > "$R/state"
sleep 2
echo start > "$R/state"

It checks for any interface under /sys/class/ieee80211/phy*/device/net/, never for a name. Under sid the interface is wld0. The restart was tested by hand. The service has not yet had to fire during a real boot.

Step 5 — Give Bluetooth an Address

The WCN3991 loads qca/crbtfw32.tlv from firmware-atheros and stays unconfigured, because nothing supplies a public address. gbgo-bt-addr sets a stable, locally administered one derived from /etc/machine-id, before bluetooth.service starts.

btmgmt from bluez 5.82 waits forever when its standard input is /dev/null, which is what systemd gives a service. Over SSH the same command returns at once. The script feeds it an empty pipe:

bm() { echo | timeout 3 btmgmt -i hci0 "$@" 2>&1; }

$ bluetoothctl show | head -1
Controller 02:33:FD:7E:34:33 (public)

Step 6 — Read the Keyboard Interrupt From Samsung's ACPI Tables

The keyboard worked, at a price. The draft tree wires its interrupt to TLMM GPIO 48, level-low, and GPIO 48 sits low permanently. The interrupt never clears, so the driver polls the keyboard about 1100 times a second:

i2c_hid_of 6-0005: i2c_hid_get_input: incomplete report (11/65535)

That was 785,000 log lines in 14 minutes and about 7 % of a CPU. Watching all 73 free input GPIOs while typing showed nothing.

Samsung's ACPI tables describe the board for the Windows drivers. Linux boots with a device tree here, so /sys/firmware/acpi/tables is empty and /dev/mem refuses the page. The UEFI shell can read it. A one-time GRUB entry chainloads Debian's shellaa64.efi, and a startup.nsh at the root of the ESP dumps the EFI-reserved block holding the tables and resets:

dmem 0xFFF22000 0xDC000 > mem.txt

acpi-from-dmem.py rebuilds the bytes and follows RSDP → XSDT → DSDT. The keyboard device, ECKB, declares:

GpioInt(Level, ActiveLow, PullUp) pin 0x140

On Qualcomm's ACPI, a pin of 0x40·k above the 120 real pins means direct-connect interrupt number k. 0x140 is interrupt 5, which is PDC pin 67, which is TLMM GPIO 104. The mapping checks out against three known pins: the SD card detect, the touchpad and the lid.

gbgo-dtb-patch builds the tree from the draft with three changed keyboard lines: interrupt 0x30 → 0x68, pin gpio48 → gpio104, bias-disable → bias-pull-up. After the change:

$ grep hid-over-i2c /proc/interrupts
 147:        799          0          0          0          0          0          0          0  msmgpio   104 Level     hid-over-i2c
 148:       4645          0          0          0          0          0          0          0  msmgpio    31 Level     hid-over-i2c

$ journalctl -k -b | grep -c "incomplete report"
0

The 799 interrupts on GPIO 104 are every key pressed since this boot, including the request that started this article.

🔎 Tip: Rehearse Firmware Scripts in QEMU

The shell dump was run first in qemu-system-aarch64 -M virt with AAVMF, the laptop's own shim, GRUB and UEFI shell, and twelve extra FAT disks. That rehearsal showed three things before any reboot of the laptop:

The shell runs startup.nsh only from the root of a filesystem.
With many FAT partitions, the ESP sits past fs9, so the script searches by a marker file.
Debian's UEFI shell has no acpiview. dmem with > redirection works.

Step 7 — Add the Power Button and the Lid, Quiet the Display Bridge

The same patch script adds two more devices to the tree and removes one wrong interrupt:

Power button. The draft disables pon@800, so pm8941-pwrkey never binds. Enabling it gives KEY_POWER and GNOME's power-off dialog.
Lid. The draft has no node. ACPI LID0 returns the raw level of GPIO 93, 1 for open, so a gpio-keys node reports SW_LID on TLMM 93, active low.
Display bridge interrupt. The draft gives the SN65DSI86 eDP bridge the Aspire 1's level-high interrupt on TLMM 11. GPIO 11 reads high here, so the kernel disabled it at every boot after 100,001 interrupts (irq 147: nobody cared), with I2C errors on the bridge's bus. Samsung's ACPI tables declare no interrupt for the bridge, so the patch removes it. Kernel warnings per boot fell from 62 to 26, and the panel works as before.

Suspend is s2idle only. One cycle from the login screen resumed in 0.6 s, woken by Shift on the built-in keyboard, with Wi-Fi, Bluetooth and Ethernet back. Wake by lid is still missing, so closing the lid locks the screen.

The PMIC clock was tested and left off. It reads 1970-01-03, because Samsung keeps the offset elsewhere, and it set the system clock to 1970 at boot. systemd-timesyncd and NTP keep the time instead.

Step 8 — Upgrade to sid Without Touching the Boot

A Lenovo Yoga C630 running the same plan lost its boot when sid's GRUB 2.14 blacked out every kernel. The Book Go upgrade was built around that:

Risk	Protection
sid replaces GRUB 2.12 and shim	`apt-mark hold` on all six `grub` and four `shim` packages
sid's `initramfs-tools` drops `qcom/` firmware	an initramfs hook adds the GPU firmware and zap shader
a new kernel becomes the default	the old kernel is held and stays default until a one-time boot passes
a rebuilt initramfs breaks boot	SAFE entries boot copies saved in `/boot/gbgo-known-good/` with `SHA256SUMS`
everything breaks	an `rsync` of `/` and the ESP on another machine, checked by hash

gbgo-boot-verify runs before every reboot. It checks the holds, hashes all nine ESP files against the backup, runs grub-script-check, confirms the default entry carries the tree and the options, reads the tree, initramfs and grub.cfg off the raw disk, and lists the storage, keyboard and display modules and the GPU firmware inside every initramfs.

The upgrade set up 1712 packages with no dpkg errors. GRUB stayed at 2.12:

$ dpkg-query -W grub-efi-arm64-bin
grub-efi-arm64-bin  2.12-9+deb13u2

Kernel 7.2.9 passed a one-time boot and became the default. On 7.2.9 the GPU loads after the root mount, so the zap shader failure seen on 6.12 is gone.

🔎 Tip: Three Names That Change Under You

Disk. The root partition was sda3 at install and is sdc3 today. Every boot entry uses UUID=.
Wi-Fi. sid's systemd renames wlan0 to wld0 a few seconds into boot.
sudo path. sid's sudo no longer includes /usr/local/sbin. Call the scripts by full path.

apt autoremove after the upgrade would also have removed the fallback kernel and mesa-va-drivers. Both are now marked manual.

Step 9 — Bring Up the Speakers

The draft tree has no sound card, no codec and the Aspire 1's audio DSP region. Samsung's ACPI tables name the headset codec, a Realtek RT5682 on i2c@a8c000, and declare no speaker amplifier on any I2C bus. The audio device AUDD drives two GPIOs, TLMM 21 and 20, so the amplifier is a plain I2S amplifier with enable pins.

gbgo-dtb-audio.py adds the audio parts to the tree:

Change	Why
DSP region `0x8e400000`, grown to `0x5500000`	covers Samsung's `0x90b00000`–`0x93900000`
RT5682 on `i2c@a8c000`	from the ACPI tables
amplifier enable on TLMM 21	`AUDD`'s first GPIO
`qcom,sc7180-qdsp6-sndcard` with four front ends	the driver for this SoC's DSP audio
speaker on tertiary MI2S, data line 1	Samsung's Windows calibration
PM6150L `bob` always on at 3.296 V	Samsung's power settings for `AUDD`

The port comes from Samsung's Windows audio package. Its calibration database maps each audio device to a port record: SPEAKER_OUT sits on TERTIARY_MI2S_RX, with the same I2S fields as the headset on PRIMARY_MI2S_RX. The supply comes from Samsung's power tables: AUDD votes the PM6150L buck-boost regulator bob on at 0x324B00 µV, 3.296 V. Linux gives bob no user, so nothing turns it on without the tree change.

With all of that in place the card comes up, and the speakers still stay silent under PipeWire. The amplifier plays 16-bit stereo, and PipeWire opens the port as 8 channels of 24-bit audio:

Stream to tertiary MI2S	Speakers
PipeWire default, 8 channels `S24_LE`	silent
`speaker-test -D hw:0,1 -c 2 -F S16_LE`	tone

A WirePlumber rule opens the speakers as 16-bit stereo:

monitor.alsa.rules = [
  {
    matches = [ { device.name = "alsa_card.platform-sound" } ]
    actions = { update-props = { device.profile = "pro-audio" } }
  }
  {
    matches = [ { node.name = "alsa_output.platform-sound.pro-output-0" } ]
    actions = {
      update-props = {
        node.description = "Speakers"
        audio.format = "S16LE"
        audio.channels = 2
        audio.position = [ FL, FR ]
      }
    }
  }
]

The mixer route TERT_MI2S_RX Audio Mixer MultiMedia1 is saved with alsactl store. After a reboot on the audio tree:

$ cat /sys/class/remoteproc/remoteproc1/name /sys/class/remoteproc/remoteproc1/state
adsp
running

$ amixer -c0 cget name='TERT_MI2S_RX Audio Mixer MultiMedia1' | tail -1
  : values=on

$ wpctl status | grep -A1 Sinks
 ├─ Sinks:
 │  *   53. Speakers                            [vol: 0.80]

The headphone jack and the internal microphone are untested. The RT5682's headphone switches start muted, and Samsung's calibration puts the microphone on quaternary MI2S, which the tree does not wire yet.

Step 10 — Use the Video Decoder

The Venus video block decodes H.264, VP8, VP9, HEVC and MPEG-2 through V4L2. Debian's ffmpeg reaches it as h264_v4l2m2m and vp9_v4l2m2m. Twenty seconds of 1080p30 test pattern, 600 frames:

Clip	Decoder	Wall	CPU
H.264	software	2.42 s	11.62 s
H.264	Venus	7.42 s	0.43 s
VP9	software	2.92 s	10.85 s
VP9	Venus	7.56 s	0.45 s

Venus decodes about 80 frames a second at about a twenty-fifth of the CPU.

The browsers decode in software. Chrome on Linux reaches hardware decoders through VA-API, and no VA-API driver exists for this chip:

$ vainfo
libva info: Trying to open /usr/lib/aarch64-linux-gnu/dri/msm_drv_video.so
libva info: va_openDriver() returns -1
vaInitialize failed with error code -1 (unknown libva error),exit

Chrome 155 and Chromium 154 never opened /dev/video0 while playing a video, with or without the V4L2 and VA-API feature flags. Firefox ESR 153 with media.hardware-video-decoding.force-enabled opens the V4L2 decoder, gets one frame back with no timestamp, and falls back to software. YouTube in Chrome runs its renderer at 47 to 81 % of a core, on a laptop with 3.4 GiB of memory.

🔎 Tip: Watch YouTube in mpv

mpv with yt-dlp uses the decoder:

mpv --hwdec=v4l2m2m-copy 'https://www.youtube.com/watch?v=aqz-KE-bpKQ'

It prints Using hardware decoding (v4l2m2m-copy). Twenty seconds each of 1080p60 H.264, 720p60 H.264 and 1080p60 VP9 decoded with 0 dropped frames, at about 7.5 s of CPU per run including yt-dlp. Those runs discarded the video output, so playback on the screen is not yet measured.

Step 11 — Charge the Battery

Samsung's ACPI tables give the embedded controller I2C resources at 0x09 and 0x0b on i2c@890000. 0x0b answers the Smart Battery word commands, and Linux's sbs-battery driver reads it once the tree has a node:

battery@b {
    compatible = "sbs,sbs-battery";
    reg = <0x0b>;
};

POWER_SUPPLY_MANUFACTURER=SDI
POWER_SUPPLY_MODEL_NAME=4432D53
POWER_SUPPLY_CHARGE_FULL_DESIGN=5320000
POWER_SUPPLY_CYCLE_COUNT=8

The battery then reported Not charging with the adapter in, and it was right: zero current, constant charge. 0x09 is a TI BQ25710 charger (manufacturer ID 0x0040, device ID 0x0089). Its readings after a one-shot conversion:

Register	Reading
`ChargerStatus`	`0x8000`, adapter present
input voltage	15.04 V
`ChargeCurrent`	0 mA
`MaxChargeVoltage`	8760 mV

The BQ25710 datasheet explains the zero: the charger "terminate[s] charging if the charger does not receive a write MaxChargeVoltage() or write ChargeCurrent() command within 175 s". The charge voltage of 8760 mV is set before Linux starts, and nothing in Linux writes the charger, so charging stops within 175 s of Linux taking over.

gbgo-charge plays the host. Every 60 s it reads the battery's own requests and passes them to the charger, capped at 8760 mV and 2048 mA, and it writes 0 when the battery reports full, asks for 0, or reads outside 0–45 °C. If the service stops, the charger's timer ends charging within 175 s.

$ journalctl -u gbgo-charge | tail -1
gbgo-charge[827]: charge: 8760 mV, 2048 mA (gauge asks 8860 mV 6916 mA; battery 7993 mV, 26 C, charger status 0x8400)

$ upower -i /org/freedesktop/UPower/devices/battery_sbs_4_000b | grep -E "state|time to"
    state:               charging
    time to full:        3.0 hours

The battery went from 30 % to 68 % in 53 minutes across a reboot, with the service restarting on its own.

Step 12 — Power the Camera

The camera never enumerated, and nothing in the tree mentions one. Samsung's ACPI tables put it on USB: ICAM sits on port 4 of the internal hub EHUB, and the hub's power-on method switches it on:

Method (_PS0, 0, NotSerialized)
{
    ADBG ("EHUB_PS0")
    \_SB.GIO0.CAME = One
}

CAME is an output on TLMM GPIO 13, which the boot firmware leaves low. Driving it high brought the camera up at once:

usb 1-1.4: Product: 720p HD Camera
usb 1-1.4: Manufacturer: Chicony
uvcvideo 1-1.4:1.0: Found UVC 1.00 device 720p HD Camera (04f2:b72f)

A gpio-hog in the tree holds GPIO 13 high from boot. The camera now registers before the video decoder, so it takes /dev/video0 and video1, and Venus moves to video2 and video3.

Step 13 — Find the Backlight Channel

Samsung's panel configuration names a PMIC-driven PWM backlight at 6250 Hz, and the boot firmware leaves the PM6150L light pulse generator (LPG) running. With pwm-backlight on LPG3, the GNOME slider moved LPG3's duty cycle and the screen stayed the same.

LPG4 was running at 50 % with nothing in Linux holding it. Exported through /sys/class/pwm, it was switched between 5 % and 100 % every 2 seconds:

C=/sys/class/pwm/pwmchip1/pwm3
for i in $(seq 1 20); do
  echo 8000 > $C/duty_cycle; sleep 2
  echo 159688 > $C/duty_cycle; sleep 2
done

The screen got darker and brighter in step. pwm-backlight now uses channel index 3, and the slider drives LPG4:

 pwm-3   (backlight           ): requested
  requested configuration:  enabled, 158889/160000 ns, normal polarity

What Works?

Hardware	State
GPU	🟢 Adreno 618, `FD618`, `simple_ondemand`
Video	⚠️ venus decoder in mpv and ffmpeg. Browsers decode in software
Wi-Fi	🟢 5 GHz, 80 MHz, with Samsung's modem firmware
Bluetooth	🟢 with `gbgo-bt-addr`. ⚠️ Pairing and audio untested
Keyboard, touchpad	🟢 keyboard on GPIO 104, touchpad on GPIO 31
Power button, lid	🟢 lid locks the screen
Suspend	⚠️ `s2idle`, one cycle tested, the lid does not wake it
USB-C	⚠️ works if plugged at boot. No UCSI node
Speakers	🟢 16-bit stereo through PipeWire
Headphones, mic	⚠️ untested. Mic not wired in the tree
Battery	🟢 `sbs-battery` on the EC's gauge. Charging through `gbgo-charge`
Brightness	🟢 PM6150L LPG4 dims and brightens the screen
Camera	🟢 Chicony 720p UVC, powered by GPIO 13

Compare and Contrast

	Yoga C630	Galaxy Book Go
SoC	Snapdragon 850 (SDM850)	Snapdragon 7c Gen 2 (SC7180)
Device tree	🥇 upstream, in Debian's kernel	🥈 community draft, patched
Firmware source	Lenovo's Windows driver releases	this laptop's Windows driver store
Audio	🥇 works	🥈 speakers work, headphones and mic untested
Battery	🥇 EC driver, the EC charges	🥈 gauge driver, a service drives the charger
Browser video decode	❌ software	❌ software
Blocking driver	`ipa` resets the SoC	none found
Boot loader	GRUB 2.12 held	GRUB 2.12 held

So, Is It a Daily Driver?

For writing, terminals, a browser and SSH, yes. This article is the test: the keyboard, the Wayland desktop, Wi-Fi and the assistant ran together for the whole session without a fault.

The battery reading, charging, brightness, speakers and camera close the gaps that kept it on a desk. What remains: browsers decode video in software, and with 3.4 GiB of memory YouTube stutters in Chrome and plays through the hardware decoder in mpv. The headphone jack and microphone are untested, USB-C works only when plugged in at boot, and every Wi-Fi, GPU, video and audio feature rests on firmware copied out of Windows.

Summary

The goal of this article was to run a current Debian with a full desktop on a laptop that has no upstream support, and then use it. The key to the solution was treating Samsung's Windows install as the reference: its driver store supplied the firmware and the audio calibration, and its ACPI tables supplied the keyboard interrupt, the camera power pin and the battery addresses that the community tree lacked. Each fix took one or more test boots. The results were:

🟢 Debian sid, kernel 7.2.9 and GNOME Shell 50.5 on Wayland, with GPU acceleration, Wi-Fi, Bluetooth, keyboard, touchpad, power button and lid.
🟢 The speakers play, configured from Samsung's Windows calibration: tertiary MI2S, bob at 3.296 V, 16-bit stereo.
⚠️ The video decoder works in mpv at about a twenty-fifth of the CPU, and browsers decode in software.
🟢 The keyboard interrupt flood is gone: 0 incomplete report lines per boot. The display bridge interrupt storm is gone too.
🟢 The sid upgrade left the boot path byte-identical, checked by script before each reboot.
🟢 This article was written, fact-checked and built on the machine.
🟢 The battery reads and charges: a service relays the battery's requests to the BQ25710 charger every 60 s.
🟢 The camera, powered by GPIO 13, and brightness, on PM6150L LPG4.
⚠️ Wi-Fi, GPU, video and audio depend on five firmware files found only in a Windows driver store.
⚠️ Suspend has one tested cycle, and the Wi-Fi restart service has not yet fired on a real boot.

Scope: one Galaxy Book Go (NP340XLA-KA3CA, Wi-Fi model, BIOS P00AJW.036), Debian sid as of 2026-10-11, Samsung firmware from the Windows driver store, brought up over two days and 26 logged boots, with the dog food session and the audio, video, battery, camera and backlight work on 2026-10-12.

The strategy for bringing up Debian on the Galaxy Book Go was validated with an incremental step by step approach.

References

Repository: https://github.com/xbill9/linux-arm-adventures
Galaxy Book Go notes: galaxy-book-go/
Debian on the Yoga C630, the companion article: c630-debian
Bring-up history: aarch64-laptops/debian-cdimage#21
Current community status: velvet-os/imagebuilder#136
Draft device tree and kernel notes (readme.qc7): hexdump0815/linux-mainline-qcom-kernel
BQ25710 charger datasheet: ti.com/lit/ds/symlink/bq25710.pdf

Small Apache-2.0 Models Against Gemma 4 E2B on One TPU v5e: Which Ones Keep Up?

xbill — Tue, 13 Oct 2026 00:57:26 +0000

This article provides a step by step guide to surveying small open-weight models on one Google Cloud TPU v5e chip through vLLM, with Gemma 4 E2B as the baseline. Every model is Apache-2.0 or MIT licensed, so any of them can be repacked and redistributed, and every one is served on the same pinned image with the same settings, then timed and scored by code.

Gemma 4 E2B is small and workable on a v5e-1, but its attention is unusual: two head sizes, 256 for sliding-window layers and 512 for full-attention layers, one KV head, and 20 of its 35 layers reading another layer's KV cache. The question here is which plainer model can stand in for it. Qwen3-4B matches Gemma on both accuracy tests at 0.57x the output. SmolLM3-3B and Granite 3.3 2B match it on tool calls at 0.76x and 0.93x. With long prompts, Gemma stays ahead of every model of comparable accuracy.

https://github.com/xbill9/model-sweep

Why One v5e Chip?

A v5litepod-1 is the cheapest TPU Google Cloud rents and is usually available as flex-start capacity in us-west4-a. It has 15.75 GiB of HBM, and vLLM gets 0.9 of that, so a bf16 model has to leave room for its KV cache in about 14 GiB.

That rules the field down to roughly 4B parameters. Gemma 4 E2B loads in 8.97 GiB on this chip, which makes it one of the larger models in the lineup.

At This Point You Should Have…

A Google Cloud project with the Cloud TPU API enabled and flex-start TPU v5e quota in us-west4-a
A GCS bucket for results, and a Hugging Face token saved in Secret Manager as hf-token
gcloud with the alpha components, and Python 3 with no extra packages
A clone: git clone https://github.com/xbill9/model-sweep

Step 1 — Build the Lineup From the Hub

survey.py reads each candidate's license, architecture and weight size from the Hugging Face API and decides whether it fits the chip. A model is in when its license is Apache-2.0 or MIT and at least 2 GiB is left for the KV cache. vLLM on TPU runs some architectures as native JAX models and the rest through a PyTorch fallback, so the survey splits the lineup into two tiers:

python3 survey.py

| `tpu-vllm-v5e1-qwen3_4b` | `Qwen/Qwen3-4B` | apache-2.0 | Qwen3ForCausalLM | 4.02 | 7.49 | 6.68 | JAX | ✅ | 1 |
| `tpu-vllm-v5e1-gemma4_2b` | `google/gemma-4-E2B-it` | apache-2.0 | Gemma4ForConditionalGeneration | 5.12 | 9.54 | 4.63 | JAX | ✅ | 1 |
| `tpu-vllm-v5e1-smollm3_3b` | `HuggingFaceTB/SmolLM3-3B` | apache-2.0 | SmolLM3ForCausalLM | 3.08 | 5.73 | 8.45 | torchax |  | 2 |
| `tpu-vllm-v5e1-qwen25_3b` | `Qwen/Qwen2.5-3B-Instruct` | other | Qwen2ForCausalLM | 3.09 | 5.75 | 8.43 | JAX |  | excluded: license |
| `tpu-vllm-v5e1-qwen3_8b` | `Qwen/Qwen3-8B` | apache-2.0 | Qwen3ForCausalLM | 8.19 | 15.26 | -1.08 | JAX |  | excluded: memory |

Tier 1: 8 models. Tier 2: 7 models. Excluded: 2.

Qwen2.5-3B is out on its license and Qwen3-8B on memory. Everything else goes into two arms files, one per tier.

Step 2 — Describe Each Model in One Line

An arms file lists one model per line: the checkpoint, its rig name, what to measure, and any serving flags it needs. The tool-call parser and the thinking switch differ by family:

google/gemma-4-E2B-it=gemma4_2b=load+long+gen=parser=gemma4,mm=image+audio+video
Qwen/Qwen3-4B=qwen3_4b=load+long+gen=parser=hermes,nothink
HuggingFaceTB/SmolLM3-3B=smollm3_3b=load+long+gen=parser=hermes,nothink
ibm-granite/granite-3.3-2b-instruct=granite33_2b=load+long+gen=parser=granite

nothink sends chat_template_kwargs {"enable_thinking": false} with every request, so Qwen3 and SmolLM3 answer directly like the others. mm= turns off the image, audio and video inputs of multimodal checkpoints. A unit test checks every line against the lineup and against the tool-call parsers registered in the pinned vLLM build.

Step 3 — Launch One Chip per Tier

tpu/launch.sh bundles the committed tree, uploads it, and creates a flex-start queued resource whose startup script serves every arm in turn, uploads results after each one, and deletes the TPU at the end:

tpu/launch.sh arms-v5e1-tier2.txt --go

arms:    tpu/arms-v5e1-tier2.txt (7 arms)
bundle:  gs://aisprint-491218-bucket/model-sweep/inputs/sweep-86d97ba02-0837b648.tgz
results: gs://aisprint-491218-bucket/model-sweep/2026-10-12-v5e1-tier2/
Created queued_resource [ms-2026-10-12-v5e1-tier2-1349].
created ms-2026-10-12-v5e1-tier2-1349; it deletes itself when the last arm finishes, or at 4h

Both tiers ran on vllm/vllm-tpu@sha256:19a1a0526476f902eb83e1057f3d8938f35b457dd7716a30e9ab4f7bee90d507, vLLM 0.29.1rc1.dev468+g0b7f11a1e, so the tier 2 models pair with tier 1's Gemma run:

[sweep 2026-10-11T17:14:28Z] start 2026-10-11-v5e1-tier1 on us-west4-a, arms arms-v5e1-tier1.txt
[sweep 2026-10-11T18:24:30Z] DONE: all arms attempted
[sweep 2026-10-12T13:51:26Z] start 2026-10-12-v5e1-tier2 on us-west4-a, arms arms-v5e1-tier2.txt
[sweep 2026-10-12T16:02:12Z] DONE: all arms attempted

Step 4 — Measure Speed and Accuracy the Same Way

Every model is served with --max-model-len 4096 --max-num-seqs 16 and prefix caching on. Each one gets three measurements:

Output speed: 256 output tokens per request at 1, 4 and 16 requests at once, median of three passes
Long prompts: the same with 1,024- and 3,072-token prompts that never repeat, so the prefix cache cannot help
Accuracy: all 1,319 GSM8K problems, greedy, with a 2,048-token answer limit, and the 400 BFCL simple tool-call records for models with a tool-call parser

Answers are scored in code. Differences against Gemma are paired on the same problems, with a 95% range from 10,000 resamples; a range that includes zero means no measured difference.

Step 5 — File the Results and Build the Report

python3 sweep_report.py file 2026-10-12-v5e1-tier2
python3 sweep_report.py report 2026-10-12-v5e1-tier2 --baseline-run 2026-10-11-v5e1-tier1

filed smollm3_3b -> tpu-vllm-v5e1-smollm3_3b/benchmarks/runs/2026-10-12-v5e1-tier2
wrote results/2026-10-12-v5e1-tier2/REPORT.md (6 served, 1 failed)

file moves each model's logs and scored answers into its own directory. report refuses to pair two runs that served different image digests.

How Accurate Is Each Model?

Against Gemma 4 E2B, in percentage points, with the 95% range:

Model	GSM8K	vs Gemma	BFCL	vs Gemma
Gemma 4 E2B	0.911	—	0.927	—
Qwen3-4B	0.915	+0.4 (-1.2 to +2.0)	0.940	+1.2 (-1.8 to +4.2)
Phi-4-mini	0.895	-1.6 (-3.2 to -0.1)	not scored	—
SmolLM3-3B	0.858	-5.3 (-7.2 to -3.5)	0.945	+1.7 (-1.0 to +4.8)
Qwen3-1.7B	0.827	-8.4 (-10.5 to -6.4)	0.935	+0.8 (-2.2 to +3.8)
Granite 3.3 2B	0.817	-9.5 (-11.6 to -7.4)	0.920	-0.7 (-3.8 to +2.2)
Qwen2.5-1.5B	0.729	-18.2 (-20.6 to -15.8)	0.910	-1.7 (-5.2 to +1.8)
Qwen3-0.6B	0.599	-31.2 (-34.0 to -28.5)	0.825	-10.3 (-14.2 to -6.5)
Qwen2.5-0.5B	0.469	-44.3 (-47.0 to -41.4)	0.818	-11.0 (-14.8 to -7.2)
SmolLM2-1.7B	0.461	-45.0 (-47.8 to -42.2)	no parser	—

Five models are level with Gemma on BFCL tool calls, down to Qwen2.5-1.5B. Only Qwen3-4B is level on GSM8K, and from 1.7B down, Qwen3-1.7B aside, the score falls by 18 to 45 points.

How Fast Is Each Model?

Output tokens per second with 16 requests at once, and against Gemma:

Model	Weights GiB	16 requests	vs Gemma	16 × 3,072-token prompts	vs Gemma
Qwen2.5-0.5B	0.92	4176	2.21x	2596	2.48x
Qwen3-0.6B	1.40	3210	1.70x	1301	1.24x
Qwen2.5-1.5B	2.88	2780	1.47x	1570	1.50x
Qwen3-1.7B	3.78	2049	1.09x	926	0.88x
Gemma 4 E2B	8.97	1887	1.00x	1049	1.00x
Granite 3.3 2B	4.92	1761	0.93x	732	0.70x
SmolLM3-3B	6.22	1440	0.76x	725	0.69x
Phi-4-mini	8.31	1136	0.60x	404	0.39x
Qwen3-4B	7.49	1080	0.57x	373	0.36x

Gemma carries the largest weights in the table and still out-serves every model with more than 4 GiB of weights. With 3,072-token prompts the gap widens: Granite and SmolLM3 drop to about 0.7x, and Qwen3-4B to 0.36x.

Why Does Gemma Hold Up on Long Prompts?

The KV cache. Each model's boot log reports how many tokens of cache fit in the HBM left after its weights:

Model	KV heads	Head size	Layers	KV cache tokens
Gemma 4 E2B	1	256 / 512	35, 20 sharing	321,280
Granite 3.3 2B	8	64	40	125,184
SmolLM3-3B	4	128	36	120,320
Qwen3-1.7B	8	128	28	100,096
Qwen3-4B	8	128	36	50,688

Gemma stores one KV head per layer, 28 of its 35 layers keep only a 512-token sliding window, and 20 layers reuse an earlier layer's cache. That leaves it 2.6 to 6.3 times the cache of the four models below it, even though its weights are the largest.

The same features that make Gemma awkward to port are the ones paying for its long-prompt speed on this chip.

🔎 Tip: Turn Thinking Off Per Request

Qwen3 and SmolLM3 think by default, which spends the GSM8K answer budget on reasoning text. Passing the chat template switch with each request turns it off without a different checkpoint:

{"chat_template_kwargs": {"enable_thinking": false}}

With it, Qwen3-4B leaves no GSM8K answer at the 2,048-token limit, and SmolLM3-3B leaves 16 of 1,319.

🔎 Tip: DeepSeek-R1-Distill-Qwen-1.5B Needs a One-Line Fix on TPU

The MIT-licensed R1 distill stops at load on this image:

FAILED deepseek-ai/DeepSeek-R1-Distill-Qwen-1.5B (exited) after 75s: AttributeError: 'Qwen2Config' object has no attribute 'text_config'. Did you mean: 'get_text_config'?

tpu_inference/models/jax/qwen2.py builds an untied lm_head with getattr(hf_config, 'hidden_size', hf_config.text_config.hidden_size). Python evaluates the default argument first, so any Qwen2 checkpoint with tie_word_embeddings: false fails. The Qwen2.5 Instruct checkpoints tie their embeddings and load.

Compare and Contrast

Model	Size	GSM8K	BFCL	16 requests	Long prompts
Gemma 4 E2B	5.12B	0.911	0.927	1.00x	1.00x
Qwen3-4B	4.02B	0.915	0.940	0.57x	0.36x
SmolLM3-3B	3.08B	0.858	0.945	0.76x	0.69x
Qwen3-1.7B	2.03B	0.827	0.935	1.09x	0.88x
Granite 3.3 2B	2.53B	0.817	0.920	0.93x	0.70x

Long prompts are 16 requests of 3,072 tokens; every ratio is against Gemma on the same image.

So, Which One?

For the best accuracy per chip, Gemma 4 E2B: no other model matches both its scores and its speed on a v5e-1. For accuracy with standard attention, Qwen3-4B, level with Gemma on both tests at a bit over half the output. For tool calling at Gemma's size with the plainest architecture, SmolLM3-3B: level on BFCL, 5.3 points lower on math, three quarters of the speed. For tool calling at close to Gemma's speed, Granite 3.3 2B or Qwen3-1.7B, if math matters less. The models from 1.7B down run faster than Gemma but, apart from Qwen3-1.7B, give up 18 to 45 points of GSM8K.

Teardown

Each queued resource deletes itself after its last model. Confirm nothing is left:

gcloud alpha compute tpus queued-resources list --zone us-west4-a

Summary

The goal of this article was to find which small Apache-2.0 or MIT model can stand in for Gemma 4 E2B on one TPU v5e chip. The key to the solution was serving every model on one pinned vLLM image with identical settings, scoring every answer in code, and pairing each model against Gemma on the same problems. The results were:

🟢 Qwen3-4B is level with Gemma on GSM8K and BFCL, at 0.57x the output
🟢 SmolLM3-3B, Granite 3.3 2B, Qwen3-1.7B and Qwen2.5-1.5B are level with Gemma on BFCL tool calls
🟢 Gemma has the largest weights in the lineup and still out-serves every model with more than 4 GiB of weights
⚠️ Gemma's shared, single-head, sliding-window KV cache holds 2.6 to 6.3 times the tokens of Granite, SmolLM3, Qwen3-1.7B and Qwen3-4B, and keeps its lead with 3,072-token prompts
⚠️ Every model under 2B parameters except Qwen3-1.7B scores 18 to 45 points below Gemma on GSM8K
❌ DeepSeek-R1-Distill-Qwen-1.5B and OLMo 2 1B do not load on this image

Scope: one TPU v5e chip (v5litepod-1) in us-west4-a, flex-start, two runs on 2026-10-11 and 2026-10-12 on the same image digest, vLLM 0.29.1rc1.dev468+g0b7f11a1e. Speed is the median of three passes; accuracy is one greedy pass. Gemma's figures come from the first run and pair with the second through the shared image. OLMo 2 1B fails to compile (it asks for 133.56 MB of the chip's 128 MB vector memory). Ministral 3 3B and Qwen3.5 0.8B and 2B served but are left out of the tables: on this image Ministral returned server errors on part of each test and Qwen3.5 returned unusable text, so their scores describe the serving path. Phi-4-mini produced no parsed tool calls with the phi4_mini_json parser and is reported on GSM8K and speed only. Parts of the analysis and writing were done with AI assistance (Claude); every figure comes from the committed report files.

The strategy for surveying small open models against Gemma 4 on a TPU v5e was validated with an incremental step by step approach.

References

Code, logs, reports and scored answers: https://github.com/xbill9/model-sweep
vLLM TPU: https://github.com/vllm-project/tpu-inference
Gemma 4 E2B: https://huggingface.co/google/gemma-4-E2B-it
Qwen3-4B: https://huggingface.co/Qwen/Qwen3-4B
SmolLM3-3B: https://huggingface.co/HuggingFaceTB/SmolLM3-3B
Granite 3.3 2B: https://huggingface.co/ibm-granite/granite-3.3-2b-instruct
GSM8K: https://huggingface.co/datasets/openai/gsm8k
Berkeley Function Calling Leaderboard: https://gorilla.cs.berkeley.edu/leaderboard.html
Cloud TPU v5e: https://cloud.google.com/tpu/docs/v5e