Most of the talk about LLM security is about the model itself. Is it jailbroken, is it biased, does it leak its training data. Those are fair questions, but they miss where the real risk sits once a model is wired into an agent.
An LLM on its own just produces text. Connect it to tools, plugins, skills, file access, network services or a shell, and it becomes something else: an interpreter that turns information into actions.
This is the first post in a short series on what that means for security. It is the thinking behind a lot of the design decisions in Marvin, and it applies to anyone running models locally with tools turned on.
The question that matters
The obvious question is "does this model contain malicious code?" The answer is almost always no, because weights are not code. The question I think matters more is this:
Can untrusted information influence an LLM in a way that causes its surrounding execution environment to perform unauthorised actions?
Put that way, the attack surface is a stack of layers:
The model and its training data.
The model distribution and conversion supply chain.
The inference engine and its extensions.
Plugins, skills and agent frameworks.
The tools and permissions exposed to the model.
The data and documents the model is asked to interpret.
None of these layers needs to contain malicious code to be a problem. The model can be influenced by instructions inside data, trained to recognise a trigger, or led into harmful actions through tool interfaces that look completely legitimate.
Two supply chains, not one
When people pick a model they look at parameter count, quantisation, VRAM, tokens per second, and how well it codes or calls tools. Provenance barely comes up.
But a model has a supply chain of its own, and it runs alongside the software supply chain we already worry about.
Software supply chain
Model supply chain
Operating system
Training data
Inference engine
Pretraining
Plugins and extensions
Fine-tuning
Dependencies
Preference optimisation
Container images
Model conversion
Runtime configuration
Quantisation
Tool implementations
Model distribution
Agent framework
Agent integration
A compromise at an early stage of the model chain can survive every later stage. It can come out the other end as a perfectly ordinary looking model file on a public repository.
The two chains also don't cover for each other. A trusted inference engine does not make an untrusted model safe. A well sourced model does not make an agent secure if that agent can load arbitrary plugins or run whatever commands it likes.
Persuading the environment
A model cannot write to disk, install a plugin or send data anywhere unless the system around it provides a route. That is true, and it is often used to wave the problem away.
The catch is that agent frameworks exist precisely to provide those routes. A permissive framework turns generated requests into real actions. A hypothetical chain looks like this:
A model or plugin supplies conditional instructions.
The model interprets those instructions.
The model generates a request to create or modify a file.
A tool executes the requested operation.
A plugin or skill lands in a persistent location.
A later process loads or executes that component.
The component uses whatever the host process is allowed to touch.
No step in that chain needs a GPU escape or an exploit in the inference engine. Every step uses a capability the framework handed over on purpose.
The model does not necessarily need to escape its execution environment if it can persuade the environment to perform the actions that escape it.
The risk climbs sharply when file writing, code execution, plugin installation, persistent skill changes and network access all sit with the same agent and nothing independent checks them.
Why this matters if you self-host
The common home lab pattern is to download a model from a public repository, plug it into an agent framework, and give it broad tool access so it can actually be useful. That one setup stacks up every risk above. Nobody knows where the model came from, nobody has reviewed the extensions, and the agent can do almost anything.
I am not saying every model is backdoored. Most are not. The point is that the architecture should not depend on that being true.
What's next
The next post looks at the model layer itself: how training data poisoning can plant conditional behaviour in the weights, why the trigger does not have to be a fixed phrase, and why there is currently no way to certify a model as clean. Read part two.
llmgo now decodes Qwen3.8-27B at 56–61 tokens a second at ~10k tokens of context on a 16 GB RTX 5060 Ti, and 44–53 tokens a second at 60–76k. It started 28 runs earlier answering in gibberish.
llmgo is a dependency-free Go inference engine for GGUF models, and it now serves Marvin through its UI and gateway router. It serves Qwen3.8-27B-UD-Q2_K_XL: 64 layers, 48 of them Gated DeltaNet (linear attention) and 16 full attention, head dim 256, a 248k vocabulary, and one MTP (multi-token prediction) layer that drafts the next tokens. The weights are 9 GB of mostly 2- and 3-bit IQ quants; the K/V cache is q4_0, with 130k tokens of context available.
The card is an RTX 5060 Ti 16 GB (Blackwell, sm_120). Its VRAM reads at 426 GB/s measured, so reading every weight once costs ~22 ms. Every speed-up below is ultimately about how close a decode step gets to that number.
Measure
llmgo today
Decode, ~10k context
56–61 tok/s
Decode, 50–76k context
44–53 tok/s
Prefill, 9.6k prompt
11.4–11.7 s
The steps that moved the ~10k number
Most of the ~10k speed came from a handful of early steps; the chart further down shows them as a curve.
All 64 layers on the GPU, and gibberish (s01–s04). The first deploy answered in nonsense. Putting 63 layers on the GPU (s02) or turning MTP off (s05) gave correct answers at ~28 tok/s. So the fault was in the MTP path.
The MTP verify fix (s06): 41 tok/s. With the output head on the GPU, the verify step recomputed every position's logits on the CPU from stale activations and overwrote the GPU's. Fixed, 88% of drafts were taken and each forward produced 2.58 tokens.
The production config (s09): 34 tok/s. Moving to the production setup (130k available context, q4_0 K/V cache, 1024-token prefill chunks, prompt cache) cost speed, but from here every number is what Marvin actually runs.
MTP drafting on the GPU, small chunks on the matvec (s10–s11): 45–46 tok/s. Priming the draft dropped from ~300 ms to 40 ms, and 8-token chunks stopped taking the dequantise-and-GEMM path.
Seven drafts, top-k on the GPU, the whole model kept on the GPU (s12): 49 tok/s. Up to 7 drafts at p-min 0.8; the output head's top-k on the device cut logits per forward from 11.6 to 4.7 ms.
k_mmv, verify on int8 tensor cores (s16): 52–54 tok/s. A verify forward checks up to 8 positions. The matvec did each position's dot products in turn; k_mmv decodes each weight once and multiplies all positions with one tensor-core mma.
The matvec reads 2 or 4 rows per lane (s18): 54 tok/s. Each position's activations are loaded once for several rows.
Why tokens per second lies a little
From s16 on, a forward at ~10k has cost 57–64 ms, and the tok/s spread around it is mostly how well each answer drafted.
With speculative decoding, tok/s = tokens per forward ÷ time per forward. Tokens per forward depends on how predictable the text is: 2.9 on one answer, 3.4 on the next. The same build (s18, restarted twice as s19 and s20) gave 54.2, 50.9 and 49.6 tok/s. A forward that verifies more positions also takes longer, so a good-drafting answer raises both numbers.
So from s21 on, a change only counted if the milliseconds per forward moved. Runs that didn't move it still went into the log, the CSV and the write-up, but not onto the progress chart. That rule cut a lot of false hope: several builds that looked 5 tok/s faster were the same kernel on a luckier answer.
Measuring instead of guessing
The middle stretch (s16–s24) produced few gains, and the lesson was to stop guessing and instrument.
Guesses that didn't pay. An L2 prefetch hint did nothing in three runs and was removed. A cp.async weight ring for k_mmv doubled its shared memory; the theory was lost occupancy, but the probe later showed 6 blocks per SM without the ring were slower than 3 with it.
Removal probes, and their limit. Startup probes timed each kernel with a part removed: no weight load, no mma, no activation staging. They showed loads and compute adding up instead of overlapping. But removing one part changes how the others overlap, so after a while they stopped being informative.
Every weight product by role and position count. The profile's CUDA events were tagged with the weight's role (ffn_down, dn_qkv, attn_k, the output head, MTP), the forward's position count and the kernel used. This showed the int8 matvec does ~80% of decode weight time, and that each extra draft position costs ~3.5 ms.
In-kernel counters. clock64 reads at phase boundaries, on one block in eight, show where warps spend their cycles. k_mmv: 41% unpacking weights, 28% mma, only 4–5% waiting for memory.
Memory vs arithmetic. Each weight shape timed three ways: its bytes read alone, the product as decode runs it, and the product with the weights served from cache. At 2 positions the arithmetic alone (26.1 ms) already exceeds the memory floor (21.9 ms): the extra positions cost instructions, not bandwidth.
A Grafana dashboard over the same counters, so a real working session shows the picture, not just a benchmark.
The last stretch: what the counters found
Once the counters were in, the fixes were small and obvious, and they took ~10 ms off a forward at 60k (66.1 → 56.6 ms).
Attention block shapes (s25). Attention by position count showed the commonest verify (2 positions) running at half its warps: 12 query rows fill one 16-row tile, so only 2 of 4 warps worked. 6–8 positions ran 3 warps. Reshaping the blocks made attention flat across positions: at 33.7k, n=2 went 11.4 → 6.5 ms.
Quantise each input once (s25). q/k/v, gate/up and DeltaNet's four projections each share an input, but every product re-quantised it. Doing it once: 1.9 → 1.05 ms per forward.
Merge tiny products (s25). DeltaNet's beta and alpha are 48 rows each and ran at 10% of the roof. They sit back to back in VRAM, so they now run as one product.
Tune the output head (s27). The 248k-row head had never been tuned and ran on the default matvec layout at 52% of the roof. Tuned: 3.2 → 2.1 ms per forward, at its floor for one position.
Long-context attention layouts (s28). A real coding session ran to 105k, where attention reached 16–20 ms of a ~70 ms forward and read the K/V cache at only 20–25% of the roof. At head dim 256 each block needed 76–93 KB of shared memory, so one block per SM, and its loads waited for its maths. Loading K, then V, through one buffer halves that, two blocks fit, and one's loads run under the other's maths. Three layouts are timed at startup and the fastest kept per position count: 35% faster for 1–5 positions.
Context
Forward, s23
Forward, s28
Attention, s26
Attention, s28
~33k
61.2 ms
52.1 ms
7.1 ms
5.1 ms
~60k
71.4 ms
56.6 ms
11.5–12.0 ms
8.5–8.9 ms
~75k
—
59–64 ms
~14 ms
10.7–11.5 ms
Where the ceiling is now
At ~10k a forward is ~58 ms, and the weights are the biggest part: 29 ms at 2 positions and 36 ms at 4, against a 22 ms floor for reading them.
The weights are arithmetic-bound past one position. Unpacking 2- and 3-bit IQ weights and the per-position dot products take longer than the memory read. The worst offenders are the most heavily compressed types: IQ1_S runs at twice its floor. Faster loads won't help; fewer instructions per weight will.
MTP drafting costs ~5–7 ms a forward, the output head ~2 ms, attention ~3 ms at 10k.
At long context, attention for 2 positions is 10.3 ms at 96k against a 4.2 ms cache-read floor, and 6–8 positions were still on the old layout. A fourth layout and attention phase counters are being tested now.
Next up: cheaper weight unpacking at 2–4 positions (up to ~7 ms a forward), and closing the long-context attention gap.
Every run, in one picture
The big jumps are early (MTP working, drafting on the GPU, k_mmv); the zig-zag from s16 on is drafting luck around a forward time that moved by 1–2 ms per change at this context. Every run's log and numbers are in the repo under docs/baselines.
There is an old saying in software: eat your own dog food. If you build something, use it yourself, in anger, on the work that actually matters to you. It is the fastest way to find out whether the thing you built is real or only looks real in a demo.
As of this week Marvin is eating mine. Every Qwen request Marvin makes now goes to llmgo, my own Go and CUDA inference engine, and llama.cpp is no longer in that path. In my last post I talked about why I wanted it gone: security, a smaller audit surface, a single binary with nothing underneath it. This post is about what it took to get it into production, and what the numbers look like now that it is there.
Two cards, two jobs
llmgo has been built across two very different GPUs, and that was deliberate.
The RTX 3050 (8 GB) lives in my workstation. It is the development card. It also drives my KDE Plasma desktop, so a slice of that 8 GB is always taken and I can't get it back. The CPU behind it is a Ryzen 7 3700X, AVX2 only. The model is Qwen3.8-27B, and on this card a third of its 64 layers do not fit in VRAM, so they run on the CPU. It is not a friendly place to run a 27B model, and that is the point. If llmgo behaves here, it will behave anywhere.
The RTX 5060 Ti (16 GB) lives in the Kubernetes cluster, with the NVIDIA container toolkit providing the driver. It runs Qwen3.8-27B at UD-Q2_K_XL with its vision projector, all 64 layers on the GPU. This is where Marvin's real work happens, and llmgo now runs there as a container in exactly the slot the llama.cpp pod used to occupy.
Develop on the constrained box, deploy to the bigger one. Every win on the 3050 turns into a bigger win on the 5060 Ti.
The rule: host memory owns the model
The design decision that shaped everything else is simple to say. Host RAM holds every layer of the model, all the time. The GPU only ever gets what it needs right now. Nothing is ever copied back from the GPU; when a GPU copy is no longer needed it is just deleted, because the authoritative version is already sitting in host memory.
That rule applies to every model type, not just the text layers. The vision projector follows the same dynamic loading path, borrowing VRAM from the text layers when they are idle and handing it back when they are needed again. One mechanism, no special cases.
The 3050: ahead of llama.cpp on every number
The first GPU profile on the 3050 was humbling. The kernel doing the decode weight products was taking 263 ms per token for 40 layers, roughly twelve times slower than just reading the weights. A GPU layer was costing as much as a CPU layer, which defeats the purpose of having a GPU.
Over one long day of runs, each with the same three requests (a short prompt, a 9.6k-token chat prompt with tools, and a medium prompt), the things that moved the numbers were:
Prefill attention on the tensor cores, FlashAttention-style tiles instead of FP32 on the CUDA cores. The 9.6k-token prefill went from 55.4 s to 45.0 s, the biggest single step of the day.
An int8 tensor-core matmul for prefill, decoding the 2-bit weights to int8 in shared memory with the tile picked per weight shape at startup. The first version was slower. The fifth, with a two-stage pipeline, was the first run to beat llama.cpp's prefill.
A new decode kernel that works on the low-bit weights directly with int8 activations, tuned per weight shape, running at up to 166 GB/s on a card with 224 GB/s of memory bandwidth.
Speculative decoding with the model's own MTP layer, plus CPU kernels that decode each weight block once for both tokens, so a two-token step costs about the same as a one-token step.
Same model, same prompts, same machine, against llama.cpp:
Then I pointed it at the bigger quant, Qwen3.8-27B UD-Q2_K_XL, the same file the 5060 Ti runs. Getting there meant fast paths for the K-quant types, which had been running on the old float-decoding kernel. Here it is beside the IQ2_XXS results:
To be fair to the numbers: that run had a little more VRAM free than usual, so 44 layers fitted on the GPU instead of 39 or 40, and only 20 ran on the CPU. Per GPU layer the two files are about the same size, so the larger quant costs nothing in placement. There is no llama.cpp run of Q2_K_XL on this card yet, so the grey bars are its IQ2_XXS figures. On an 8 GB card that is also driving a desktop, a 27B model at over 6 tokens per second after a 9.6k-token prompt is a result I am happy with.
The rule I set is that no GPU kernel change is allowed to make the 3050 slower. The 3050 is an Ampere card and the 5060 Ti is Blackwell, and they do not always want the same kernels. If a change helps one and hurts the other, the kernel gets tuned or split per card rather than accepting the regression. The 3050 stays an honest benchmark.
The 5060 Ti: where the dog food got eaten
Marvin talks to its models through an LLM router, and every backend behind that router was llama.cpp. The router probes backends through /props and /health, so llmgo had to answer those the way the router expects before it could even be considered. Then it had to handle tool calling, streaming and thinking (at low reasoning effort, matching the llama.cpp config), with multiple requests isolated from each other, behind TLS and API keys.
Against the llama.cpp deployment it replaced, same card, same model:
Decode is well ahead, and holds at 31.7 to 32.4 tokens per second even at around 35k tokens of context. A big part of that is something llama.cpp leaves on the table. Qwen3.8 ships with a multi-token prediction (MTP) layer that drafts tokens ahead of the main model. llama.cpp ignores it. llmgo uses it, and in Marvin's tool loops 68 to 84 percent of the drafts are accepted, at 2.2 to 2.5 tokens per step.
Prefill on the 5060 Ti is still behind, and I am not going to pretend otherwise. The matrix multiplies are most of the remaining gap, and that is the next target. It has already been closed on the 3050, so I know it can be done.
What only daily use finds
This is the real reason to eat your own dog food. Benchmarks tell you how fast something is. Daily use tells you whether it is right.
The first time I watched Marvin run a long tool loop on llmgo, every turn took about 65 seconds, and 90 percent of that was prefill. Nothing was broken. Each turn just added 100 or so tokens (the tool call and its result) and llmgo re-read all 35,000 tokens of the conversation every time. llama.cpp keeps the previous prompt and only processes what is new. No synthetic benchmark I had written would have shown that, because benchmarks do not have conversations.
So llmgo got a prompt cache, with a twist: Qwen3.8 is a hybrid model, and its DeltaNet layers carry state that has to be checkpointed along the way, not just at the end of the prompt. The result, from a real session: ten turns in a row at 70 to 73 thousand tokens of context all hit the cache, each processing its 105 to 634 new tokens in 1.4 to 2.7 seconds.
Real use kept finding things:
Marvin's UI and gateway router fire two or three small title, tag and follow-up requests after every answer. Each one was pushing the big conversation off the GPU, only for the next turn to pull it straight back. Those small requests now run alongside it instead.
In a long session the cache grew until it squeezed the host, and one 105-token prefill took 66 seconds instead of 1.5. The cache now takes half of what is free beyond a fixed reserve.
When the gateway router changes the tool list mid-conversation (2, then 11, then 29 tools), nothing can be reused, because the tools sit at the very start of the prompt. llama.cpp has the same limit. Good to know it is not just me.
One early build produced gibberish. llmgo now checks itself at startup, running a token through the GPU and the CPU and comparing the results, and refuses to trust a layer that disagrees.
Every one of those came from Marvin doing Marvin's normal work. None of them came from a test suite.
Owning the whole path
It also closes a loop I care about. The same person writing the CUDA kernels is writing the container, the Kubernetes deployment, the router integration and the platform that consumes it. That is the continuum I keep coming back to: not developer, not ops, just the whole path from the silicon to the cluster, owned end to end. When something breaks, there is nobody else to hand it to, and that is exactly how I want it.
What's next
Closing the prefill gap with llama.cpp on the 5060 Ti, starting with the matrix multiplies.
Long-context decode. At 60k tokens, decode drops to 23.6 to 27.6 tok/s, with attention reading the cache far below what the card can do. A fix is built and waiting for its first run on the card.
Faster small cache hits, so adding a short tool result to a 60k-token conversation takes well under a second.
Gemma 4 and a CPU-only dual Xeon server that will serve the same 27B model with no GPU at all.
Back in April I was told I wouldn't be any good in AI because I didn't know GPUs. Marvin now runs on an inference engine I wrote, on kernels I tuned, faster than llama.cpp on every number on my desk and at the part you wait for while it types in the cluster. I like dog food.
