gfx1100 (RX 7900 XTX, ROCm 6.4.1): dflash_server ~2x slower than llama.cpp, DDTree speculation never engages, SIGSEGV at long context

[L0GYKAL]

Environment

• GPU: AMD Radeon RX 7900 XTX (gfx1100, RDNA3), 24 GB
• ROCm: 6.4.1 (rocm/dev-ubuntu-22.04:6.4.1 builder + runtime)
• lucebox-hub: commit 96100834eca003c2bc7cf76eb943b95fcd6d06a0 (+ pinned llama.cpp submodule 30c9d7da…)
• Build: DFLASH27B_GPU_BACKEND=hip, DFLASH27B_HIP_ARCHITECTURES=gfx1100, SM80_EQUIV=OFF (Phase-1 HIP, matches the default in Dockerfile.rocm)
• Model: unsloth Qwen3.6-27B-Q4_K_M.gguf (target) + Lucebox/Qwen3.6-27B-DFlash-GGUF/dflash-draft-3.6-q4_k_m.gguf (drafter)
• Built/run in rootless podman with --device /dev/kfd --device /dev/dri --group-add keep-groups.

Summary

On gfx1100, dflash_server is ~2× slower than upstream llama.cpp for the same model on the same GPU, its DDTree speculative decode never engages (silently falls back to autoregressive decode), and it SIGSEGVs at long context. Filing in case gfx1100 is intended to be supported — or to confirm the HIP speculative path is gfx1151-only.

Benchmarks — Qwen3.6-27B Q4_K_M, f16 KV cache, full GPU offload

Engine	Decode	Prompt eval	56k-token prefill
llama.cpp (Vulkan/RADV, b64739ea3)	36.4 tok/s	67 tok/s	prefill 641 tok/s, decode 32 tok/s, stable (20.8 GB)
dflash_server (HIP, Phase-1)	17.6 tok/s	—	SIGSEGV during prefill

Same 300-token generation prompt, temperature matched, both with f16 KV.

1. DDTree speculation never engages on gfx1100

Started with the matched drafter and DDTree on:

dflash_server /models/Qwen3.6-27B-Q4_K_M.gguf \
  --draft /models/dflash-draft-3.6-q4_k_m.gguf \
  --ddtree --ddtree-budget 8 --cache-type-k f16 --cache-type-v f16

Startup confirms draft = …dflash-draft-3.6…, ddtree = ON, ddtree_budget = 8, but every request logs:

[ar-decode] tokens=300 time=17.06 s speed=17.58 tok/s
[server] chat DONE … decode=17.1s(17.6tok/s)

i.e. plain autoregressive decode — identical speed with or without the drafter. No DDTree acceptance stats, no error. It looks like the draft-tree verify path (the rocWMMA flashprefill kernels) isn't available on gfx1100, so speculation silently no-ops.

2. SIGSEGV at ~56k context

A 56k-token prompt crashes during prefill:

…/ggml-hip/libggml-hip.so.0(+0x2729b2)
…ggml_backend_graph_compute…
ggml_abort  →  exit 139 (SIGSEGV)

(llama.cpp handles the same 56k prompt fine on the same card.)

3. SM80_EQUIV=ON (rocWMMA Phase-2) won't compile against ROCm 6.4.1

Enabling DFLASH27B_HIP_SM80_EQUIV=ON to get the flashprefill kernels fails:

server/src/flashprefill_kernels.hip.cu:452:47: error: cannot assign to return value
    because function 'operator[]' returns a const value
server/src/flashprefill_kernels.hip.cu:493:34: error: (same)
rocwmma/rocwmma_impl.hpp:96:16: error: non-const reference cannot bind to vector element

