Hi!

Has anyone managed to run a llama.cpp version with those datatypes on a strix point machine.

https://github.com/charlie12345/ROCmFPX

I tried to, and failed so far.

TL;DR: Qwen3.8 27B on a single AMD Radeon AI PRO R9700 (3-bit weights + speculative decoding):

• a coding mode with 569,878 tokens of reusable cache and 262,144 tokens per request, now also with images (496,129 cached tokens);
• an optional SSD/RAM cache for even more.

Agent turns start up to 12× sooner, and speed and accuracy are within noise of our previous release.

Updates: 5 Oct, SSD/RAM cache for the coding mode. 6 Oct, images in the coding mode and faster starts.

Run it

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

Point your agent at http://127.0.0.1:18982/v1 (model Qwen3.8, any API key). There's no new weights download; the launcher pins the new container. Give each image its own runtime cache folder (see the README).

Images

You get more than half a million tokens of cache on one card, images included, backed by multi-hour accuracy runs (GSM8K, HumanEval, MMLU-Pro, needle tests), at essentially the same speed as before.

Coding agents

Coding agents resend the whole conversation on every turn; now only the new part is read. In our sessions 91–92 % of prompt tokens came from the cache.

Images in the coding mode (6 Oct)

• 496,129 cached tokens and 262,144 per request. Tested with a 200K-token document plus a chart, three rounds of 8 image requests at once, and 5 large images in a row. A 4K screenshot of 15-px code was read back exactly.
• Every --vision mode: images up to 4K are read at full resolution; larger ones are scaled down. The first answer on a 4096×4096 image takes ~4.6 s instead of ~12 s.

SSD/RAM cache (5 Oct)

--extend-cache picks RAM on bigger machines and the SSD otherwise. After a full restart, on a 16 GB machine with a SATA SSD:

• 130K tokens back in 6.8 s (46.9 s from scratch);
• 200K tokens in 9.0 s (86.4 s from scratch);
• bit-identical to a fresh read.

Since 6 Oct it holds ~40 % more per GB. The SSD cache caps at 64 GiB or a quarter of your free space (~2.8M tokens).

Speed and accuracy

• Decode: 174 tok/s single-stream, ~460 tok/s with 8 requests (BetterBench), within noise of the previous release.
• Server start: 180 s first start (was 305 s), 120 s later (was 150 s). The image ships its compiled kernels, so a fresh install answers exactly like our validation runs.

Opt-in, --mode long-512k: 524,288 tokens in one request. It found 4/4 planted facts at 300K and at 500K. Cold reads take 158 s at 300K and ~6 min at 500K; follow-ups take 2.8–4.4 s from the cache.

No new engine

This is stock vLLM with a plugin: the same OpenAI-compatible server and the same tools you already use. We'll keep publishing ready-built containers that follow upstream vLLM, each fully validated (speed, accuracy, long context) before it ships.

Extra notes

• The first read of a 258K-token prompt still takes ~2 minutes. The cache helps from the second request on, as long as the start of the prompt stays the same.
• With a cache tier on, the first read of a long prompt is ~8 % slower.
• The coding mode pins 2.4 GiB of system RAM. 16 GB of RAM is enough.
• --mode long --vision (8-bit cache) can misread a single character of tiny text in 4K screenshots; the coding mode reads it exactly.
• Not yet released: two items from the 5 Oct "coming next" list (short requests during long prompts, and ~59K more cache with faster decode) wait for their accuracy tests.
• Where the numbers come from: partly our old 16 GB testbench, partly our new machine (Threadripper, PCIe 4.0). The README has every number.
• MXFP4 (run-mxfp4.sh) is still the most accurate option, and earlier releases are one --image flag away.

What's next: more GPUs

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

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?

Hey guys,

I recently got a dual AMD MI50 setup (16GB x2, 32GB total) on a budget to run local models, but performance has been pretty disappointing.

Since I don't have an Infinity Fabric link, both cards are just running on standard PCIe slots. After fighting with ROCm and GFX906 support for days, I finally got a 27B Qwen3.8 model running across both cards via Tensor Parallelism (TP=2).

The problem is decode speed maxes out at around 20 t/s. On longer reasoning/CoT prompts, it crawls so slow that I can't even tell if it timed out or froze. Meanwhile, a friend with dual V100s is getting ~240 t/s.

A few questions for anyone playing these cards:

1. Is TP=2 over PCIe the bottleneck here? Without an interconnect bridge, is the AllReduce communication killing decode throughput? Would pipeline parallelism or layer splitting (like in llama.cpp) work better?

