I'm planning a dual RX 9700 AI build with a Ryzen 9 for LLM inference.

I'm unsure if 64GB RAM (2×32GB) is enough, or if I should go for 128GB (4×32GB), since 2×64GB is really expensive.

My main concern is the RAM speed drop with 4 sticks. Would it noticeably affect performance with inference engines like Strata and Qwen models?

Any recommendations?

I've been tinkering / vibing myself an app that I think might be useful - since I've now used it on my own cases for some time and fixed most of the bugs I noticed, I'm risking a public announcement :)

https://pwilkin.github.io/ROCprofGUI / https://github.com/pwilkin/ROCprofGUI

The ROCprofGUI (aka rocprof studio) is a tool I made to be an alternative to ROCprof Compute Viewer and other tools in the ROCm ecosystem - one that is user-friendly, has a modern UI and can be used with agentic help.

Despite its name, the studio can analyze traces from all of the engines that AMD hardware supports: ROCm traces from rocprofv3, Vulkan traces from RADV (needs a patch to the driver currently) and HRX traces for Loom kernels.

Coupled with the studio is an MCP server that allows tracing from your favorite LLM agent. The studio works in two modes - as a GUI (then, if you use the MCP, a popup appears telling you that the MCP server is currently driving operation) or fully headless - in which case there's no UI and just the service running the MCP server.

I've been running it for over a week now while doing my optimization work on llama.cpp, so hopefully the most annoying bugs have been ironed out. I hope it's of use to someone :)

I have a Alienware 16 aurora AC16250 and am using Eden|v0.2.1|Clang 22.1.4 and im on windows

Came accross a website called AMD-Pro-Bench.com Seems like a cool LLM benchmark aggregator for AMD cards, specifically R9700. Anyone heard or use it? Any other sites people using to find performance comparison with their local setups?

There is so much data spread out there...

Okay, the 4-bit version we mentioned in the previous post is here :)

This time, we focused on accuracy.
We wanted to see how close the quant could stay to the original 16-bit model on math, code and knowledge tasks.

The short version: ~190 tok/s, with scores within 1 percentage point of BF16 across the three benchmarks below.

The hardware: 2× Radeon AI PRO R9700 (64 GB VRAM total) + 256 GB system RAM. Some experts are offloaded to RAM.
The original BF16 weights alone are around 360 GB (335 GiB), so there’s quite a bit of model to squeeze in here.

Accuracy against the 16-bit model

We ran both models ourselves: same questions, prompts and sampling seeds, with thinking enabled.

On GSM8K and HumanEval, the differences were smaller than the variation we saw between two runs of BF16 itself.
We're quite happy with that.

MMLU-Pro is 1 percentage point lower.
That difference wasn’t statistically significant in this comparison (p = 0.16), though that alone doesn’t establish equal accuracy.

The gap was concentrated in answers that hit our 16K-token thinking limit.
The 4-bit model tends to think about 12% longer; on questions where neither model hit the limit, the scores were 91.1% for BF16 vs 91.0% for 4-bit.

How fast does it run?

BetterBench 0.6.0, standard profile, using the exact image and launcher we publish:

For anyone coming from the 3-bit post: single-request decode goes from ~216 to ~190 tok/s.
This version uses more bits per expert weight and offloads some experts to RAM, so there is a speed trade-off.

For people using coding agents

There’s an optional --mode 4bit-cache that keeps conversation prefixes in RAM.

In our repeated 64K-token conversation test, it brought time to first token down from ~12 seconds to ~1.1 seconds.

It costs roughly 22% of multi-user throughput, so it’s opt-in. Useful if your agent keeps sending back a long conversation and you care about how quickly the next turn starts.

What “4-bit” means here

It’s mixed precision:

• Per layer, the 410 most-used experts are INT4, using GPTQ in a rotated basis.
• The 102 least-used experts use the 3-bit codebook format from our previous release.
• That averages 3.93 bits per expert weight. The rest of the model uses 6-bit and 8-bit weights.

Our benchmark setup uses about 96 GB of pinned system RAM for experts and tables.
Smaller hosts can use a fallback memory mode, though it can pause briefly under heavy disk activity.
The numbers above are from the 256 GB RAM setup.

A few deeper accuracy checks

For those who want more than the task scores:

• Perplexity: increased by 0.80–0.87% on WikiText-2 and 0.23–0.38% on held-out text, tested at 512 and 4,096 tokens. The held-out set contains 22 long FineWeb-Edu documents that were never used to build this quant.
• Long documents: likelihood stayed within 0.16% of BF16 across 44 held-out documents, 8K–32K tokens long.
• Needle test at 131,072 tokens: 40/40, with every answer matching BF16 word for word.
• Next-token agreement with BF16: 91–93%. For context, the two BF16 implementations we compared agreed 96–97% of the time.
• KL divergence against BF16: 0.031, with 0.008 between the two BF16 implementations. Using standard MXFP4 round-to-nearest on the same experts, we measured 0.052. That’s 41% lower KL divergence on this evaluation set, with about 8% fewer bits per weight.

Still regular vLLM, with our Paiton plugin and our own kernels inside. ;)

Try it

Model / weights
Setup instructions

git clone https://github.com/Eliovp-BV/paiton-vllm-plugin && cd paiton-vllm-plugin
bash models/Qwen3.8-Flash-Next/run-flashnext-4bit.sh

If you give it a try, let us know how it behaves on your own coding tasks or long conversations.
And if something looks off, share your setup and logs.

Edit:
Almost forgot our ai generated image :)

Hi guys. Reposting here hopefully someone can help. Thanks!

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)

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.

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.