feat(rocmfp4): backend for the ROCmFP4/ROCmFPx 4-bit formats on AMD RDNA3.5 APUs (Strix Halo) - #11636
feat(rocmfp4): backend for the ROCmFP4/ROCmFPx 4-bit formats on AMD RDNA3.5 APUs (Strix Halo)#11636walcz-de wants to merge 8 commits into
Conversation
A llama.cpp-fork backend carrying the ROCmFP4 / ROCmFPx weight-quantization
formats (ggml types 100-107) for AMD RDNA3.5 APUs - gfx1150/1151/1152/1153,
which is Strix Point, Strix Halo, Gorgon Point and Gorgon Halo. Stock llama.cpp
rejects those tensor types, which is precisely why this is a separate backend
rather than a flag on llama-cpp.
Follows the bonsai pattern: the fork's additions are model weight types decoded
inside libllama and therefore transparent to the shared gRPC server, so
backend/cpp/llama-cpp's grpc-server.cpp, CMakeLists.txt and Makefile are reused
verbatim and only the fetched repo and commit are swapped. No grpc-server.cpp
allow-list patch is needed (unlike turboquant, whose KV-cache types do need one).
Measured on a Radeon 8060S (gfx1151) with the published Strix Halo weights:
qwen35 27B Q4_0_ROCMFP4_STRIX 13.74 GiB pp128 335 t/s tg32 14.2 t/s
What the format buys on this hardware is memory, not arithmetic: 22% smaller at
perplexity parity with Q4_K_M. There is no FP4 matrix instruction on RDNA3.5 or
RDNA4 - only CDNA4 has one - so the matmul runs as int8 dot products and the win
is bandwidth. That belongs in the backend description rather than in the
footnotes, and llama-cpp stays the recommendation for everything else.
Deliberately NOT wired into engineNamePreferenceRules: an entry there would make
the gallery prefer FP4 variants on every AMD host, which is displacement rather
than addition. The guide covers this - a backend left out ranks below every known
engine, which is the intent for a niche format.
Backend sources: https://github.com/walcz-de/llama.cpp-ROCmFP4
Format and kernels ported from https://github.com/charlie12345/ROCmFPX (MIT).
Assisted-by: Claude:claude-opus-5 [Claude Code]
Signed-off-by: Stefan Walcz <stefan.walcz@walcz.de>
Completes the wiring the previous commit left out, so the backend is actually
discoverable, buildable and installable rather than just present on disk:
- backend/index.yaml: meta + rocm image entries (latest and development). The
description says plainly what the format does and does not buy - memory, not
arithmetic - and that llama-cpp stays the recommendation otherwise. Same tone
as the vllm-cpp entry.
- .github/backend-matrix.yml: hipblas/amd64 entry with the prebuilt gRPC base,
tag-suffix matching the index uri exactly.
- scripts/lib/backend-filter.mjs: path filter ABOVE the generic llama-cpp
suffix, plus the shared-input rule so a change under backend/cpp/llama-cpp/
retriggers this backend too. Omitting this has no effect on the PR that adds
the backend and silently breaks the next one.
- .github/workflows/bump_deps.yaml: nightly pin bump. The Makefile keeps the
ROCMFP4_VERSION?= form the bot greps for; empty Docker ARG values are unset
in the compile script instead, since an empty-but-defined variable would
otherwise beat Make's ?= and clobber the pin.
- core/gallery/importers/llama-cpp.go: preference-only, extending the existing
drop-in list. ROCmFP4 GGUFs carry ordinary .gguf names, so there is no safe
auto-detect signal - without an explicit preference an import stays llama-cpp.
Linux/AMD only, and deliberately so: the format's kernels are ROCm/HIP for
RDNA3.5, and there is neither ROCm nor the target hardware on Apple silicon.
No includeDarwin entry for that reason.
Still deliberately absent from engineNamePreferenceRules - see the previous
commit.
make test-ci-scripts: 41 pass, 0 fail.
Assisted-by: Claude:claude-opus-5 [Claude Code]
Signed-off-by: Stefan Walcz <stefan.walcz@walcz.de>
…-server)
The wrapper was modelled on bonsai, whose fork predates upstream renaming the
RPC binary. Building against a current tree therefore failed with
gmake[2]: *** No rule to make target 'rpc-server'. Stop.
Note this is ggml's distributed-inference RPC backend, not the gRPC server that
talks to LocalAI - the two are easy to confuse because the build flavour is
called "grpc" while the target it builds is "ggml-rpc-server". backend/cpp/
llama-cpp/Makefile already uses the new name; this now matches it.
Assisted-by: Claude:claude-opus-5 [Claude Code]
Signed-off-by: Stefan Walcz <stefan.walcz@walcz.de>
…model Closes the two gaps the adding-backends checklist calls out: the table-driven importer test (preference 'rocmfp4' swaps the emitted backend, default stays llama-cpp) and test-extra-backend-rocmfp4, which loads the published Strix Halo ROCmFP4 build - a weight quant only this fork decodes, following the bonsai pattern. The importer spec passes in isolation; the suite's live-HF specs for unrelated importers flake identically with and without this change. Assisted-by: Claude:claude-opus-5 [Claude Code] Signed-off-by: Stefan Walcz <stefan.walcz@walcz.de>
localai-org-maint-bot
left a comment
There was a problem hiding this comment.
The backend wiring needs a cleanup before this can run through CI:\n\n- In .github/workflows/bump_deps.yaml, the new walcz-de/llama.cpp-ROCmFP4 item is indented at the matrix level instead of under matrix.include. The following ds4 item is therefore left unexpectedly over-indented, making the workflow YAML structurally invalid.\n- backend/cpp/rocmfp4/patches/README.md is still the Bonsai/Prism copy: it identifies this as the bonsai backend, points at the PrismML prism fork, and describes Bonsai CI inputs. Please rewrite it for rocmfp4 and the walcz-de/llama.cpp-ROCmFP4 fork.\n\nDCO is also failing, but that attestation must be fixed by the human contributor.
… README Addresses the maintainer review on mudler#11636: - The bump_deps matrix item sat at matrix level instead of under matrix.include, leaving the following ds4 item over-indented and the workflow YAML structurally invalid. Re-indented to match its neighbours; YAML validated. - backend/cpp/rocmfp4/patches/README.md was still the Bonsai/Prism copy it was templated from. Rewritten for rocmfp4 and walcz-de/llama.cpp-ROCmFP4, including why this directory is expected to stay empty (the fork tracks the same upstream pin the shared gRPC server is written against). Assisted-by: Claude:claude-opus-5 [Claude Code] Signed-off-by: Stefan Walcz <stefan.walcz@walcz.de>
4696c51 to
3a0c54c
Compare
|
Thanks for the careful review — all three addressed:
|
localai-org-maint-bot
left a comment
There was a problem hiding this comment.
The earlier workflow, patches README, and DCO issues are fixed. One copied Bonsai artifact remains in the same helper: backend/cpp/rocmfp4/apply-patches.sh still says it applies Bonsai patches to the PrismML prism fork, and its success message says all bonsai patches applied successfully. Please rewrite those comments and the message for rocmfp4 / walcz-de/llama.cpp-ROCmFP4. This is documentation and diagnostic output only; no behavior change is needed.
…backend Last of the Bonsai template artifacts: the header comment still described applying Bonsai patches to the PrismML prism fork, and the success message said 'all bonsai patches applied successfully'. Rewritten for rocmfp4 and walcz-de/llama.cpp-ROCmFP4, including the expectation that patches/ normally stays empty since the fork tracks the same upstream pin as the shared gRPC server. Documentation and diagnostic output only, no behavior change. The two remaining bonsai mentions in the Makefile are deliberate comparisons to the pattern this backend follows, not leftover identity. Assisted-by: Claude:claude-opus-5 [Claude Code] Signed-off-by: Stefan Walcz <stefan.walcz@walcz.de>
|
Done — |
localai-org-maint-bot
left a comment
There was a problem hiding this comment.
The latest commit removes the remaining Bonsai/Prism text from apply-patches.sh and keeps the helper behavior unchanged. DCO passes. Good from my side, @mudler.
|
Review pass. Mechanically this is one of the more complete backend additions I have checked against Correctly wired, verified individually: matrix entry, The decision, for a maintainer: Secondary scope note: a whole new ROCm-only image for one published model, justified in the body but still a per-release CI cost. Smaller items:
DCO: |
… builds The backend shipped grpc-server only, so the FP4 formats could be consumed but not produced: usable weights were limited to the one published model. The fork already carries the quantization recipes and CLI arguments, they were simply never built. Adding --target llama-quantize to the grpc build and copying the binary out means LocalAI's own quantization service finds it -- it resolves the tool from PATH or its backend directory -- so /app/quantize can drive FP4 with its custom-type field, on any GGUF the user already has. While here, the dead paths a review flagged: rocmfp4-avx/avx2/avx512 and rocmfp4-cpu-all are built by no matrix entry (the single entry is hipblas amd64), the CPU-image branch in rocmfp4-compile.sh is unreachable for the same reason, and run.sh carried llama-cpp's Intel Level Zero block, which a ROCm-only image never needs. rocmfp4-fallback and rocmfp4-grpc stay -- the hipblas entry does build those. Also adds the backend-filter test case the suite was missing for this suffix, covering that a change under backend/cpp/rocmfp4/ resolves to this backend and rebuilds nothing else. node --test scripts/lib/backend-filter_test.mjs: 42/42. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01PHcZ8CTpZEoK3qgroNnkB4 (cherry picked from commit 2286566)
The build that produces llama-quantize had every CPU feature switched off, inherited from the portable fallback profile. That is the wrong default here: quantization is pure CPU work over every weight in the model, and this backend only targets gfx1150-1153 -- integrated GPUs whose CPU sits on the same package and is always Zen 4/5. AVX, AVX2, AVX512, FMA, F16C and BMI2 are therefore guaranteed present on any machine that can run this backend at all, so nothing portable is lost by using them. The portable rocmfp4-fallback build, which serves inference, is untouched. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01PHcZ8CTpZEoK3qgroNnkB4 (cherry picked from commit b784a35)
|
Thanks — that is a genuinely useful review, and the ordering of concerns is right. Everything below is pushed except where noted. On the supply chainOne correction about the mechanism, because it changes the shape of the decision rather than just the wording. The nightly job does not commit anything. So no unreviewed commit reaches a published image on a cron; a maintainer merges, exactly as for Your underlying point still stands on its own, though: with ~25 bump PRs a night that review can get thin, and "personal fork" is fair to weigh separately from "third-party fork". So, your call and we are fine either way:
Say the word and it is out in the next push. We would rather you picked than have us pre-empt a decision that spends your risk budget. On fork freshness, since it is the same topic: the fork sat on llama.cpp Scope: one model, or one architecture?Your note about a whole ROCm image for a single published model was fair as stated, and this push changes the fact underneath it. The backend shipped That is not just a CLI in a directory. LocalAI's own quantization service resolves the tool from Measured on a Radeon 8060S / Ryzen AI MAX+ 395 (gfx1151), not derived:
Output GGUFs carry The other half of the scope answer is hardware, not weights: gfx1150/1151/1152/1153 — Strix Point, Strix Halo, Gorgon Point, Gorgon Halo. Four shipping families on unified memory, which is exactly where smaller weights decide whether a model fits at all. CPU features in the grpc buildWorth calling out because it changes an existing build line: that build had every CPU feature off, inherited from the portable fallback profile. Quantization is pure CPU work over every weight, and this backend targets integrated GPUs whose CPU sits on the same package and is always Zen 4/5 — AVX, AVX2, AVX512, FMA, F16C and BMI2 are guaranteed present on any machine that can run this backend at all. So they are on now, and nothing portable is lost. The portable Measured limitations — please read this before mergingHaving produced weights with the shipped quantizer, we ran our own evaluation harnesses, and the result is not uniformly good. Reporting it here because the PR body currently quotes perplexity parity, and perplexity does not measure what broke. Base model
Generation is preserved; discrimination is not. Free chat is fine, deterministic coding is lossless, throughput is up ~11% single-stream. But choosing correctly between similar options degrades badly — and the classification row is the one that matters, because it uses no tools at all. So this is not a tool-calling quirk; it is the model's ability to pick.
We are not proposing a change from this — the format does what the PR says it does, and for generation-shaped and throughput-bound work it is a real win. But anyone pointing this backend at an agent should measure first, and that belongs in the documentation rather than in each user's afternoon. Happy to add a paragraph to the Related, and filed separately as an issue on our fork: Smaller items — all addressed
Two notes backDCO is green. The check completed at
Disclosure: prepared with AI assistance (Claude). All numbers above are from runs on the hardware named, not estimates. |
What this adds
A new backend,
rocmfp4, carrying the ROCmFP4 / ROCmFPx 4-bit weight-quantization formats (ggml types 100–107) for AMD RDNA3.5 APUs — gfx1150/1151/1152/1153, i.e. Strix Point, Strix Halo, Gorgon Point and Gorgon Halo. Stock llama.cpp rejects these tensor types (unknown type q4_0_rocmfp4_fast), which is precisely why this is a separate backend rather than a flag onllama-cpp.The formats originate in charlie12345/ROCmFPX (MIT). Because that tree predates the MMQ config refactor (#24127 in llama.cpp) by several weeks, the formats were ported onto a current llama.cpp base — including re-expressing the FP4 MMQ load-tile kernels against the new
ggml_cuda_mmq_config/sram_layoutarchitecture — and are maintained at walcz-de/llama.cpp-ROCmFP4, with copyright headers and attribution intact.Published weights this serves today: kingjones777/Qwen3.8-27B-ROCmFP4-STRIX-MTP-GGUF.
What the format buys — honestly
Memory and bandwidth, not arithmetic. There is no FP4 matrix instruction on RDNA3.5 or RDNA4 (only CDNA4 has one), so the matmul runs as int8 dot products. The win is ~22% smaller files than Q4_K_M at perplexity parity (5.8877 vs 5.8926 on wikitext-2 per the published card) — on unified-memory APUs that can decide whether a model fits at all.
Measured end-to-end on a Radeon 8060S (gfx1151, Strix Halo), 27B model, via this backend inside LocalAI:
spec_type:draft-mtp, n_max 4)llama-cppstays the recommendation for everything else; the backend description inindex.yamlsays so explicitly (same tone as thevllm-cppentry).Non-displacement, by design
engineNamePreferenceRules— an AMD entry would make the gallery prefer FP4 variants on every AMD host. Per the adding-backends guide, leaving it out ranks it below every known engine, which is the intent for a niche format.llama-cppdrop-in list, likeik-llama-cpp/turboquant): ROCmFP4 GGUFs carry ordinary.ggufnames, so there is no safe auto-detect signal. Without an explicit preference an import staysllama-cpp.rocmfp4installed, a plain Q4_K_M GGUF withoutbackend:still auto-loads onllama-cpp(backends=[llama-cpp vllm]— the new backend never enters the candidate list).Scope: Linux/AMD only, deliberately
The kernels are ROCm/HIP for RDNA3.5; there is neither ROCm nor the target hardware on Apple silicon, and NVIDIA already has native FP4 paths upstream. Hence no
includeDarwinentry and a singlehipblasmatrix entry — stated here rather than omitted silently, per the checklist.Checklist coverage
backend/cpp/rocmfp4/wrapper following the bonsai pattern (weight types decoded inside libllama; shared gRPC server reused verbatim, no allow-list patch needed)Dockerfile.rocmfp4+.docker/rocmfp4-compile.sh(from-source and prebuilt-base stages). The Dockerfile carries a documented local-build escape hatch (LLAMA_REPO/ROCMFP4_VERSIONARGs, empty defaults) so the fork can be built from a local mirror; empty values are unset in the compile script so the Makefile pin always wins in CI.github/backend-matrix.ymlentry (hipblas/amd64, prebuilt gRPC base, tag-suffix matchesindex.yamlURIs exactly)scripts/lib/backend-filter.mjs: path filter above the genericllama-cppsuffix + shared-input rule (make test-ci-scripts: 41 pass, 0 fail)bump_deps.yamlregistration; pin kept inROCMFP4_VERSION?=form the bot greps forbackend/index.yamlmeta + latest/development image entries (YAML validated)test-extra-backend-rocmfp4loading a real ROCmFP4 GGUF (fork-only decodable, bonsai pattern; artifact is 13.75 GiB — workstation test, noted in the target comment)docs/content/features/backends.mdHappy to adjust scope, naming or the escape hatch if the maintainers prefer a different shape. We run this backend in production on Strix Halo and intend to maintain it (pin bumps via the registered bump_deps entry, rebases of the format fork against upstream llama.cpp).
🤖 Generated with Claude Code
https://claude.ai/code/session_01SeuYBDnA3qH3ngNagKDNjA