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.

Sharing some results from my own PC running [Strata](https://github.com/Niko1221/Strata) with the HIP backend. The model is **Qwen3.8-Flash-Next-GSQ-RCO-IQ3_XXS**. These are engine timings from real chat sessions, including generated reasoning tokens.

My setup:

• AMD Radeon RX 7800 XT, 16 GB VRAM (`gfx1101`)
• Ryzen 7 5800X3D, 8 cores / 16 threads, AVX2
• 64 GB DDR4 RAM
• Ubuntu 26.04.1 LTS, kernel 7.0.0-38-generic
• Strata 0.1.42, built locally for HIP; system HIP 7.15 and hipBLASLt 1.4.1

In my latest session, the prompt/context grew from **41,709 to 75,100 tokens** across **54 completed requests**. The engine generated **18,806 tokens in 289.772 seconds of decode time**:

| Measurement | Result |

| Aggregate decode throughput, total tokens / total decode time | **64.9 tok/s** |

| Median throughput per request | **66.4 tok/s** |

| Per-request decode range | **56.1–82.8 tok/s** |

| Fresh 41,709-token prompt, zero reused tokens | **985.8 tok/s prefill** |

Most later turns reused the conversation prefix, so their prompt timings should not be treated as full-context prefill benchmarks. The expert cache was also warmed by the conversation.

Other fresh-prompt measurements with this setup were **912.9 tok/s at 6,731 tokens**, **1,004.6 tok/s at 76,599 tokens**, and **791.0 tok/s at 116,392 tokens**. I also had a slower 38,790-token read at **509.8 tok/s**. Across all 61 completed tail-skip-enabled requests in today's log, including the slower sessions and much longer generations, aggregate decode was **52.7 tok/s**. The ~65 figure refers specifically to the latest 54-request session.

The settings used for these runs:

--expert-cache auto
--prefill auto
--max-context 131072
--kv int8
--spec 4
--spec-min-p 0.70
--pcie-frac 0.35
--adapt-every 1
--adapt-swaps 160
--adapt-decay 0.97

MTP was enabled. The default CPU pool used 7 workers plus the host thread. The PCIe probe measured about 27.9 GB/s, but decode calibration selected `--pcie-frac 0.35` rather than the automatic 0.55. The engine loaded **39.97 GiB of expert weights into host RAM** and cached **5,048 experts / 8.19 GiB** in VRAM at startup.

I also used:

export STRATA_SH_STREAM=1
export STRATA_ROUTE_TAIL_SKIP=7
export STRATA_HIPBLASLT_TUNING="/path/to/strata-gfx1101-hipblaslt.txt"

The hipBLASLt table was generated locally for **gfx1101 + hipBLASLt 1.4.1**, with 26 rows. The log confirms that it was loaded. Those solution IDs are tied to that architecture and library version.

The project's **experimental speed projection** was also enabled at strength 1.0, layers 4–44, using `--cvec-mode project --cvec-dir per-layer`.

One informal comparison on the same accounting/Python prompt gave **50.8 tok/s without route tail skip** and **56.7 tok/s with it**. Both visible answers were correct, but the generated lengths and expert-cache hit rates differed, so this is not a controlled estimate of the switch's benefit. Tail skip changes the computation and can change the answer; I have not run a comprehensive quality benchmark.

I hit one watchdog stall during an earlier stock-config run. The latest 54-request session completed without a reported stall. This is a report of my own measurements, rather than a claim that these settings will be fastest or stable on every AMD setup.

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)

Okay, here we are again :)
We received our second R9700 earlier this week, so we decided to jump on tensor parallelism.
We already had prepared a few things for this, and the TP kernels we already had in Paiton did their work.

We spent a lot of time (significant amount of GPU time (MI355x)) on the 3-bit quantization of this model and we are pretty happy with the accuracy.

What "3-bit" means here.

This is not a full 3-bit model. The 512 experts in every layer, that is where almost all the weights are, are stored at 3.125 bits per weight (3-bit codes, rotated, with a small scale per 64 weights). Everything else stays bigger: the attention and the shared parts are 8-bit, the speculative layer is 4-bit, and the vision part is plain bf16.
So the big part is 3-bit, the sensitive parts are not, and that is why the accuracy holds up. Counted over all weights it is a bit above 3 bits on average. Our own GPTQ run, no shortcuts.

Accuracy against the bf16 model (same evaluation corpus, text and assistant positions; two correct bf16 implementations against each other are the floor):

KL to bf16	KL to bf16, same expert routing forced	Top-1 agreement with bf16
bf16 vs bf16 (the floor)	0.0077	n/a
This 3-bit model	0.060	0.041

Long context check	Result
Needle test, 131,072 tokens	40 / 40, same answers as bf16
Needle test, 200,000 tokens	40 / 40, same answers as bf16

The whole language model (experts, trunk, speculative layer, draft head, vision tensors) sits in the VRAM of the two cards; only the model's n-gram embedding table lives in pinned system RAM. No expert offload, no weights streamed from the host during decoding.

And again, like before: we did not build a custom engine for this. It is still regular vLLM, with our Paiton plugin and our own kernels inside. ;)

Performance.

BetterBench 0.6.0, standard profile, measured on our own machine against our OpenAI-compatible server, with the exact image and launcher we publish. Decode mode: 98K context, speculative decoding on. Prefill with exact bf16 activations, no compression.

Decode, single request	tok/s
Weighted over the 8 categories	216.3
chat	202.1
code	214.3
file edit	236.0
json	259.2
math	255.5
prose	183.3
reasoning	190.1
summarization	239.8

Concurrent requests	Aggregate tok/s
1	204.6
2	315.1
4	449.2
8	565.1 (48 of 48 requests ok)

Prefill, prompt length	tok/s
2K	7,925
8K	8,321
16K	8,311
32K	8,232
64K	8,161
128K	7,928

Latency
Time to first token, short prompts (p50)
Update p99 (the longest gap between two tokens)

Long context. We tested 200K. The launcher has a 200K mode, and prefix caching is an opt-in flag that gives exactly the same output on a hit.

200K mode	Result
Needle test at 200,000 tokens	40 / 40, same answers as bf16
Single-stream decode after a 32K / 100K / 190K prompt	105 / 101 / 100 tok/s, flat
Time to first token, 190K prompt	25 s
Prefix cache hit, 64K repeated prompt	9.6 s cold, 0.35 s on the hit
Prefix cache hit, 128K repeated prompt	20 s cold, 0.41 s on the hit

