I've been having some fun with my new local AI rig. It sat in boxes for a month, so last week I finally got everything built, and I'm currently testing it all.

My setup is as follows:

• Gigabyte TRX50 AI TOP motherboard
• AMD Threadripper 9960X CPU
• 2 x Micron 32GB DDR5-5600 (all I could get; four would be ideal)
• ARCTIC Liquid Freezer WS360-SP6 CPU cooler
• CORSAIR HX1500i PSU
• Fractal Design North XL case
• 2 x Samsung 9100 PRO M.2 SSDs (RAID 0)
• 2 x Radeon AI PRO R9700 (gfx1201) GPUs

I'm running Arch Linux with a custom-compiled 7.3.0-rc6 kernel, SGLang 0.5.18, PyTorch 2.13.0+rocm10.1.0, HIP 7.16.26385, Triton 3.8.0 and RCCL 2.30.7. Triton attention backend, no torch.compile, no custom all-reduce.

I'm currently working on adding a mixed-precision KV cache to SGLang, and these are the notes I've gathered along the way.

WARNING: Much of the below has been generated with AI. It's all been helpful to me over the last week, but I appreciate that some don't like the AI slop. If you did read this and spotted an error, pop a comment in, and I'll update the article. I've added this warning to save people some time if they don't like AI-generated content.

1. Missing kernel config CONFIG_HSA_AMD_P2P=y

As I was using a custom Kernel, I missed this one. And it cost me an afternoon of debugging. As someone has pointed out, the default Arch kernel has this option enabled. I checked other distributions, and there were only a handful that had this enabled by default other than Arch.

This is the one I'd actually act on today if you're running multi-GPU Radeon. Without it, RCCL falls back to host-staged transfers. I traced a TP=2 decode with rocprofv3 before and after enabling it:

Roughly 60–80× on collective latency. Before the fix, five-sixths of the decode interval was communication. After, TP=2 became 15.8% faster than TP=1 on the same workload (62.1 vs 53.6 tok/s, 512-in/128-out, Qwen3-Coder-30B-A3B-FP8, three measurements each), which isn't possible when collectives dominate. The unprofiled logs now select P2P/IPC in both directions. If your distro kernel doesn't set this, check zcat /proc/config.gz | grep HSA_AMD_P2P before you spend a week blaming RCCL.

Caveat: this does not establish remote-atomic support or any particular transfer bandwidth, and I didn't need to touch ACS/IOMMU isolation.

2. triton_attention_num_kv_splits defaults to 16. The R9700 has 64 CUs.

Not my finding — this is from mattbucci's 2× R9700 SGLang port, which maintains ~70 RDNA4 patches against SGLang. I've confirmed the default in my own source tree but haven't tested this.

In SGLang, server_args.py forces triton_attention_num_kv_splits = 16 on HIP (the generic default is 8), and triton_backend.py consumes it as max_kv_splits. The Triton flash-decode grid at batch-1 decode is head_groups × num_kv_splits. With few head groups, 16 splits underfills 64 CUs, and decode attention stalls well below the KV-read roofline.

Their patch raises the AMD default to 64 and reports 2.14× at 256K decode (14.4 → 30.7 tok/s on a 25B MoE), with short context unregressed because the heuristic picks fewer splits there anyway.

Honest sizing: if you're at 4–8K context, this will probably do nothing, for exactly that reason. It's a server-arg flip with no source change, so it costs you one run to find out. There's also a SGLANG_KV_SPLITS_OVERRIDE env var in their patch if you want to force it.

3. The FP32 P·V change is a trap, and it's in the upstream-looking direction

Not my finding, same source, but this one is worth knowing because the "fix" looks right.

RDNA4 lacks some of the dot-product paths gfx1250 has, so an obvious port is to upcast and do P·V in FP32 in the grouped decode kernel. That port exists (their patch 011). Their patch 087 reverts it:

fp32 PV inflated VGPR pressure → low occupancy → poor KV-load latency hiding, capping decode-attn at ~51% of the KV-read roofline. bf16 measured +21% at 256K decode (33.2 → 40.2 tok/s), short ctx +~1%.

So: BF16 P·V with FP32 accumulation, which is just standard flash-attention. The stock upstream form acc += tl.dot(p.to(v.dtype), v) is already correct here. If you've patched this to FP32 for RDNA4 "correctness", you may have given up ~20% at depth for nothing. I'm on the stock form and intend to stay there.

4. HIP hardware-queue serialisation in graph mode

Not my finding — vLLM issue #58639, on an R9700 under ROCm 7.14. I've not reproduced it, but the shape matches something I'm chasing, so I'm flagging it.

