One level down: #003 what CUDA has that Vulkan lacks

The previous article ended on a question: what CUDA could teach us that we cannot do in Vulkan. I had written that Vulkan lacks two instructions, without naming them. Five of the seven reviews of that article picked it up. Here is the answer: it is not missing two, it is missing one.

To measure them, I needed a Herbert that speaks CUDA. It writes its kernels in PTX, NVIDIA's intermediate language, at startup, for the shape of the model and for the card. It loads the driver at run time. The driver contains its own assembler, the equivalent of the CUDA toolkit's ptxas: it is the driver that translates this PTX into machine code for the card, allocates the registers, then keeps the result in its cache. ptxas itself only serves to check the modules offline. No nvcc, no NVRTC, no cuBLAS: the same principle as on the CPU and under Vulkan, with the same numerical contract (weights, activations and cache in bf16, accumulators in f32).

In short

• What CUDA has that Vulkan lacks comes down, at the machine-code level, to one instruction: cp.async. Removed from the CUDA engine, it costs 5 to 14% of prefill. The other one, ldmatrix, the Vulkan driver already emits on its own. At decode, neither matters.
• On an RTX 4080, one request at a time, Herbert on CUDA decodes faster than TensorRT-LLM, NVIDIA's own engine, on all three models measured (by 2 to 7%). It is also faster than vLLM, SGLang and llama.cpp. At decode, the lead holds after a 40,000-token prompt. At prefill it is level within a few percent on long prompts, both ways. On the 2B, Herbert is behind only for lack of a measured setting for that model.
• I had predicted I would get "wrecked" on very long prompts. That was true: at 60,000 tokens, Herbert decoded at half the speed of vLLM. I found the cause and fixed it. It is now ahead at decode at that length, and level at prefill.

Two instructions

On an NVIDIA GPU, a well-written matrix multiplication does two things at once: it computes on one tile while it loads the next. Two PTX instructions serve that purpose.

cp.async copies a tile from the card's memory into the block's shared memory without going through registers and without waiting: computation goes on during the copy. ldmatrix loads in one instruction, from shared memory, the matrix fragments the tensor cores need, in the layout they expect.

Here is what Herbert writes at startup for the GEMM of the 0.8B's gate_up projection (file gemm_1024x7168_bf16_64x128x32_w2x2_s2.ptx, raw lines, cut at the ...): the asynchronous copy of the next tile, then the computation on the current one, ldmatrix then mma on the tensor cores.

	cp.async.wait_group 0;
	bar.sync 0;
	cp.async.cg.shared.global [%r83], [%rd9], 16, %r16;
	cp.async.cg.shared.global [%r84], [%rd13], 16, %r25;
	...
	cp.async.commit_group;
	ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r91, %r92, %r93, %r94}, [%r89];
	ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%r99, %r100, %r101, %r102}, [%r90];
	mma.sync.aligned.m16n8k16.row.col.f32.bf16.bf16.f32 {%f0, %f1, %f2, %f3}, {%r91, %r92, %r93, %r94}, {%r99, %r100}, {%f0, %f1, %f2, %f3};
	...

The rest of the kernel is in the engine's README.

Under Vulkan, SPIR-V has neither. Asynchronous copy exists in SPIR-V, but it relies on types reserved for OpenCL kernels, forbidden in a Vulkan shader. ldmatrix has no explicit equivalent: cooperative matrices let the driver choose how to load. That is what I had in mind when I wrote the previous article.

The machine code proved me half wrong. I did not see that coming. I read the SASS the card actually runs for Herbert's Vulkan kernels, in the driver's disk cache. The cooperative GEMM does contain LDSM: the driver emits ldmatrix on its own. It contains no LDGSTS, the instruction behind cp.async: every tile goes through registers. At the machine-code level, what CUDA has that Vulkan lacks is one instruction.

What it is worth

It remained to know what it is worth. I measured it the other way round: in the complete CUDA engine, I removed one setting at a time, with the same binary, on the 4080, alternating the variants over five rounds. Prefill from 128 to 4,000 tokens, on the 0.8B and the 4B. The tokens produced are identical in every variant: only the speed changes.

cp.async weighs the most. It is precisely the one that is missing. At decode, no GEMM setting matters: only the CUDA graph, which records the step of one token once and for all, weighs 4.7% on the 0.8B and 1.1% on the 4B. Vulkan has the equivalent of the graph, with its recorded command buffers.

One detail: this measurement ran on October 18 with the card in P2, the state the driver keeps it in under CUDA (more on that below). These are gaps between variants of the same engine in the same state, not speeds to compare with other engines.

