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:

Median collective call	Collective share of decode interval
Before (host-staged)	~1,001–1,020 µs
After (CONFIG_HSA_AMD_P2P=y)	12.7–16.2 µs

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:

R9700 forward (median)	Decode
Side streams (default)	~16.9–17.0 ms
Everything on main stream	7.18–7.23 ms

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:

Effective BW	% of peak
My Qwen3-Coder-30B-A3B-FP8, TP=1 (53.6 tok/s)	~177 GB/s
mattbucci's dense FP8 27B, TP=2 (22.5 tok/s)	~608 GB/s aggregate

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.