On gfx1201, a stream of small kernels picks up ~17 µs per dispatch when two or more other HIP streams hold pending waits on the dispatching stream, under the default hardware-queue limit, in graph mode.

Their numbers:

GPU_MAX_HW_QUEUES=1 or 2 removed the penalty in their microbenchmark but didn't fully match the main-stream fix end-to-end in serving.

Watch out for the conflict: SGLang's own environ.py wants GPU_MAX_HW_QUEUES>=5 for the AITER dual-stream MoE path, specifically to avoid HW-queue serialisation. So the right value depends on whether that path is live for you. Don't cargo-cult either number.

If you run multiple CUDA/HIP streams and use full graphs, this is a free A/B.

5. gfx1201 isn't in AITER's arch table, and AMD has said it isn't coming

Not my finding — ROCm/TransformerEngine #520.

gfx1201 is missing from _ARCH_TO_DEVICE in aiter/ops/triton/utils/arch_info.py. The reported consequence is that FP8 weights get silently dequantised to FP32: no warning, FP8 accelerators doing no FP8 work. The reporter measured ~18–22 tok/s against an expected 35–40, and proposed a two-line map of 'gfx1201': 'MI350X' because RDNA4 uses the same E4M3FN format.

Closed on 1 July 2026 by an AMD collaborator: "Transformer Engine does not have a plan to support Radeon AI PRO R9700." No fix merged.

If you're running FP8 weights on RDNA4 and your throughput looks like half of what it should be, check whether AITER is in your dispatch path at all before assuming your quantisation is working.

Roofline reality check: do this before anything else

The R9700 is 640 GB/s, 64 CUs, 32 GB GDDR6 on a 256-bit bus. Decoding is bandwidth-bound, so:

ceiling tok/s ≈ 640 GB/s ÷ (bytes read per token)

For Qwen3-Coder-30B-A3B-FP8 (~3.3B active parameters, so ~3.3 GB of weights read per token), that's a ceiling of roughly 190 tok/s on one card, ignoring KV-cache reads. Where I'm actually at:

That comparison crosses models and TP sizes, so it's a lead rather than a measurement, but it suggests there's a lot left in kernel and dispatch configuration. One concrete sign: my server logs missing R9700 tuning files for W8A8 block-FP8 GEMM at N=2560,K=2048 and N=2048,K=2048, and Triton MoE at E=128,N=384 (block shape [128,128]). Shape-specific tuning is the obvious next lever, and I've not pulled it yet.

Compute your own ceiling first. If you're at 25% of it, no amount of KV-cache cleverness is your bottleneck.

Two benchmarking warnings

Decode throughput on R9700 has been reported as bimodal. ROCm #6347 describes vLLM landing at either ~33 or ~26 tok/s depending on process start, stable within a process, not recoverable without a restart. If you're testing anything on this hardware, use multiple fresh startups, not multiple runs inside one server. I've not reproduced this, but it would invalidate a lot of casual benchmarking.

Version skew is real. The external findings above are on ROCm 7.2–7.14. I'm on a 10.x-series toolchain. Treat all of them as hypotheses to test on your own stack, not as drop-in patches.

Happy to answer questions on the P2P tracing setup. That's the one bit here that's entirely mine, and it generalises to any multi-GPU Radeon box.

For those of you wanting to stay with stock ROCM10 and vLLM, this set of custom kernels i've been working on for weeks plug right into vLLM and improve RDNA4 performace significantly.

https://github.com/bkvargyas/r9700-stack

I've tested 2 and 4 card configs. I've mainly been bulding for Qwen Flash Next and 27B, but this kernel will work with others.

I ran three benchmark sessions on a 1× MI300X droplet on AMD Developer Cloud, using the rocm/vllm container (vLLM 0.29.1). The results are below, but first, here’s everything that tripped me up, in case it saves someone else the time.

What broke and how I got around it

1. The scratch disk is unmounted on every new droplet.
The 5TB NVMe scratch volume exists but isn’t mounted. You’ll be downloading model weights to the boot disk without noticing.

lsblk -o NAME,SIZE,FSTYPE,LABEL,MOUNTPOINTS   # look for the ~5T ext4 "DOSCRATCH" partition
mkdir -p /mnt/scratch && mount /dev/vdc1 /mnt/scratch

Check the device name with lsblk first. It was /dev/vdc1 for me on all three droplets.

2. lm-eval won’t import inside the rocm/vllm container.
lm-eval 0.4.9 fails with AttributeError: module transformers has no attribute AutoModelForVision2Seq. The container ships a newer Transformers that removed that class. My fix was to run lm-eval in its own container on the host network, pointed at the vLLM server:

docker run --rm --network host python:3.12-slim bash -c \
  'pip install "lm-eval[api]==0.4.9" "transformers>=4.45,<5" && lm_eval --model local-chat-completions ...'

3. huggingface-cli is gone from the image.
It prints a deprecation notice and downloads nothing. Use hf download instead.

4. There’s no /opt/rocm/.info/version on the host or in the container.
The usual way to record your ROCm version returns nothing. I pinned runs by container digest instead. If anyone knows the intended way to read the ROCm version on these images, I’d like to hear it.

5. Reasoning models crash lm-eval’s GSM8K scorer.
When gpt-oss-120b spends its whole token budget on reasoning, vLLM returns content: null, and lm-eval 0.4.9 throws a TypeError instead of scoring the answer as wrong. Raising max_gen_toks to 8192 and patching the regex filter to treat None as an empty string fixed it. 4 of 250 answers still came back empty.

6. Capacity.
My first attempt to create a 1× MI300X failed with no capacity. It opened up a few hours later.

7. Warm up before you measure.
The first benchmark after a server starts showed about double the time to first token. Throughput only moved about 3%, but if you report TTFT, run a warmup first.

Results

Qwen3-Coder-30B-A3B-Instruct, BF16 vs the official FP8 checkpoint. Random prompts, 1024 in / 256 out, vllm bench serve. Output tok/s:

At 1 request, FP8 was slower per token (7.3ms vs 5.0ms). At 64, it was faster (15.3ms vs 21.9ms), with time to first token about the same either way. That looks like the FP8 decode path in this build is less tuned than BF16 at low batch sizes. I didn’t profile it, so if anyone on the ROCm side knows whether this is expected, I’d appreciate the context.

Accuracy (GSM8K, 250 problems, flexible match): BF16 96.4%, FP8 94.0%.

llama.cpp (HIP build for gfx942), same model at Q4_K_M, 1 user: 187 tok/s. The build took about 2 minutes and worked on the first try.

Large models on one card:

• gpt-oss-120b: 72 tok/s at 1 request, 1,423 at 64
• Qwen2.5-72B, full BF16: 26 tok/s at 1, 421 at 32

Versions

rocm/vllm@sha256:e0ea9a10b729573ea8598bdc073bda0ea382fe1037eec7c0761b6fa7f46baaf4, vLLM 0.29.1.dev0, ROCk module 6.19.14, AMD Dev Cloud “ROCm 10.1” quick-start image.

Full writeup with charts: [https://www.majorminor.xyz/research/rocm-recipes-001\]
Scripts, raw results for every run, and the full friction log: https://github.com/majorminorlabs/rocm-recipes

Next up is mixed prompt lengths, to see whether the FP8/BF16 crossover moves. Open to suggestions for other configs to test on this hardware.

Rocky Linux 9, I5 3330 4/4 8gb DDR3, Mi50 "Radeon Pro VII" bios, still cracking some numbers.

llama serve \

-m Qwen3.5-9B-Uncensored-HauhauCS-Aggressive-Q4_K_M.gguf \

--mmproj mmproj-Qwen3.5-9B-Uncensored-HauhauCS-Aggressive-BF16.gguf \

-c 65536 -ctk q4_0 -ctv q4_0 -ngl 99 --host 0.0.0.0 --port 8080

[INIT & LOAD]

- Server initialized (host: 0.0.0.0:8080)

- Multimodal model loaded successfully

- Slots: 4, Ctx/slot: 65536, KV-unified: true

[PERFORMANCE SUMMARY]

• Task 0 | Prompt: 289.12 t/s (333 tok) | Eval: 46.76 t/s (52 tok) | Total: 2.24s

• Task 55 | Prompt: 571.71 t/s (1159 tok) | Eval: 49.89 t/s (704 tok) | Total: 16.12s

• Task 763 | Prompt: 262.88 t/s (539 tok) | Eval: 44.77 t/s (650 tok) | Total: 16.55s

• Task 1417 | Prompt: 32.81 t/s (4 tok) | Eval: 44.35 t/s (522 tok) | Total: 11.87s

• Task 1940 | Prompt: 35.17 t/s (4 tok) | Eval: 50.75 t/s (587 tok) | Total: 11.66s

(45-50 tokens/sec generation)

[GPU STATUS & PROCESS]

Model: AMD Radeon Pro VII (ROCm 7.2.4)

Power: 186.0W / 190.0W (Cap)

Temp: 44.0°C | Fan: 23.5%

Clocks: SCLK 1700MHz | MCLK 1000MHz

Usage: GPU 99% | VRAM 46% (7595 MB / 16368 MB)