There is another lever specific to CUDA, which I did not take: accumulating in f16. The previous article already said it, the format alone changes almost nothing. Under CUDA too, Herbert in f16 runs exactly at the speed of bf16, within 0.5%. The only gain comes from accumulating the products in f16 instead of f32: 6 to 21% at prefill. But that is another numerical contract: on an isolated GEMM, 14% of the outputs change. I did not take it.

My predictions, then the measurements

Before the campaign, on the morning of October 18, I dictated my predictions. They are published as they are with the data, in a dated file. They are not blind: I knew the first measurements of Herbert on CUDA against an older version of vLLM. But I knew neither the new versions of the engines nor their best settings. I saw myself behind them, everywhere. The argument was simple: these engines have existed for years, written by extremely skilled engineers. I saw no chance that my code would be at least as fast.

The bench: an RTX 4080, driver 595.99.02, locked clocks (core at 2,805 MHz, memory at 11,201 MHz), Qwen3.5-0.8B and 4B in bf16, one request at a time. Against it: llama.cpp b11718, vLLM 0.31.0, SGLang 0.5.22 and transformers 5.20.1, each at its best setting, searched for before measuring. Every CUDA engine in P0. Two instruments, as in the previous article: each engine's own bench and a common HTTP client.

To my great surprise, it is the opposite, almost everywhere.

The common HTTP client says the same, with a sometimes different order. At a 4,096-token prefill on the 4B, SGLang, Herbert and vLLM stand within 5% of one another.

I am not one to show off: I have to say where a good part of the lead comes from. The bare engine, without the automatic choice of GEMM tiles, is 15 to 28% behind vLLM and SGLang at 4B prefill. The tables give the engine with that choice, which has been its default since October 19, in the published version.

Against NVIDIA's own engine

One engine was missing from this bench: TensorRT-LLM, the one NVIDIA develops for its own cards. I added it on the evening of October 24, for this article, along with a third model, the 2B. Same card, same locked clocks, everything in bf16. Herbert at its default settings: the engine measured emits exactly the same PTX as the published version.

I measured prefill from 128 to 40,000 tokens, then decode after prompts of 759 to 40,000 tokens. The measurement had just started when I dictated one more prediction: "level at decode, maybe a bit better, but worse with longer pp ... and a bit worse at prefill, all the time ... it's NVIDIA's code, even if I imagine they care more about the H100, B200 etc ... I think they know better than me how to use the power of their hardware".

These figures do not read quite like those of the previous section. TensorRT-LLM, vLLM and SGLang each ran in a single pass, with their own repetitions, after the alternated rounds of Herbert and llama.cpp. Their prompts are random tokens, Herbert's a real text. For TensorRT-LLM and vLLM, prefill is the latency of a whole request that generates one token, scheduling included: at 128 tokens that fixed cost weighs heavily, and that is mostly what the ×3 on the 0.8B says. From 2,000 to 40,000 tokens, where computation dominates, the gaps are a few percent, both ways. On the 2B, TensorRT-LLM, SGLang and vLLM keep the edge. On these long prompts the card reaches its 320 W limit and its clock drops at times to 2,625 MHz; vLLM splits the prompt into 8,192-token chunks, and TensorRT-LLM has to do the same on the 4B at 40,000 tokens.

At decode, on the other hand, the measurement leaves no doubt. There too, I was surprised: on an NVIDIA card, without any NVIDIA library, Herbert runs faster than NVIDIA's engine.

All things considered: the engines I compare myself with are built to serve many users at once; this bench only measures one request at a time. I have not measured the data center cards their teams work on. I plan to do it with RunPod, and I think the balance will be restored there. But for me, this is an important moment. I did not expect to reach such a result so soon.

That left the 2B, the only model where Herbert is behind at prefill. The reason is in the code: the automatic choice of GEMM tiles was measured only for the widths of the 0.8B and the 4B. The 2B, which has none, falls back to a basic setting, with smaller blocks. With the tile of the 0.8B and the 4B and blocks of 2,048 rows, set through the options of the published version, alternating the variants over five rounds, its prefill goes from 25,684 to 29,056 tokens per second at 2,000 tokens, then from 24,524 to 27,706 at 10,000: 5 and 6% ahead of TensorRT-LLM. The tokens produced are the same. This setting is not the default yet; it will be in a coming version. To replay it:

--tune gemm.auto=0 --tune gemm.tile=64x128w2x2 --tune gemm.bk=32 --tune gemm.stages=2 --tune prefill.block=2048

The long prefill

