Squeezing vLLM on AMD Strix Halo: 2.3× llama.cpp under load

By Mario CamposPublished
7 min readES
A small glowing cube on a hilltop under a starry sky, with a stream of light spiralling up from it.

vLLM on the Strix Halo iGPU, measured against llama.cpp with the same 35B model: 2.1× faster per token with graphs, 2.3× the throughput with 16 concurrent requests, startup from 9 minutes to 97 seconds, and the one case where llama.cpp still wins.

AMD’s Strix Halo pairs a large integrated GPU (gfx1151) with unified memory the GPU can use almost entirely: with the kernel tuned, PyTorch sees 124 GB of GPU memory. That’s enough to serve a 35-billion-parameter model without a discrete card.

The model here is Qwen3.6-35B-A3B: a mixture of experts with 3 billion parameters active per token and a hybrid architecture that mixes normal attention with Gated DeltaNet, a linear attention with a recurrent state. llama.cpp runs it well. The question was what vLLM adds on the same hardware, and what it takes to get there.

The short version, with the same model and the same prompts:

  • 2.3× llama.cpp’s throughput with 16 concurrent requests (2.07 vs 0.89 requests per second).
  • 2.1× faster per generated token with graphs instead of eager mode (39 vs 81 ms).
  • Startup from 9 minutes to 97 seconds.
  • Identical answers for identical requests.
  • And one honest exception: on a single long request, llama.cpp is still a little faster.

Four levers got there. None of them touches the model or its kernels.

Lever 1: an autotuning loop that ran for hours

On first launch, vLLM sat in its warm-up with the GPU at 100 % and no log output. py-spy dump (through podman unshare, since ptrace isn’t allowed inside a rootless container) showed where: _warmup_prefill_kernels → Triton’s autotuner → triton.testing.do_bench.

Why vLLM looked frozen 1. Estimate: Time 5 launches, → ≈ 0 ms on this iGPU → 2. Repetitions: rep ÷ estimate, → hundreds of thousands → 3. Each repetition: Clear a 256 MiB buffer,, then run the kernel. × hundreds of thousands, per autotune config. The fix: 3 warm-up runs + 10 timed runs → autotuning in seconds. Why vLLM looked frozen Triton's do_bench on a kernel that finishes almost instantly 1. Estimate Time 5 launches → ≈ 0 ms on this iGPU 2. Repetitions rep ÷ estimate → hundreds of thousands 3. Each repetition Clear a 256 MiB buffer, then run the kernel × hundreds of thousands, per autotune config The fix: 3 warm-up runs + 10 timed runs → autotuning in seconds
Why vLLM looked frozen 1. Estimate: Time 5 launches, → ≈ 0 ms on this iGPU → 2. Repetitions: rep ÷ estimate, → hundreds of thousands → 3. Each repetition: Clear a 256 MiB buffer,, then run the kernel. × hundreds of thousands, per autotune config. The fix: 3 warm-up runs + 10 timed runs → autotuning in seconds. Why vLLM looked frozen Triton's do_bench on a kernel that finishes almost instantly 1. Estimate Time 5 launches → ≈ 0 ms on this iGPU 2. Repetitions rep ÷ estimate → hundreds of thousands 3. Each repetition Clear a 256 MiB buffer, then run the kernel × hundreds of thousands, per autotune config The fix: 3 warm-up runs + 10 timed runs → autotuning in seconds

do_bench decides how many times to run a kernel by timing five launches and dividing its budget by that estimate, and it clears a 256 MiB buffer before every run. The warm-up autotunes tiny kernels whose estimated time on this iGPU is practically zero, so the division gives hundreds of thousands of repetitions for every configuration being compared. It would have finished, in hours.

Replacing do_bench from a sitecustomize.py in the image, which Python loads before anything else, fixes it:

# Triton's do_bench computes repetitions as rep / estimate. For tiny kernels
# the estimate is ~0 ms on gfx1151, so it runs hundreds of thousands of times.
# Three warm-ups and ten timed runs are enough to rank autotune configs.
import triton.testing as tt

def do_bench(fn, warmup=25, rep=100, grad_to_none=None, quantiles=None, return_mode="mean", **_):
    import torch
    for _ in range(3):
        fn()
    torch.cuda.synchronize()
    times = []
    for _ in range(10):
        start, end = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
        start.record(); fn(); end.record(); torch.cuda.synchronize()
        times.append(start.elapsed_time(end))
    return tt._summarize_statistics(times, quantiles, return_mode)

tt.do_bench = do_bench

Ten timed runs are plenty to rank autotuning configurations. The same autotuning now takes seconds (one of the kernels, merge_16x16_to_64x64_inverse_kernel, in 4.07 s).

Lever 2: stop redoing work on every start