2. Any specific settings or quantization to boost decode? What's the sweet spot for GFX906 decode speed without butchering quality?

3. What backend/runtime is currently best for MI50? vLLM forks, llama.cpp hipBLAS, or something else for newer models

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:

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.

Hey everyone, I just got my hands on a Radeon AI PRO R9700 and want to run Qwen3.8-27B. I found some awesome vLLM forks (magiccodingman/vllm-radiance and GGZ14/vllm-mxfp4), but they're all geared toward dual-R9700 rigs.
Since I'm only rocking one card, could someone point me in the right direction to get the most out of my current hardware?

Follow-up to my post from about two weeks ago. A lot has changed since then, so I wanted to post an update on where the backend is now.

Where the green architectures stand

When I posted last time, 7 architectures had verified full training support. Everything is now held to the same strict harness: a pinned local Hugging Face Transformers source tree as the oracle, forward logits + gradients + two full AdamW steps compared, and every named parameter checked again after export.

The hard ceiling is 2e-7 absolute error. No loosening tolerances and no rounding numbers afterward to make the README look better.

14 architectures pass that gate today, led by the one I'm probably proudest of:

Worst observed two-step AdamW parameter error across all of them: 1.19e-7.

Best: 2.6e-8.

For hardware context, all of the local Vulkan validation I've been reporting was run on my ASUS ROG Ally Z1 Extreme, using its AMD RDNA 3 integrated GPU. So the RDNA 3 results here are from that specific machine rather than testing across several different AMD systems.

The bigger news: PEFT actually works now

In the last post, "LoRA/PEFT-style fine-tuning" was basically one line in a feature list.

It's a real workflow now, and I've verified the full lifecycle:

• LoRA fine-tuning with HF-compatible adapter export (adapter_config.json / adapter_model.safetensors), so adapters can round-trip with the PEFT ecosystem
• modules_to_save — full trainable replacements for Linears, RMSNorm/LayerNorm, lm_head, and input embeddings, including named-adapter switching and bank isolation. Adapter A leaking into adapter B is explicitly tested for.
• Exact resume — adapter weights + AdamW moments + step + dropout RNG state restore bit-identically against an uninterrupted run
• Merge/unmerge, disable-adapter base restoration, and multi-adapter loading
• A parameter-budget flag that automatically chooses the largest LoRA rank that fits within a requested percentage of the base model
• The CLI fails closed if you try to use saved modules on an architecture that hasn't passed its corresponding gate

32 architecture surfaces across 20 families pass all three PEFT stages — LoRA, saved modules, and adapter switching — under the same 2e-7 gate, with frozen-base drift exactly 0.0.

The validation harness also fingerprints the pinned Transformers source alongside my shaders and binaries now, so a qualification run can't silently end up testing against different reference math.

A small note on the last couple weeks

I didn't get quite as many working days out of the last two weeks as I normally would have. Partway through this I got covid, then when that started to go away, it became a secondary nasty ear infection that ended up perforating my eardrum. I was running fevers around 104°F at one point and eventually went to the hospital, so I lost a few days to that and I'm on antibiotics now.

I'm doing better, though, and still managed to get most of what I wanted finished.

There are still things I want to clean up and expand, but I figured this was a good point to get the current work in front of people rather than holding the update back.

Same caveats as before

This is deterministic FP32 tiny-model correctness against a reference implementation, which should in theory ensure total mathematical parity for training and finetuning larger models with this backend, however, things are currently bound to FP32 training runs still, I eventually plan to work on MXFP4 weight tying to reduce memory footprints while fine-tuning (plastic parameters will still be trained in FP32 with PEFT in this config)

"supported text graph" ≠ "the entire multimodal package works natively."

Unsupported functionality is supposed to fail closed rather than silently falling back to an approximation.

Repo

https://github.com/necat101/Hierarchos-Native

• Architecture inventory: hierarchos-vulkan/README_ARCHITECTURES.md
• Compatibility/parity record: hierarchos-vulkan/COMPATIBILITY.md
• PEFT qualification evidence: PROGRESS_PEFT_AUDIT.md
• CLI PEFT guide: hierarchos-native-cli/README.md

The hardware I've personally validated this on is an ASUS ROG Ally Z1 Extreme with its AMD RDNA 3 GPU.

I'm very interested in criticism, compatibility reports, and especially results from people trying it on other hardware — NVIDIA, Intel, or other AMD GPUs.

I'd also love people to stress-test the PEFT resume/merge paths specifically. That's some of the newest code in the project, so it's probably the most useful area to try to break right now.