Overall, not bad for a first release :)
Oh, right, of course we generated an "AI Slop" image as well. (we know some of you like them (some of you don't ;-)))

https://preview.redd.it/836rjbm2cmuh1.png?width=1672&format=png&auto=webp&s=033741b4cfe64c03fdd966a400d847150ca275de

Weights
Image and launcher: github.com/Eliovp-BV/paiton-vllm-plugin (directory models/Qwen3.8-Flash-Next)

This is our first release of this model, more will follow.
Next is the RAM and SSD cache tier, like we did for the 27B model, so cached prompts can live outside the VRAM and much longer contexts become possible.
After that the 4-bit version, with even better accuracy.
We keep tuning, we keep optimizing.
A longer blog post with all the details is being worked on as well.

And of course, feedback is very welcome!

I've been tuning llama.cpp on my Strix Halo box (Radeon 8060S, 128 GB) and the numbers are good enough to share. It's a branch called rdna35-perf on top of halo-box/strix-llama.cpp, a llama.cpp fork aimed at Strix Halo:

https://github.com/asf0/strix-llama.cpp/tree/rdna35-perf (10 commits, ROCm/HIP only)

TL;DR

• Prompt processing: 21–54% faster than Vulkan, and 8–174% faster than stock ROCm. The lead over ROCm grows with context, because stock ROCm flash attention falls off hard at depth.
• Generation: ties Vulkan on the dense 27B. On the MoE and the 2B it's 2–3% behind Vulkan at short context and within 1% from 64k up. 5–29% faster than stock ROCm.
• Quality: outputs stay closer to stock ROCm than Vulkan's do, and perplexity is unchanged.

Model	PP vs Vulkan (0k → 128k)	PP vs stock ROCm (0k → 128k)	TG vs Vulkan (0k → 128k)
Qwen3.6 35B-A3B (MoE)	+54% → +36%	+38% → +88%	−3% → 0%
Qwen3.8 27B (dense)	+38% → +25%	+8% → +174%	0% → 0%
Qwen3.5 2B (dense)	+21% → +34%	+9% → +87%	−3% → 0%

Results

llama-bench, Q8_0 weights, q8_0 KV cache, flash attention on, tok/s. "Depth" is how many tokens are already in the KV cache when the test runs. Vulkan and ROCm are the same strix-llama.cpp master that rdna35-perf is built on.

Qwen3.6 35B-A3B (MoE, ub 4096)

Test	Depth	Vulkan	ROCm	rdna35-perf	vs Vulkan	vs ROCm
pp4096	0	2625	2935	4045	+54%	+38%
pp4096	8k	2304	2430	3452	+50%	+42%
pp4096	16k	2065	1966	3004	+45%	+53%
pp4096	32k	1716	1460	2388	+39%	+64%
pp4096	64k	1279	959	1695	+32%	+77%
pp4096	128k	786	568	1068	+36%	+88%
tg64	0	58.2	52.3	56.3	−3%	+8%
tg64	8k	56.1	50.4	54.3	−3%	+8%
tg64	16k	54.6	48.5	53.2	−3%	+10%
tg64	32k	51.8	45.3	50.7	−2%	+12%
tg64	64k	47.1	39.7	46.4	−1%	+17%
tg64	128k	39.9	32.2	39.9	0%	+24%

Qwen3.8 27B (dense, ub 2048)

Test	Depth	Vulkan	ROCm	rdna35-perf	vs Vulkan	vs ROCm
pp2048	0	484	614	666	+38%	+8%
pp2048	8k	442	542	611	+38%	+13%
pp2048	16k	421	476	575	+37%	+21%
pp2048	32k	385	348	513	+33%	+47%
pp2048	64k	326	204	422	+29%	+107%
pp2048	128k	250	114	312	+25%	+174%
tg64	0	8.0	7.6	8.0	0%	+5%
tg64	8k	7.9	7.4	7.9	0%	+7%
tg64	16k	7.8	7.2	7.8	0%	+9%
tg64	32k	7.6	6.8	7.6	0%	+12%
tg64	64k	7.3	6.2	7.3	0%	+18%
tg64	128k	6.8	5.2	6.7	0%	+29%

Qwen3.5 2B (dense, ub 2048)

Test	Depth	Vulkan	ROCm	rdna35-perf	vs Vulkan	vs ROCm
pp2048	0	6481	7179	7852	+21%	+9%
pp2048	8k	5618	6253	7158	+27%	+14%
pp2048	16k	5000	5172	6611	+32%	+28%
pp2048	32k	4155	3888	5580	+34%	+44%
pp2048	64k	3220	2597	4277	+33%	+65%
pp2048	128k	2160	1547	2886	+34%	+87%
tg64	0	99.3	91.5	96.8	−3%	+6%
tg64	8k	95.9	88.2	93.2	−3%	+6%
tg64	16k	92.9	84.5	91.4	−2%	+8%
tg64	32k	88.3	79.1	87.1	−1%	+10%
tg64	64k	79.5	69.2	79.2	0%	+14%
tg64	128k	67.6	56.2	67.7	0%	+21%

f16 KV gives the same picture: PP 20–56% faster than Vulkan and up to +171% over stock ROCm at 128k, TG within −3% / +2% of Vulkan.

Quality

rdna35-perf changes kernels, so I checked it doesn't change outputs: KL divergence against stock ROCm on wikitext-2 (6 × 4096-token chunks), plus perplexity at 32k context.

Model	KLD vs stock ROCm (Vulkan / rdna35-perf)	Same top token (Vulkan / rdna35-perf)	PPL @ 32k (Vulkan / ROCm / rdna35-perf)
Qwen3.6 35B-A3B	0.0038 / 0.0018	97.8% / 98.5%	6.028 / 6.023 / 6.024
Qwen3.8 27B	0.0066 / 0.0013	98.9% / 99.7%	6.003 / 6.004 / 6.007
Qwen3.5 2B	0.0008 / 0.0001	98.6% / 99.5%	10.046 / 10.049 / 10.049

rdna35-perf is closer to stock ROCm than Vulkan is, and perplexity is the same across all three builds.

What changed

• Flash attention kernel for RDNA3.5 (fattn-rdna35). This is the biggest win and the reason long-context PP holds up. GGML_FA_RDNA35=0 turns it off if you want to compare.
• Prefill GEMM tuning for the quantized matmuls.
• MoE: routing helper, fused reduction + residual, grouped mat-vec for speculative-decoding verify batches.
• Gated DeltaNet (the linear-attention layers in Qwen3.5/3.6/3.8): fused SiLU-gated output norm and fused decode.
• Decode: plain mat-vec for rows longer than 4096 (this closed the 27B decode gap to Vulkan), plus decode fusions for the Qwen3.6 MoE.
• Misc: HIP graphs off by default on RDNA3.5 (slower here), a faster penalties sampler, host-side prefill for Qwen3.5/3.6 with a draft model, and DFlash drafts read target features in place on the GPU.

I have only tried those 3 models so far.

Hi everyone!

Well this is my first contribution for community with the help of vibe coding :)

TL;DR: links at bottom that boost Qwen Image Edit clip encoding up to x7 times!
(51.3 s → 6.7 s (~7.6x) for the two custom encode/clip nodes.)
____

Context:
I am running comfyui-rocm Patientx-cfz build,
ComfyUI 0.39.0 / ComfyUI_frontend v1.53.10
Win 10, AMD 7900XTX, 5900x, 64Gb ram
PyTorch Version 2.15.0a0+rocm10.2.0a20261006
Default comfy WF with kitchen attention and 6-step lora applied.

So I really was upset with how painfully slow Qwen 2.1 edit was with image refs. CLIP taking much longer than sampler itself. I saw people suggested lowering input images resolution, but that's always led to details loss on my tests. And re-encoding images on prompt change was making it even worse. So I asked OpenCode and DeepSeek4.1 Flash to look at it and ask to add separate image encoding node and clip encode, so encoded references at least could be cached.

But Ai found even more optimizations!
So my test workflow with 1 ref image with resolution set to 1248 in qwen clip node and final render res of 1088x1440:

Orig - 66s
Mod - 21s

That's some real big win I think! And this actually made possible to take my ref image with original size of 1440x1920 and be as fast as 40s!
As a bonus we made separate Qwen Image 2.1 Reference and Text Encode Qwen Image 2.1 Fast, so now if you don't change ref images and change only prompt - images doesn't needs to be re-encoded!

2nd run with only Prompt changed:
Orig - 62s
Mod - 18.5s

These custom nodes doesn't modify any ComfyUI files. When you run one of these nodes it patches the Qwen VL vision encoder in memory. The QwenVL Vision Linear Patch toggle switches the vision patch-embed between the original Conv3d and an equivalent linear op (bit-identical output), so you can A/B in the same workflow.

Currently ONLY TESTED ON MY MACHINE!

What I don't know - is if that work on someone's else hardware, especially NVidia, so if you interested to try it out - you can do it and tell here :)