With autotuning fixed, the first start took 9 minutes and later ones 5 to 9. Two costs were paid on every start and thrown away.

  • aiter compiled its core module every time. AMD’s kernel library builds module_aiter_core with hipcc the first time it’s imported: about 46 seconds, inside the container, lost when it stopped. It’s now compiled once, at image build time. The catch: aiter detects the GPU architecture by querying the device and ignores GPU_ARCHS, so the build needs the GPU (podman build --device /dev/kfd --device /dev/dri). 22 seconds, once.
  • Triton cached kernels, not verdicts. Compiled kernels were cached, but every autotuning measurement was repeated on each start. TRITON_CACHE_AUTOTUNING=1 stores which configuration won. The caches live outside the container: 239 MB for Triton and 642 MB for torch.compile.

vLLM startup time Horizontal bars on a linear scale starting at zero: before, 540 s (9 min); 5–9 min, 9 min the first time. After, 97 s. Expanded breakdown of the 97 s: weights 33 s; compile + autotune + profiling ~15 s; graph capture ~15 s; rest ~34 s. vLLM startup time Qwen3.6-35B-A3B AWQ 4-bit, from launch to ready Scale: 0–540 s Before: every start 9 min 5–9 min, 9 min the first time After: aiter prebuilt + Triton autotune cached 97 s 97 s breakdown (expanded) 1 2 3 4 1. weights 33 s 2. compile + autotune + profiling ~15 s 3. graph capture ~15 s 4. rest ~34 s Measured on an AMD Strix Halo (gfx1151)
vLLM startup time Horizontal bars on a linear scale starting at zero: before, 540 s (9 min); 5–9 min, 9 min the first time. After, 97 s. Expanded breakdown of the 97 s: weights 33 s; compile + autotune + profiling ~15 s; graph capture ~15 s; rest ~34 s. vLLM startup time Qwen3.6-35B-A3B AWQ 4-bit, from launch to ready Scale: 0–540 s Before: every start 9 min 5–9 min, 9 min the first time After: aiter prebuilt + Triton autotune cached 97 s 97 s breakdown weights 33 s compile + autotune + profiling ~15 s graph capture ~15 s rest ~34 s Measured on an AMD Strix Halo (gfx1151)

Ready at 97 seconds: weights in 33 s, compilation, autotuning and profiling in about 15 s, graph capture in about 15 s. The final image is 14.7 GB, against 61.8 GB for AMD’s official one.

Lever 3: graphs instead of eager, 2.1× per token

The usual advice for vLLM on iGPUs is --enforce-eager. In eager mode every generated token cost 81 ms, five times what llama.cpp needs on the same hardware: Python dispatches every kernel of every layer for every token, and with only 3 billion active parameters that overhead is most of the work.

Without --enforce-eager, vLLM compiles the model (44 s the first time, then cached) and captures seven piecewise graphs, so each step replays a recorded graph instead of launching every kernel one by one.

Time per generated token Bars share a linear scale starting at zero. Eager (--enforce-eager): 81 ms; CUDA/HIP graphs (piecewise): 39 ms. 2.1× faster. Time per generated token vLLM, same model; lower is better Scale: 0–90 ms Eager (--enforce-eager) 81 ms CUDA/HIP graphs (piecewise) 39 ms 2.1× faster Measured on an AMD Strix Halo (gfx1151)
Time per generated token Bars share a linear scale starting at zero. Eager (--enforce-eager): 81 ms; CUDA/HIP graphs (piecewise): 39 ms. 2.1× faster. Time per generated token vLLM, same model; lower is better Scale: 0–90 ms Eager (--enforce-eager) 81 ms CUDA/HIP graphs (piecewise) 39 ms 2.1× faster Measured on an AMD Strix Halo (gfx1151)

39 ms per token. The answers on my evaluation set didn’t change, and a repeated request still came back bit for bit identical.

Lever 4: bigger batches, 2.3× llama.cpp under load

Throughput under concurrency is where vLLM earns its place. llama.cpp serves a fixed number of slots and, on Vulkan, doesn’t batch the prefill of several sequences together. vLLM schedules every running request into each forward pass.

The last lever was --max-num-batched-tokens: up to 8,192 tokens per step instead of 2,048. With 8 concurrent requests throughput went from 0.98 to 1.18 requests per second; with 16, from 1.77 to 2.07. Latency for a single request didn’t get worse.