ROCm 6.4.1 ships rocWMMA 1.7.0, whose fragment operator[] is const-returning; the kernel assigns through it. So Phase-2 is unavailable on this stack, and (per #1) that seems to take DDTree with it.

Questions

1. Is gfx1100 (RX 7900 XTX) a supported target, or is the HIP speculative/flashprefill path intended for gfx1151 (Strix Halo) only?
2. Is there any way to get DDTree speculation working on gfx1100, or is autoregressive Phase-1 the ceiling there?
3. Any plan to update flashprefill_kernels.hip.cu for rocWMMA 1.7.0 (ROCm 6.4.x) so SM80_EQUIV=ON builds on RDNA3 consumer cards?

Happy to provide full logs or re-run with other flags.

[weicj]

PR #456 should fix the rocWMMA issue in point 3. Please rebuild on current main and retest. If it still fails, please also try the latest ROCm 7.2.x (there's also some performance improvement over ROCm 6.x).

[L0GYKAL]

Update — retested with PR #456 on a 7900 XTX (gfx1100)

Thanks for the pointer to #456. I rebuilt from current main and tested both the compile fix and the runtime paths. Summary: #456 fixes the build, but only on ROCm 7.2.x, and the two runtime issues (DDTree decode + long-context) remain.

1. Compile (rocWMMA FlashPrefill)

• ROCm 6.4.1 + SM80_EQUIV=ON → still fails identically: rocwmma_impl.hpp:96 operator[] returns const, won't compile.
• ROCm 7.2.4 + SM80_EQUIV=ON → ✅ compiles cleanly — flashprefill_kernels.hip.cu.o and dflash_server link with only warnings.

So the fix lands, but the floor is now ROCm 7.2.x — 6.4.1 users are still blocked. Could the README/build matrix note the minimum ROCm for SM80_EQUIV=ON?

2. DDTree decode still no-ops on gfx1100

Even with the kernels compiled (7.2.4 build) and the matched dflash-draft-3.6 drafter loaded (--draft … --ddtree --ddtree-budget 8), decode stays plain [ar-decode] at ~19 tok/s — unchanged from before, and ~half of llama.cpp-vulkan's 36 tok/s on the same model (Qwen3.6-27B Q4_K_M, f16 KV, full offload). No DDTree fallback message is logged. Is DDTree decode expected to engage on gfx1100, or is it gated to gfx1151?

3. Long context: SIGSEGV gone, now OOM

The previous ggml-hip SIGSEGV at ~56k is fixed — a 49k-token prefill now runs. But it then OOMs (cudaMalloc failed: out of memory / CUBLAS_STATUS_ALLOC_FAILED) on the 24 GB card with both f16 and q4_0 KV. The log shows pflash=false.

4. PFlash compression — couldn't enable it (drafter loading)

I tried --prefill-compression auto --prefill-drafter <gguf> to get PFlash engaged, but neither drafter on hand works:

• Qwen3-0.6B (unsloth GGUF) → tensors failed to load in the dflash loader.
• dflash-draft-3.6 → rejected: missing tokenizer.ggml.tokens metadata.

Could you point to a known-good --prefill-drafter GGUF (one that both loads in dflash's loader and carries full tokenizer metadata)? That would let me actually exercise PFlash compression for the long-context case.

Environment

• AMD Radeon RX 7900 XTX (gfx1100 / RADV NAVI31), 24 GB, rootless podman
• Model: Qwen3.6-27B-Q4_K_M (MTP), full GPU offload
• Built from main @ 1b11c50d, SM80_EQUIV=ON, ROCm 7.2.4 (Ubuntu 24.04 base)

Happy to run more tests if useful.

[weicj]

Point 1: For now, ROCm 7.x is the recommended stack for AMD testing.

Point 2: DDTree is not gated to gfx1151; it is expected to cover RDNA3 as well. Please check:

• whether the request was explicitly greedy (temperature: 0), since sampled requests fall back to AR by default unless sampled-verify is enabled
• whether startup printed warning: feature mirror init failed, spec decode will use AR fallback

Point 3: OOM with Q4_0 KV on the plain AR path is unusual, but it can still happen depending on the configured context/window and prefill workspace size. Please try:

• shrinking --max-ctx to 65536, and also reducing the prefill chunk / ubatch size (--chunk, or DFLASH27B_PREFILL_UBATCH=256 / 128 / smaller)
• if possible, minimizing desktop / display usage on the 7900 XTX while testing (for example by driving the display from the iGPU instead)

Point 4: For --prefill-drafter, please try the exact unquantized BF16 GGUF documented here:

• https://huggingface.co/unsloth/Qwen3-0.6B-GGUF/blob/main/Qwen3-0.6B-BF16.gguf