Also - Ai told me that this may also help other models that use the Qwen VL vision tower (e.g. MiniMax H3), untested so far.

LINK to nodes.
Workflow (with original AND fast nodes in one wf) Simply turn on/off groups you don't want to run (use original OR fast)

How to use:

-Unzip into custom_nodes, restart ComfyUI, hard-refresh the browser. Nodes appear under QwenVL Vision.

- Put QwenVL Vision Linear Patch after your CLIP loader and set it to ON. It swaps the vision Conv3d for a mathematically identical linear op — output is bit-identical. Toggle OFF to compare in the same workflow.

- Use Qwen Image 2.1 Reference → Text Encode Qwen Image 2.1 Fast, so changing only the prompt no longer re-encodes the reference images.

- Original Text Encode Qwen Image 2.1 node is still slower, as it encodes every image twice: once for pos and neg prompt - mod does it once, as original encode is same for both anyway, independent from text.

Sharing a ComfyUI node I've been working on, since video gen on Strix Halo is where a lot of us keep hitting the wall.

Credit first: Veda is a learned sparse-attention method for video diffusion models (ICML 2026). A small distilled predictor says in advance which tiles of the attention map actually carry the result; only the top ~10% get computed. The original project is by veda-sparse:

- Project page: https://veda-sparse.github.io/

- Original repo: https://github.com/veda-sparse/Veda-on-ComfyUI

- Predictor model file (275 MB): https://huggingface.co/Veda-Sparse/Minimax-H3-T2VA-Veda-8NFE-600Step-Preview — goes in ComfyUI/models/veda/ (opening a Veda template offers it in the missing-model dialog)

Upstream supports NVIDIA (SM80+) and Apple silicon, and lists ROCm as "no kernel, the model runs its own attention". So I ported it to ROCm and tuned it for Strix Halo. What that took:

- Backend layer: report the real gfx family (gfx1151) and unified memory on HIP, and admit ROCm devices to the Triton INT8 backend (kernels still self-test on the GPU before use, same as upstream)

- Triton's ROCm backend compiles the INT8 kernel for gfx1151: Q and K are block-quantized to int8, QK^T runs on the hardware int8 WMMA instructions (v_wmma_i32_16x16x16_iu8), P*V stays fp16 WMMA

- The launch config upstream shipped (4 warps / 3 stages) was tuned on an RTX 5070 and is wrong on this chip: measured on the real portrait grid, 8 warps / 2 stages is 2x on the attention kernel (35.3 -> 17.6 ms per call)

- A few more gfx1151-specific kernel tweaks, e.g. the quantize pass runs 23% faster with 2 warps, bit-identical output

- Accuracy verified on this hardware: kernel vs the INT8 reference is ~0.14% relative L2, below the quantization noise that's already there

Numbers on my machine, MiniMax H3, 11-second clip at 480x864:

- dense attention sampler: 691s

- sparse (Veda): 366s

- 1.89x on sampling, 2.15x on the full prompt (1016s -> 473s)

- only ~27% of the attention is computed, output looks the same to me

My ROCm port lives in rocm-ninodes (ComfyUI Manager, or comfy node install rocm-ninodes); the node is VedaSparseAttention, drop it on the MODEL wire last before the sampler:

- https://github.com/iGavroche/rocm-ninodes

If you try it on your Halo I'd be curious about numbers on other resolutions/lengths. Happy to answer questions about the sparse selection or the ROCm port.

tl;dr: Veda Sparse Attention on ROCm (optimized for Strix Halo)

TL;DR: Found a fix for 100% CPU usage of llama-server on AMD RX 6600 GPUs. Fixes in github link at the end!

Disclaimer: Claude Opus 5.5 was used to run traces and find the fixes! Posting here so that others in the same boat might find it helpful.

---

I've been using gemma4:e4b-qat locally on my server for a couple months now.

Here's my setup:

CPU: i3-10100F

RAM: 32GB (4x 8GB)

Board: MSI Z490 Gaming Plus

GPU: 7x AMD RX 6600

I had this rig from ETH mining a couple years ago when that was a thing! It was sitting there collecting dust for a long time, so I had an idea to use it to run some local LLM experiments and started tinkering with it. After testing different models, settled on gemma4:e4b-qat at 128k context which runs comfortably at 5.5/8GB VRAM (tell me if you found a better model for the 8GBs).