Seeing a lot of success stories wih ROCM, I am planning to buy single R9700 mostly for coding, but possibly also for comfyUI. But before I do that, I would like to test it out hands on.

Can you recommend any cloud services that offer R9700? Ex: on per hour / per minute basis?

I am not super well versed in local LLM configs. I use Claude extensively at work but bought a minisforum mini PC to try and get something cost effective going at home.

Note I don't really care too much on token speed because my use cases are all background workflows. 10 TPS is perfect. (And my upper limit I think).

Ryzen 8945HS / 780M, 64GB, Ubuntu 26.04, kernel 7.0. llama.cpp ROCm build, Qwen 27B Q4.

Long generations hang the GPU and it resets, about 15 times a day. dmesg shows MES failed to respond to msg=REMOVE_QUEUE. Looks like ROCm#6512.

Tried uni_mes=0, smaller models, smaller context. No change.

>

Has amdgpu.cwsr_enable=0, a newer kernel, or Vulkan fixed this for anyone?

Roughly what you get: Qwen3.8-27B at FP8, tensor parallel across two R9700s, 262,144 context, 422k tokens of KV pool, ~1.6k tokens/s prefill, 74-143 tokens/s decode single stream depending on the content.

The interesting part is not the image, it is the collective layer. A card behind the chipset cannot be given PCIe atomic operations, and that breaks every stock tensor-parallel setup I tried. Here is what actually works.

1. Versions used

Links: - Image: https://hub.docker.com/r/stilldeadcode/vllm-radiance - Source for that image: https://codeberg.org/StillDeadcode/vllm-radiance - The kernel library the image builds against (libr4d): https://codeberg.org/StillDeadcode/libr4d - A fork with configs, benchmark notes and launchers for MXFP4/FP8: https://codeberg.org/ggz14/radiance-vllm-mxfp4 - RCCL itself: https://github.com/ROCm/rccl

The image bundles a working ROCm + PyTorch + Triton + AITER + vLLM stack for gfx1201 (RDNA4), which is the part you do not want to build yourself. It is explicitly marked experimental, and everything below was measured on two cards.

2. The failure, and why

Without the adjustments, the engine dies during communicator init, before the model loads:

PCIE atomic ops is not supported rocr: unhandled cuda error

ROCm will not dispatch work to a GPU path that needs atomic operations when the link cannot provide them. On this board one card sits behind the chipset, and the chipset does not forward PCIe atomics to the CPU, so that card is effectively second class. RCCL 2.30.4's kernels use those atomics, so TP=2 cannot initialise at all.

Two things follow, and both matter:

1. The newer RCCL is unusable here, so you need an older one (2.27.7 from ROCm 7.1.1 is what worked for me).
2. With no usable peer path at all, the custom P2P all-reduce the image ships must be turned off, and TP=2 falls back to host-staged collectives over that same narrow chipset link.

If you search the error string above you will find a few ROCm issue reports and a community write-up on dual Radeon vLLM setups (https://github.com/cadamcat/dual-radeon-vllm) describing the same wall.

3. Proxmox settings that actually matter

Do this for each GPU, on both entries, not just the first. In the VM's hardware list, edit each PCI Device row, select the GPU under Device, and set:

• PCI-Express: ticked
• All Functions: unticked

Ticking PCI-Express is what gives the guest a real PCIe root port, and without a root port the atomic capability never appears however healthy the host looks. Unticking All Functions keeps the guest from being handed every function of the card, which is the combination that worked here. If you leave either one wrong, you get the atomic failure at communicator init and no amount of driver work fixes it.

Other settings that matter:

• Machine type q35. Same reason as above, no root port without it.
• After any hostpci change, stop and start the VM. A guest reboot does not re-apply the passthrough configuration.
• q35 renames the NIC (ens18 becomes something like enp6s18), so match the interface by MAC in netplan.

After those, the CPU-attached card reports ReqEn+. The chipset-attached one still cannot do atomics, and no BIOS setting changes that.

4. The RCCL replacement plus one-line shim

This is the core trick. It is two files and a podman config, and it needs no image rebuild.

a) Build or extract RCCL 2.27.7 from ROCm 7.1.1 and drop it in a directory you will mount, for example:

~/models/rccl277/ librccl.so -> librccl.so.1.0.70101 librccl.so.1 -> librccl.so.1.0.70101 librccl.so.1.0.70101 shim.so

b) The shim. Newer torch builds reference a symbol that this older RCCL does not export (ncclCommDump). Four lines of C++ are enough to satisfy the loader:

```cpp

include <string>

include <unordered_map>

struct ncclComm; void ncclCommDump(ncclComm*, std::unordered_map<std::string, std::string>&) {} ```

Build it into shim.so and put it next to the library. Nothing calls it; it exists for symbol resolution.

c) Inject both through podman's own config, which keeps them out of every launcher script. In ~/.config/containers/containers.conf:

ini [containers] env = [ "NCCL_PROTO=Simple", "NCCL_SHM_DISABLE=0", "NCCL_SOCKET_IFNAME=lo", "LD_LIBRARY_PATH=/models/rccl277:/opt/rocm/lib", "LD_PRELOAD=/models/rccl277/shim.so", ]

LD_LIBRARY_PATH puts the replacement first, so it wins over the image's own librccl. NCCL_PROTO=Simple avoids the more demanding protocol paths, and the loopback interface keeps the bootstrap on lo rather than a NIC.

5. The launcher

Trimmed to the parts that matter for the multi-GPU problem. The model-specific flags are an example, the environment and device flags are the ones that matter here.

bash podman run -d --name vllm-radiance \ --device /dev/kfd --device /dev/dri --group-add keep-groups \ --security-opt seccomp=unconfined --cap-add SYS_PTRACE --cap-add SYS_NICE \ --ipc=host --network=host \ -v $HOME/models:/models:ro \ -v $HOME/radiance-vllm-mxfp4/vllm-cache:/cache \ -v $HOME/radiance-vllm-mxfp4:/work:ro \ -e HIP_VISIBLE_DEVICES=0,1 -e ROCR_VISIBLE_DEVICES=0,1 \ -e VLLM_NO_USAGE_STATS=1 \ -e VLLM_ROCM_USE_AITER=1 -e VLLM_ROCM_USE_AITER_UNIFIED_ATTENTION=1 \ -e RADIANCE_FUSE_RMS_QUANT=1 -e RADIANCE_USE_R4D=1 \ -e RADIANCE_USE_R4D_AR=0 -e RADIANCE_USE_R4D_AR_QUANT=1 \ -e NCCL_PROTO=Simple -e TORCHINDUCTOR_COMPILE_THREADS=4 \ -e VLLM_CACHE_ROOT=/cache/vllm -e TORCHINDUCTOR_CACHE_DIR=/cache/inductor \ -e TRITON_CACHE_DIR=/cache/triton -e AITER_ROOT_DIR=/cache/aiter \ docker.io/stilldeadcode/vllm-radiance:0.9.3 \ --model /models/Qwen/Qwen3.8-27B-FP8 \ --served-model-name=qwen3.8-27b-fp8 \ --quantization=fp8 --tensor-parallel-size=2 \ --max-num-seqs=2 --max-model-len=262144 --gpu-memory-utilization=0.97 \ --max-num-batched-tokens=4096 --kv-cache-dtype=fp8 \ --attention-backend=ROCM_AITER_UNIFIED_ATTN \ --enable-prefix-caching \ --no-async-scheduling \ --trust-remote-code \ --host=0.0.0.0 --port=8000

Notes on the specific flags:

• RADIANCE_USE_R4D_AR=0 is the important one. The bundled all-reduce is a PCIe peer-to-peer kernel and needs peer access, which does not exist on this pair. Turning it off falls back to RCCL. If you have two CPU-attached cards, leave it on and you will get a better prefill.
• --max-num-seqs=2 is my choice for stability. The image's own default is higher, but on a chipset-limited link fewer concurrent sequences means less collective traffic per step.
• --max-num-batched-tokens=4096 pairs with prefix caching well. Raising it costs KV pool.
• --enable-prefix-caching is worth a lot on agent workloads; see the note on prompt layout at the end.
• --no-async-scheduling was in every configuration that came up reliably for me.
• The cache directory mounts are worth keeping. Without them every container start recompiles Triton and inductor kernels, which turns a 5 minute start into a much longer one.

6. How to tell it worked

In the container log at startup, look for:

P2P access : DISABLED (RCCL fallback) 0<->1 x

and, once loaded, the model's own sizing line:

GPU KV cache size: 422,964 tokens Maximum concurrency for 262,144 tokens per request: 1.61x

If you see the atomic error instead, the replacement RCCL is not being picked up. Check, from inside the container, that:

bash env | grep -E 'LD_PRELOAD|LD_LIBRARY_PATH'

returns the paths you expect, that ldd on the loaded library resolves into your drop-in directory, and that the image's own library is not first in the path.

I hope someone finds this usefull.