On October 19, before measuring prompts from 10,000 to 60,000 tokens, I dictated one more prediction: "I think I'm going to get wrecked on long prefills. And at decode too, by the way."

I was right. At 60,000 tokens on the 0.8B, Herbert did 0.84 times SGLang and 0.87 times vLLM at prefill. At the decode that follows a prompt of that length, it was worse: 0.49 times vLLM. From 10,000 to 60,000 tokens, Herbert lost 55% of its decode speed, vLLM and llama.cpp about 20%.

The profile pointed to a single kernel, the decode attention. Between 10,000 and 60,000 tokens, it went from 22 to 62% of the time of each token; the rest did not move. Reading the code, the cause was simple: this kernel split the work into 64 blocks whatever the context length, on a card that has 76 multiprocessors. Each block went through its share of the positions alone. And the four heads that share the same K and V, in grouped attention, each read them again. The kernel held 149 GB/s, a fifth of what the card's memory can deliver. I hated myself for letting that slip through: reading each piece of data only once is the very heart of my approach.

The fix took one night. One block per K/V head and per slice, K and V read once for the four heads, 64 slices instead of 8. At prefill, the copy of the next tile starts during the computation of the previous one. On the 0.8B at 60,000 tokens, decode goes from 133 to 278 tokens per second, and prefill gains 17%.

Then the measurement against the others, in the same session, alternated, five rounds, at 60,000 tokens:

At decode, Herbert is now ahead at that length. At prefill it is level, within 4%, with llama.cpp and vLLM still ahead on the 4B, SGLang on the 0.8B. These fixes change the order of some sums: on one 0.8B prompt, the token sequence now departs from the reference at the 172nd token, where it used to be identical; on the 4B, another one becomes identical. These gaps fall on near-ties between two tokens, as in the previous article. The fix became the default setting on October 21. The former attention remains available through three options, to replay the measurements from before.

The memory clock

The previous article left a question open: under CUDA, the 4080's memory ran at 10,801 MHz instead of 11,201. I did not explain it.

The driver does it on purpose. As soon as a program opens a CUDA context, it keeps the card in the P2 state, one notch below its maximum, on Linux as on Windows. The environment variable you find everywhere to avoid it is actually meant for the opposite: it turns the boost off. On my machine, it changed nothing. What works is an application profile that NVIDIA ships with its drivers, CudaNoStablePerfLimit. With it, the card stays in P0 and 4B decode gains 3.6%, the ratio of the two clocks. The whole campaign of this article ran in P0, for the sake of fairness, for every CUDA engine, not just for Herbert. Without the profile, the engines measured on October 18 (Herbert, vLLM, SGLang, llama.cpp and transformers) all lose 2.7 to 3.7% at decode. TensorRT-LLM, added later, only ran in P0: the clock log confirms it second by second during its measurements, but I did not measure it in P2.

What I do not explain yet

Prefill beyond 20,000 tokens: the prefill attention takes 62% of the time there, with a single block of four warps per multiprocessor, 8% of what the card could run at once. I have not yet found how to fill it without changing the results.

The decode of Herbert on Vulkan on the 0.8B with a long context, which already fell behind in the previous article: the cause found under CUDA is probably there too. I have not measured it yet.

The limits

One card, the 4080. Three dense models, in bf16. One request at a time. No B200, no H100, no FP8, no quantization. vLLM and SGLang are built to serve many requests at once: this bench does not measure what they do best.

These figures hold from October 18 to 25, 2026, with these versions and this driver (TensorRT-LLM 1.4.0rc1). Engines evolve: some of these figures may no longer be true when you read this.

Reproducing at home

Two repositories are online, under the MIT license. The engine: github.com/xigh/herbert-cuda-bf16, tag v0.1.0. It builds without the CUDA SDK: all it needs is the NVIDIA driver.

git clone --branch v0.1.0 https://github.com/xigh/herbert-cuda-bf16.git
cd herbert-cuda-bf16 && cargo build --release --locked
target/release/herbert-cuda-chat ~/models/Qwen3.5-4B --no-think --ask 'What is the capital of France?'

The default setting is the fixed attention. To replay the measurements from before the fix, add --tune sdpa.gqa=0 --tune sdpa.split=8 --tune sdpa.pipe=0.

The campaign scripts, as they ran, my predictions, the protocol and the raw data: github.com/xigh/herbert-cuda-bench, tag v1-2026-10-25c. If a figure looks wrong to you, you have what you need to check it.

What comes next

The next article moves to Metal, on Apple's chips. The same question, once more: does the same principle hold?