I am running an app that takes a list of stock tickers and runs the TradingAgents (https://github.com/tauricresearch/tradingagents) pipeline on them twice a day (market open and market close) as a fun side project to see how good that is.

So, for running TradingAgents which makes about 30-40 LLM calls per ticker, I am running a custom llama-pool which runs a llama-server per GPU with gemma4:e4b-qat model at 128k context. The thing I noticed was that my CPU is at 100% while it's running the pipeline so I started digging into it using Claude Opus 5.5. After a couple hours of testing, traces, LLM calls, it found a fix and applied them which reduced the CPU usage down to ~20% for 7 parallel llama-servers.

The fixes and instructions are all published here: https://github.com/nandyalu/rocm-llamacpp-cpu-fix

I built LlamaStash to scratch a personal itch: I run local models through llama.cpp on AMD Strix Halo and got tired of writing the same llama-server wrapper script for the tenth time.

Ollama and LM Studio both wrap llama.cpp but hide too much (and cost real performance). Raw llama-server is fast but tedious. LlamaStash is the middle ground.

What it does:

• llamastash init — first-run wizard. Detects your hardware (CUDA / ROCm-HIP / Metal / Vulkan / CPU), installs llama-server, scans your existing HuggingFace / Ollama / LM Studio model caches, recommends a GGUF that fits your VRAM, downloads it, writes a tuned config, smoke-launches it.
• TUI + CLI + daemon + OpenAI-compatible proxy in one Rust binary. The proxy at 127.0.0.1:11435/v1 lets OpenCode, Cline, the OpenAI SDKs, and llm-cli work as-is. There's also an opt-in --ollama-compat mode that takes port 11434 and answers the byte-exact "Ollama is running" handshake.
• Multi-model concurrency with per-model port allocation, /health-probed state machine, intelligent context auto-fit (sidesteps llama.cpp's --fit collapse on Linux iGPUs).
• Agent-friendly CLI: every TUI capability has a CLI subcommand, --json is a stable agent contract, documented exit codes per failure class.
• In-TUI HuggingFace browser with search, sort, paginate, per-file hardware fit, download with cancel.

On performance — this is the part that matters for this sub.

LlamaStash spawns the unmodified upstream llama-server. So the wrapper should add zero overhead. I measured it. Across AMD APU (Ryzen AI Max+ 395), Apple Silicon, and NVIDIA, on four model sizes (small E2B Q4, mid 31B Q4, large 27B Q8, large MoE 35B-A3B Q8), every cell matches raw llama-server within ≤1%.

Cross-tool numbers on AMD APU (decode tok/s / TTFT ms on chat_turn):

Tool	small	mid	large_dense	large_moe
LlamaStash	86.9 / 51	9.8 / 467	7.4 / 417	42.6 / 181
raw llama-server	86.0 / 51	9.9 / 468	7.4 / 414	42.7 / 186
LM Studio 2.16.0	91.1 / 187	11.6 / 1477	7.9 / 1274	37.0 / 683
Ollama 0.24.0	50.4 / 223	4.8 / 1092	2.6 / 1745	12.1 / 476

LM Studio wins decode on small/mid/large_dense (their Vulkan path is well-tuned on gfx1151) but loses on the MoE and pays a 1-1.5s TTFT tax from its OpenAI shim. Ollama is consistently slower, and its RAG prefill is catastrophic (cold prefill every rep — 4 min on a 31B). Mac and NVIDIA tables are in the benchmarks page.

Methodology, variance gates, fairness rules, and per-cell JSONs are all checked in. The harness is reproducible: make bench-end-to-end. Tear it apart.

What it's not:

• Not an Ollama fork or replacement (though --ollama-compat exists for tools that auto-detect Ollama).
• Not a model hub.
• Not a llama.cpp fork. Same upstream binary.
• Not a hosted service. Loopback-only in 0.0.2. LAN + auth + TLS are on the roadmap.

Install:

curl -fsSL https://llamastash.dev/install.sh | sh   # macOS + Linux one-shot
irm https://llamastash.dev/install.ps1 | iex        # Windows 11 (PowerShell, no admin)
scoop bucket add llamastash https://github.com/llamastash/scoop-llamastash && scoop install llamastash
brew install llamastash/llamastash/llamastash       # Homebrew (macOS + Linuxbrew)
yay -S llamastash                                   # Arch Linux (AUR — source build)
yay -S llamastash-bin                               # Arch Linux (AUR — prebuilt binary)
yay -S llamastash-git                               # Arch Linux (AUR — main checkout)
cargo install llamastash                            # any Rust toolchain

Then llamastash init and you're up.

Platform: Linux (x86_64, aarch64), macOS (Intel, Apple Silicon), Windows 11 (x86_64). aarch64-pc-windows-msvc and Windows AMD GPU detection on the roadmap.

Honest tradeoffs: Single-author project. Bug reports especially welcome on hardware I don't own. The OpenAI-compat surface covers chat/completions, embeddings, rerank; Anthropic /v1/messages shim is coming.

Repo: https://github.com/llamastash/llamastash

Blog post with the full story: https://deepu.tech/introducing-llamastash

Benchmark methodology: https://deepu.tech/benchmarking-llamastash

Happy to answer questions in the thread.

Tune my Strata setup for longer context: 131K, 262K first, then 512K if it holds up. I'm on Strata v<FILL VERSION, e.g. 0.1.39>.

My hardware:
- GPU(s): <model + VRAM each, e.g. 1x RTX 4090 24 GB / 2x RTX 5090 32 GB>
- System RAM: <e.g. 32 GB / 192 GB>
- CPU: <model, cores>
- Platform: <bare metal or VM (Proxmox/ESXi/etc.)>, OS: <distro/kernel>
- Storage for model files: <NVMe/SATA, model>

Model: Qwen3.8-Flash-Next. Suggest a quant after reviewing my hardware.
REPO: https://github.com/Niko1221/Strata

Ground rules:
- Read docs/DETAILS.md in the Strata repo first, especially the sections on --max-context, --kv, --kv-resident, --prefill, the expert cache, multi-GPU split, and "Context extension past 262K (rope scaling, EXPERIMENTAL)". Go by the docs and `strata --help` for my installed version, not by memory. If a flag below doesn't exist in my build, tell me; don't invent one.
- Change one thing at a time. Before you change anything, record a baseline of my current config: decode tok/s on a short prompt, decode tok/s and prompt-read speed at depth (one ~100K-token prompt, then one near the target context), plus a needle-recall check at depth. Run every candidate the same way and alternate baseline and candidate at least twice, so drift doesn't show up as a gain.
- Watch host RAM and VRAM on every GPU during each run (nvidia-smi, free -g). Report the minimum free RAM, not just the speed.
- Don't run a big compile while the Strata server is loaded. Stop it first.
- Back up the working config before every edit, and give me the exact undo.

What a single-4090, 32 GB RAM box (IQ2_XS) learned, as starting points to test, not answers. Scale these to my hardware:

1. 262K is native for this model. At 262K, --kv q4_0 kept the KV cache in VRAM and stayed fast; recall held at ~238K. If I have more VRAM than that box (24 GB), check whether int8 or k8v4 KV fits at 262K. Measure both against q4_0; higher-precision KV is worth it if it fits.
2. Past 262K needs rope scaling: --rope-scaling yarn --rope-scale 2 for 512K, together with --max-context 524288 and a streamed KV cache (--kv-resident <tokens>, e.g. 32768) so most of the KV lives in host RAM. A 477K-token prompt read in ~150 s and decoded at ~85 tok/s at that depth on that box. It's marked experimental, so run a recall test at 400K+ before trusting it. Estimate how much host RAM streamed KV needs at 512K for my setup and tell me up front if my RAM can't support it.
3. --prefill auto (default) beat --prefill auto:32768 by about 3x on long prompt reads for us. Measure prompt-read speed for any prefill setting you try.
4. "fit_max_tokens": true in the server config keeps requests within the context.
5. If I have more than one GPU: check how my version splits layers, experts and KV across the cards, and whether auto placement accounts for each card's PCIe link (confirm each GPU's link width/gen with nvidia-smi -q and lspci -vv). Tell me where the KV ends up at each context size. Skip this if single-GPU.
6. If I'm in a VM: check whether the guest sees RAM as one NUMA node, whether hugepages or THP are in use for the expert and KV arenas, that vCPUs are pinned, and that GPUs show their full PCIe link inside the guest. Report what you find; don't change the hypervisor host without asking me. On bare metal, still check NUMA layout and THP/hugepages.

If I mention a stat that used to show up and is now missing (e.g. "VRAM hit rate"), check the changes between my previous and current version (git log, CHANGELOG, server /health, /props, /metrics, and the engine log) and tell me whether it was renamed, moved, hidden when experts are all resident, or removed. Quote the commit or line that answers it.

Deliver: a table of each config tried (context, KV type, kv-resident, rope, prefill) with short and deep decode tok/s, prompt-read speed, recall pass/fail, min free RAM and per-GPU VRAM; a recommended daily config and an optional 512K config (or a clear "not viable on this hardware" with the reason); and the exact config diff plus undo for each.

Posting for all the new users and current.

how to use the tuning prompt

1. copy the whole prompt block above.
2. fill in the hardware section at the top: gpu model and vram per card, system ram, cpu, bare metal or vm, os, and what drive the model lives on. also put your strata version in (run `strata --version` if you're not sure).
3. paste it into an agent that can run commands on your box. codex or claude code both work. use opus 5.5 or astra if you have access. this prompt has it reading the repo docs, running a bunch of benchmark passes, checking numa/pcie/hugepages, and diffing git history, and the smaller models tend to lose the thread halfway through or start making up flags.
4. let it run the baseline first. don't skip this. without a baseline you can't tell if a change helped or if your box just warmed up.
5. expect it to take a while. every config gets run at least twice against the baseline, and long-context prompt reads at 262k+ aren't fast.
6. when it's done you get a table of every config it tried, a recommended daily config, an optional 512k config (or a straight answer that your hardware can't do it), and the exact diff plus undo for each change.

notes

- the numbers in the prompt come from a single 4090 with 32 gb ram. they're starting points, not targets. your results will differ.

- 512k uses rope scaling and is marked experimental. trust it only if the recall test at 400k+ passes.

- streamed kv at 512k leans on system ram. if you're on 32 gb or less, 262k is probably your ceiling.

- if you're in a vm, it'll report on the host side but won't touch your hypervisor without asking.

- back up your config before you start anyway. the prompt tells it to, but don't rely on that alone.

EDIT: if you use the guide, please just post a quick update if it helped your setup. Reach out if you have any issues please.

I want to thank everyone for the kind words, trust, and encouragement in the development of this project. I am doing my best to make the halogen UX amazing for you. I also want to thank AMD's Jeremy Fowers. u/jfowers_amd has been nothing but supportive and sent me a development Strix Halo machine which has been taking a beating over the last couple weeks.

halogen-flash-server 0.16.0 is out.

It runs Qwen3.8-Flash-Next on the Strix Halo GPU, and this release puts the chip's other accelerator to work too.

The Ryzen AI NPU now serves small models beside the Flash model, on the same port, while the GPU keeps generating.

What's new

Small models on the NPU, with one variable. Set HALOGEN_NPU_MODELS and the server runs them on the NPU. Leave it out and nothing changes.

• decider-0.8b makes decisions. Give it a text, a question and 2 to 10 options, and it returns a probability for every option from one pass.
  • Use it as a guard, a router or a classifier.
• qwen3-embedding-0.6b answers `/v1/embeddings`, in the OpenAI request shape.
• qwen3-reranker-0.6b answers `/v1/rerank`, in the Cohere and Jina request shape.

Your own fine-tunes. Put a fine-tune of any of the three on the models volume and name its path. The server converts it for the NPU at the first start and serves it under the directory's name.

A smaller checkpoint, opt-in. qwen38-flash-next-ht43.hgn pins about 8 GiB less memory than the default checkpoint. Prefill is several percent slower, and decode with the default drafter is a few percent slower. The default checkpoint does not change. Select it with HALOGEN_CHECKPOINT=/models/qwen38-flash-next-ht43.hgn.

96 GB: On a 96 GB machine, run ht43 with HALOGEN_MAX_TOK=16384 and HALOGEN_CTX=131072. That leaves room for one conversation of 131k tokens or four of 32k.

Fixes. A long prefill can be cancelled now (#122, thanks u/semidark). The draft counters count every drafted token (#124, thanks u/wszgrcy).

How fast is the NPU?

• A decision with Qwen3.5-0.8B (the architecture `decider-0.8b` uses) takes 49 ms at 64 tokens, 84 ms at 256 and 183 ms at 1k. Eight decisions of 256 tokens in one pass take 41 ms each.
• Batched embeddings run at about 9,400 tokens a second through the server, eight texts of 512 tokens at a time. One text of 512 takes about 97 ms.
• The reranker scores 19 pairs a second, eight pairs of 512 tokens at a time.

The NPU and the GPU together

Something worth knowing even if you never run halogen. On Strix Halo, GPU work and NPU work at the same time can hang the whole machine while the GPU's fabric clock changes speed. Holding that clock at its top level stops it, at a cost of about a watt. The server holds it itself when it runs as root with /sys mounted. For rootless podman, a small systemd unit in the repo holds it at boot. If you run an NPU program next to a GPU program and have seen freezes, this may be why.

The Flash model runs somewhat slower while the NPU works, since the two share the memory and the power budget. With the NPU idle there is no cost.

Setup

The host needs the amdxdna driver, XRT with its NPU plugin, the NPU firmware, and the IOMMU on. If you added amd_iommu=off for GPU speed, it also turns the NPU off. iommu=ptkeeps it. The full setup, every request shape and the fine-tune rules are in https://github.com/peonist-ai/halogen-flash-server/blob/main/docs/NPU.md

podman pull ghcr.io/peonist-ai/halogen-flash-server:0.16.0

Repo: https://github.com/peonist-ai/halogen-flash-server

Weights: https://huggingface.co/peonist-ai/halogen-qwen3.8-flash-next

Discord: https://discord.gg/bcm6QknaV6

How to read this. FastFlowLM's numbers are from its own published benchmark pages. They were measured on a different chip (Ryzen AI 7 350) with older versions than its current one, in its default power mode. FastFlowLM lists no embedding model, so its Qwen3-0.6B row (the same backbone) stands in for it. Ours are medians from our own testing on Ryzen AI Max+ 395 machines, and where we measured two machines, the slower one is shown. This is not a same-machine test. If you run both on one machine, we would like to see your numbers.

TL;DR: Qwen3.8 27B on a single AMD Radeon AI PRO R9700, with 3-bit weights and speculative decoding, now keeps 569,878 tokens of reusable cache in its new coding mode.
Every request still gets 262,144 tokens of context, and two full-length requests fit at once. Coding agents resend the whole conversation on every turn; now only the new part is read.
Later agent turns start up to 12× sooner, whole agent sessions finish about 6× faster, and a 258K-token document the server has already read comes back in 2.7 s instead of 134 s. Decode speed and accuracy are within noise of our previous release.

We read your comments and we listened

You asked for more context. You got more than half a million tokens of cache on one card, with multi-hour accuracy runs on every configuration (GSM8K, HumanEval, MMLU-Pro, needle tests), at essentially the same speed as our previous release.

	Previous release (2 Oct)	This release, --mode long-kv4
Reusable (prefix-cached) tokens in the 4-bit mode	0 (no prefix cache)	569,878
Tokens per request	262,144	262,144
Full-length requests at once	1	2
Decode, single stream	174.3 tok/s	173.4 tok/s (−0.5 %)
Aggregate, 8 requests	462.4 tok/s	478.5 tok/s (+3.5 %)
Time to first token, short prompts (p50)	86.7 ms	87.1 ms
GSM8K / HumanEval / MMLU-Pro	95.53 / 92.07 / 60.57	95.45 / 92.68 / 59.93

Our previous release's FP8 --mode long already had a prefix cache, holding 281,665 tokens. The new 4-bit coding mode caches about twice as much.

Run it

git clone https://github.com/Eliovp-BV/paiton-vllm-plugin && cd paiton-vllm-plugin
# set the PAITON_* model paths as shown in the README (already set up? just `git pull`)
bash models/Qwen3.8-MXFP4-DFlash2/run-3bit.sh --mode long-kv4

Point your agent at http://127.0.0.1:18982/v1
The model name is Qwen3.8, any API key works, and the context window is 262,144 tokens. Already running our 3-bit weights? No new download. The launcher pins the new container and Docker pulls it on first start. Every mode is in the README.

No new engine

This is stock vLLM with a plugin: the same OpenAI-compatible server, the same API, and the same tools and workflows you already use. Nothing to relearn, nothing to migrate.

We will keep publishing new ready-built containers that follow upstream vLLM. Each one goes through the same full validation (speed, accuracy, long-context) before it ships. You get the upstream improvements without lifting a finger. One command and you're serving.

What it does for coding agents

Without a prefix cache, the server rereads the whole prompt on every turn: at 250K tokens that is about two minutes before the first token. Now finished requests stay cached until the space is needed, turn 20 reads only what changed, and in our sessions 91–92 % of prompt tokens came from the cache.

Workload (3-bit, thinking off)	Previous 4-bit mode (no prefix cache)	--mode long-kv4	Faster
20-turn conversation growing from 50K to 253K tokens, whole session	1,237 s	204 s	6.1×
… average wait for the first token, turns 2–20	63.6 s	7.8 s	8×
3 agents sharing a 100K-token repo, 5 turns each, whole session	655 s	111 s	5.9×
… average wait for the first token, later turns	84 s	7.0 s	12×
New question about a 258K-token document already read	134 s	2.7 s	49×

The cached answer to the 258K-token document matched the cold read exactly.

Speed details (BetterBench, full 20-pass run): decode within 0.5 % single-stream; +3.5 % with 8 requests (−2.5 % at 4); time to first token unchanged.
A new long prompt reads about 3 % slower, and short prompts 7–13 % slower (a fraction of a second). That's the price of ending each prefill step where a later request can pick up from the cache.

Accuracy (greedy, paired per question): the coding mode is within noise of the previous release, and the 512K mode is within noise of the coding mode.

	MXFP4	Previous release (3-bit)	--mode long-kv4	--mode long-512k
GSM8K 5-shot (1,319)	95.68	95.53	95.45	95.53
HumanEval pass@1 (164)	95.12	92.07	92.68	94.51
MMLU-Pro subset (1,400)	62.57	60.57	59.93	60.64

Opt-in: 524,288 tokens in one request

--mode long-512k extends the position range with the model's official long-context scaling. That scaling applies to every request in this mode, so use it only when a single request needs more than 262K.

• It found 4/4 planted facts at 300K and 4/4 at 500K tokens.
• A cold read takes 158 s at 300K and about 6 minutes at 500K. Follow-up questions take 2.8–4.4 s from the cache.
• The cache holds 594,290 tokens, and weighted decode is within 0.6 % (in one run the chat category was 16 % slower).

Unchanged

MXFP4 (run-mxfp4.sh) is still there and still the most accurate option. --vision works in the default 65K mode and in --mode long (images with up to 245K tokens of context).
The previous release is one --image flag away (see the README).

Caveats, honestly

• System RAM: the coding and 512K modes pin 2.4 GiB of system RAM for the embedding table, and 16 GB of RAM is enough (tested). Below about 13.5 GiB of total RAM, --mode long-kv4 keeps the table on the GPU and caches 451,879 tokens. You still get 262K per request and the prefix cache. --mode long-512k needs the RAM.
• Cold reads: the first read of a 258K-token prompt still takes over two minutes. The cache helps from the second request on, as long as the start of the prompt stays the same (same system prompt, no timestamp at the top). usage.prompt_tokens_details.cached_tokens shows every hit.
• Vision isn't in the coding or 512K modes yet. Use --mode long --vision for images with long context.
• RAM/SSD cache tier: the experimental system-RAM tier for the prefix cache (--host-cache-gib, plus an SSD tier behind it) spills even more context off the card. For now it works with the FP8 --mode long; support for the 4-bit coding mode lands in the next release. The coding mode's 569,878-token GPU cache works today.
• The agent-session numbers and the accuracy columns were measured on pre-release builds of these configurations; the README has every number.

Our testbench is ancient, slow CPU, limited and slow memory (16GB), so your results will most probably be even better!

What's next: more GPUs

More R9700s are on the way, and we're going after tensor parallelism next.
Others are working on multi-card setups too, so expect some healthy competition on that front. Good for everyone running AMD at home!

More: Paiton

Currently models are crashing.

Also note I'm a little bit over my head right now as I'm just getting my feet wet with GPU and model architecture. Most of this help comes from AI.

It doesn't seem like there's support for my iGPU. I found this issue on GitHub which makes me believe I made a mistake in buying this mini PC.

Hardware:

Model I'm working with: Qwen 3.8 27b q4m

Minisforum UM890 Pro

Ryzen 9 8945HS

Radeon 780M iGPU (gfx1103)

64GB DDR5-5600

OCuLink PCIe 4.0 x4 available

Ubuntu

Flash Attention causes GPU memory faults, while disabling it causes rocBLAS/Tensile gfx1103 kernel errors.

Vulkan works, but is slow.

Building my own llama.cpp seems suboptimal if I'm wanting tried and tested Builds.

Considering selling and buying a mac studio.

Thoughts?

I've released WHIRL (Windows HIP Inference for RDNA LLMs), an open-source (Apache-2.0) LLM inference engine for the AMD Radeon AI PRO R9700 that runs natively on Windows. It's pure C++ and HIP. There's no WSL, no Linux VM and no llama.cpp underneath. You need only the AMD Adrenalin driver to run it. The HIP SDK is only needed if you build from source.

• GitHub: https://github.com/tsaipifong/whirl-llm
• Release (zip, ~28 MB): https://github.com/tsaipifong/whirl-llm/releases/tag/v0.1.0

The idea. WHIRL isn't a general engine that runs every model reasonably well. It's a specialised one: a small set of models on a specific GPU, run as close to the hardware limit as possible. This follows the philosophy of NInfer, and zynfer took the same idea to RDNA 4. WHIRL shares only the idea. The code base is its own.

What's in it

• Hand-written RDNA 4 kernels: WMMA prefill and MXFP4 MoE grouped GEMM.
• MTP + n-gram speculative decoding, on by default. The output is identical to plain greedy decoding, and we verify this bit-exact.
• An OpenAI-compatible server with prefix caching and a tiered KV cache (VRAM → RAM → SSD). A long shared system prompt is cached once. A conversation can be restored after a server restart.
• Vision (mmproj) support for the dense model.

Supported architectures are qwen35 (dense) and qwen35moe (MoE). Architectures are added one at a time, so this isn't meant to be a Qwen-only engine. It's just where we started.

Numbers vs llama.cpp (b11214, same GGUF, same prompts, greedy)

Setup: R9700 32 GB, Windows 11, Adrenalin 26.8.1. llama.cpp gets its fastest configuration on every row. Prefill uses the best -ub for each length. Decode takes the fastest of plain / MTP / MTP+n-gram.

WHIRL vs llama.cpp	Qwen3.8-27B Q4_K_M (dense, standard unsloth quant)	Swift-1.5 (Qwen3.8-27B fine-tune) MXFP4 (dense)	Ornith-1.5 35B-A3B MXFP4 (MoE, ~3B active)
Prefill 8k tok/s	1,689 vs 1,223 (1.38×)	3,278 vs 1,338 (2.45×)	10,858 vs 4,637 (2.34×)
Prefill 32k tok/s	1,479 vs 1,086 (1.36×)	2,595 vs 1,174 (2.21×)	7,978 vs 3,778 (2.11×)
Decode, Chinese coding prompts, WHIRL MTP+n-gram vs llama.cpp fastest	98.5 vs 56.0 (1.76×)	107.5 vs 60.8 (1.77×)	244.5 vs 118.8 (2.06×)
Decode, file-edit prompts, MTP+n-gram vs llama.cpp fastest	305 vs 137 (2.23×)	328 vs 145 (2.27×)	618 vs 239 (2.58×)
Decode, no MTP vs llama.cpp plain	34.4 vs 30.9 (1.11×)	37.5 vs 33.4 (1.13×)	166.7 vs 118.8 (1.40×)
Server, 4 concurrent users, aggregate tok/s	134 vs 59 (2.29×)	199 vs 66 (3.02×)	397 vs 192 (2.07×)
Warm TTFT, new chat sharing a 26k system prompt	0.14 s vs 0.55 s	0.12 s vs 0.54 s	0.07 s vs 0.29 s
Restore after server restart (from SSD)	0.74 s (N/A)	0.73 s (N/A)	0.30 s (llama.cpp: N/A)

Vision encode for a 1920×1088 image with Swift's mmproj takes 254 ms in WHIRL and 1,330 ms in llama.cpp mtmd. Both measurements are warm.

Where WHIRL doesn't lead by much, which is why I list it here:

• Decode without speculation on the dense models is only 1.1–1.18× faster. Plain decode is bound by memory bandwidth, and both engines are near the limit. Most of the decode gap comes from MTP + n-gram.
• Q4_K_M prefill with a very short prompt (88 tokens) is roughly a tie: 1.03×.
• Swift past 16k context is only 1.09× faster. On that model, MTP alone is currently a bit faster than MTP + n-gram (72.4 vs 69.5 tok/s). The n-gram cost model needs tuning.
• VRAM. The CLI uses 3–5 GiB more VRAM than llama.cpp on the dense models. The server pre-allocates the rest of VRAM as KV pool, so it will show about 30 GB in Task Manager.
• Hardware. It runs only on the R9700, one GPU, on Windows 11. A Ryzen AI Max+ 395 (Radeon 8060S) version is planned.

About my setup. My R9700 is connected as a USB4 eGPU. That slows model loading and the RAM/SSD cache restores in both engines. It doesn't affect prefill or decode, which stay in VRAM. On a direct PCIe slot, restores should be faster than what I measured.

The full methodology, every table, the exact commands and the raw-data description are in docs/benchmarks.md.

Models

WHIRL runs standard GGUFs of the supported architectures. The best results come from the MXFP4 quants I published:

• Swift-1.5-Qwen3.8-27b MXFP4: https://huggingface.co/tsaipifong/Swift-1.5-Qwen3.8-27b-MXFP4-GGUF
• Ornith-1.5-35B-A3B MXFP4: https://huggingface.co/tsaipifong/Ornith-1.5-35B-A3B-MXFP4-GGUF

Both also run in llama.cpp.

Quick start

1. Install AMD Software: Adrenalin Edition 26.8.1 or newer.

2. Download the zip from Releases and unzip it.

3. Run:

   whirl devices whirl chat path\to\model.gguf whirl-server path\to\model.gguf

The server exposes an OpenAI-compatible API on localhost. The first run tunes kernels once: about 1.5 minutes for the 27B models and a few seconds for the MoE. The result is cached.

The exe is unsigned, so SmartScreen will ask the first time ("More info → Run anyway"). If Smart App Control is on, it may block unsigned apps. The README explains how to check.

Development note

WHIRL was designed, implemented, optimised and benchmarked with Claude Opus 5.5 (Anthropic) working under my direction. Every result above is measured, and outputs are checked bit-exact against plain greedy decoding. The docs, in English and Traditional Chinese, include a long pitfalls list of the Windows/HIP/RDNA 4 problems we ran into. I hope it saves someone time.

Feedback, bug reports and measurements on direct-PCIe R9700s are very welcome.

TL;DR: A drop-in sageattention build for RX 9070 / 9070 XT on Windows. The common case (fp16, head_dim 128) runs on an fp8 attention kernel I wrote by hand in HIP. Against the existing gfx12 port (SageAttention PR #368), each attention call is 1.07–1.21× faster non-causal and 1.72–1.96× faster causal. In ComfyUI with Krea2 that works out to a 5–14% faster sampling step than PyTorch SDPA. Install the wheel, keep using --use-sage-attention, done. Caveats below, and there are a few, so please read them.

Repo + wheel: https://github.com/IxMxAMAR/SageAttention-RDNA4

Why I did this

I run ComfyUI on a 9070 XT under Windows. SageAttention on RDNA4 works thanks to the gfx12 port in PR #368, so huge credit to DELUXA for that. But I wanted to see how far this card can actually go. So I started with a Triton kernel, lost to #368 on most shapes, spent a long time figuring out why, and ended up writing the kernel by hand.

What finally made the difference:

• 64 keys per loop iteration instead of 16.
• Computing Q·Kᵀ transposed. The result then lands in exactly the register layout the next matrix multiply needs, so the attention weights never leave registers.
• One scale per token for Q, K and V, instead of one per 64-token block.

The whole story, including everything that didn't work (most things didn't, lol), is in the repo's docs/JOURNEY.md and docs/FINDINGS.md.

Who this is for, right now

• ComfyUI users on an RX 9070 or RX 9070 XT, on Windows. This is what I've tested and what I use daily.
• Anyone else who calls sageattn from PyTorch on the same card should be able to use it too; it's the same package API. But I haven't tested any other apps or workflows yet. If you try it outside ComfyUI, tell me how it goes.

Speed (image 1)

Same tensors, same process, interleaved runs, full call including quantization:

• Non-causal (what image/video diffusion uses): 1.07–1.21× faster than #368.
• Causal: 1.72–1.96× faster.
• These are kernel-level numbers, and kernel-level isn't what you feel. Image 2 is what you feel.

In an actual ComfyUI render (image 2)

Krea2, 8 steps, fp16, called exactly the way stock ComfyUI calls it:

Backend	1 MP	2 MP
PyTorch SDPA	1.03 s/step	2.39 s/step
SageAttention 1.x	1.00 s/step	2.17 s/step
This build	0.97 s/step	2.06 s/step

Attention itself is about 2× faster than SDPA per call. But most of a step is the model's matmuls, so the end-to-end gain is 5–14% over SDPA and 2–5% over SageAttention 1.x. The gain grows with resolution.

Accuracy, honestly (image 3)

Everything is fp8 with per-token scales. I measured against exact attention on real Q/K/V captured from inside Krea2 and Flux2-Klein:

• On "calm" layers, #368 is more precise. It uses int8 for Q·K, which has finer steps than fp8.
• On layers with a few extreme keys, #368's error blows up, by up to ~200×. Krea2's first block is one such layer. One outlier wrecks the scale for its whole 64-token block. Per-token scaling never has that bad case.
• smooth_k is on by default and should stay on. It made my kernel more accurate on 10 out of 10 real captures.

In my Krea2 tests the images were closer to full-precision attention than with SageAttention 1.x. I can't tell them apart by eye. Like any quantized attention, though, your images won't match SDPA pixel for pixel: a turbo 8-step sampler amplifies tiny differences into different details.

Caveats (please read)

• RX 9070 / 9070 XT (gfx1201) only. On any other GPU it falls back to #368's kernels: no harm, no gain.
• Windows only, as far as testing goes. Linux is untested.
• The prebuilt wheel needs exactly Python 3.12 + PyTorch 2.13.0+rocm10.0.0**.** The compiled parts are tied to that PyTorch version. Anything else means building from source, which takes about 50 minutes.
• fp16 models only for now. bf16 models silently fall back to #368's kernel. Still works, just no speed-up from my kernel.
• head_dim 128 only. That covers Flux-style models such as Krea2, Flux and Klein. Other head sizes fall back.
• Tested end to end on Krea2 only. The kernel accepts the input layout video models like Wan use, and it's checked bit for bit, but I haven't benchmarked a video workflow yet.
• Very small images (well under 1 MP) can come out a hair slower than SageAttention 1.x. The win is at normal and high resolutions.
• No attention masks. Masked calls go to the fallback.
• LLMs: only prompt processing (prefill) through PyTorch would benefit. Token-by-token generation isn't covered, and llama.cpp / LM Studio don't use this at all.
• One bug already found and fixed. A pre-release build could turn a render black when a single value got huge enough to overflow fp8. It's fixed in this release, and every conversion in the kernel is now stress-tested with extreme values.

Install

1. Close ComfyUI.
2. Download the wheel from the GitHub release.
3. path\to\ComfyUI\venv\Scripts\python -m pip install --no-deps sageattention-2.2.0+amd.gfx12.1-cp312-cp312-win_amd64.whl (portable ComfyUI: use python_embeded\python.exe instead)
4. Start ComfyUI with --use-sage-attention like always.

--no-deps matters: it stops pip from touching your PyTorch. To switch my kernel off without uninstalling, set SAGEATTN_SK1_BACKEND=0.

What's next

• bf16 support, so bf16 models get the speed-up too. This is the big one for coverage.
• int8 Q·K with per-token scales, aiming for #368's precision on calm layers and no blow-ups on outlier layers. Already in the works.
• head_dim 64, for SDXL-era models. Also in the works.
• An end-to-end test on a video workflow (Wan), since the layout support is already in.
• An experiment to overlap the two halves of the kernel, which might be worth another few percent. A quick test decides whether it's worth building.

If you've got a 9070 / 9070 XT, I'd love numbers from your setup: your model, resolution and s/it before/after. Bug reports are very welcome too. Credit to the SageAttention team (thu-ml) and to DELUXA for #368; this builds directly on their work. Apache-2.0.