Throughput with 16 concurrent requests Bars share a linear scale starting at zero. llama.cpp (Q4, 4 slots): 0.89 req/s; vLLM, eager: 1.56 req/s; vLLM, graphs: 1.77 req/s; vLLM, graphs + 8,192-token batches: 2.07 req/s. 2.3× llama.cpp. Throughput with 16 concurrent requests Requests per second; higher is better Scale: 0–2.2 requests/s llama.cpp (Q4, 4 slots) 0.89 req/s vLLM, eager 1.56 req/s vLLM, graphs 1.77 req/s vLLM, graphs + 8,192-token batches 2.07 req/s 2.3× llama.cpp Measured on an AMD Strix Halo (gfx1151)
Throughput with 16 concurrent requests Bars share a linear scale starting at zero. llama.cpp (Q4, 4 slots): 0.89 req/s; vLLM, eager: 1.56 req/s; vLLM, graphs: 1.77 req/s; vLLM, graphs + 8,192-token batches: 2.07 req/s. 2.3× llama.cpp. Throughput with 16 concurrent requests Requests per second; higher is better Scale: 0–2.2 requests/s llama.cpp (Q4, 4 slots) 0.89 req/s vLLM, eager 1.56 req/s vLLM, graphs 1.77 req/s vLLM, graphs + 8,192-token batches 2.07 req/s 2.3× llama.cpp Measured on an AMD Strix Halo (gfx1151)

2.07 requests per second with 16 at a time: 2.3× llama.cpp with the same model (llama.cpp with a Q4 GGUF and 4 slots). Every lever adds up: eager vLLM already beats llama.cpp under load, graphs add 13 %, and bigger batches another 17 %.

One request at a time: where llama.cpp still wins

One request at a time Grouped bars share a linear scale from 0 to 2,000 ms. llama.cpp (Q4, 4 slots) first; vLLM (AWQ, graphs, 8,192 batch) second. Short request: 876 ms / 629 ms; Medium: 1,377 ms / 1,353 ms; Long: 1,748 ms / 1,868 ms. One request at a time Latency in ms by request size; lower is better llama.cpp (Q4, 4 slots) vLLM (AWQ, graphs, 8,192 batch) Scale: 0–2,000 ms Short request llama.cpp 876 ms vLLM 629 ms Medium llama.cpp 1,377 ms vLLM 1,353 ms Long llama.cpp 1,748 ms vLLM 1,868 ms Measured on an AMD Strix Halo (gfx1151)
One request at a time Grouped bars share a linear scale from 0 to 2,000 ms. llama.cpp (Q4, 4 slots) first; vLLM (AWQ, graphs, 8,192 batch) second. Short request: 876 ms / 629 ms; Medium: 1,377 ms / 1,353 ms; Long: 1,748 ms / 1,868 ms. One request at a time Latency in ms by request size; lower is better llama.cpp (Q4, 4 slots) vLLM (AWQ, graphs, 8,192 batch) Scale: 0–2,000 ms Short request llama.cpp 876 ms vLLM 629 ms Medium llama.cpp 1,377 ms vLLM 1,353 ms Long llama.cpp 1,748 ms vLLM 1,868 ms Measured on an AMD Strix Halo (gfx1151)

With a single request, vLLM wins on short prompts, they tie on medium ones, and llama.cpp is still a little faster on long ones. Alone, the iGPU is bound by memory bandwidth, and llama.cpp’s Q4 kernels are tuned precisely for that. Under concurrency, batching wins by a wide margin. If your traffic is one user at a time with long prompts, llama.cpp is still a great choice; if it’s many users, vLLM is.

Same request, same answer

The number that doesn’t fit in a throughput chart: in vLLM, the same request repeated gives the same output, bit for bit. In llama.cpp, with this family of hybrid models, the same prompt at temperature 0 moved by up to 0.44 in log-probability between runs. For anything that reads log-probabilities, or that must give the same answer twice, that alone decides it. Why llama.cpp drifts deserves a post of its own.

What didn’t pay off

  • MTP (speculative decoding with the model’s own draft head). 76 % of drafted tokens were accepted with complete log-probabilities, but prefix-cache hits dropped to zero, this build lacks the fused kernel for verifying the draft so that step falls back to Triton, and with 16 concurrent requests throughput fell to 0.80 per second. Not yet.
  • More prefix-cache reuse through tuning. In a hybrid model vLLM sizes attention cache pages to match the recurrent state, which fixes attention blocks at 1,056 tokens. Prompts of ~1,200 tokens rarely reuse more than one block: 28 % of tokens hit. --mamba-block-size doesn’t change it.

Takeaways

  • Profile a stall before blaming the hardware. A process at 100 % with no logs is busy. py-spy dump says with what in one command.
  • Give timing loops a floor. Anything that divides a budget by a measured time explodes when the measurement is almost zero.
  • Cache the verdicts, not only the artifacts. Autotuning results and aiter’s build were tens of seconds on every start.
  • Try graphs, even where eager is the advice. Here they halved the time per token.
  • Benchmark concurrency and single requests separately. They crown different winners.
  • Batch size is a throughput knob. 8,192 tokens per step was a free 17 % under load.

With four levers and no changes to the model, vLLM on the Strix Halo went from not finishing its warm-up to 2.3 times llama.cpp’s throughput, twice as fast per token, and ready in 97 seconds.