Skip to content

feat(rocmfp4): backend for the ROCmFP4/ROCmFPx 4-bit formats on AMD RDNA3.5 APUs (Strix Halo) - #11636

Open
walcz-de wants to merge 6 commits into
mudler:masterfrom
walcz-de:feat/rocmfp4-backend
Open

feat(rocmfp4): backend for the ROCmFP4/ROCmFPx 4-bit formats on AMD RDNA3.5 APUs (Strix Halo)#11636
walcz-de wants to merge 6 commits into
mudler:masterfrom
walcz-de:feat/rocmfp4-backend

Conversation

@walcz-de

Copy link
Copy Markdown
Contributor

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 on llama-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_layout architecture — 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:

result
prefill pp128 317–335 t/s
decode tg32 14.0–14.3 t/s
decode with MTP draft head (spec_type:draft-mtp, n_max 4) 30.4 t/s (2.13×)
MTP losslessness identical routing accuracy (43/60) and coding pass rate (7/8) with and without MTP

llama-cpp stays the recommendation for everything else; the backend description in index.yaml says so explicitly (same tone as the vllm-cpp entry).

Non-displacement, by design

  • Deliberately NOT in 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.
  • Importer is preference-only (extends the existing llama-cpp drop-in list, like ik-llama-cpp/turboquant): ROCmFP4 GGUFs carry ordinary .gguf names, so there is no safe auto-detect signal. Without an explicit preference an import stays llama-cpp.
  • Verified at runtime: with rocmfp4 installed, a plain Q4_K_M GGUF without backend: still auto-loads on llama-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 includeDarwin entry and a single hipblas matrix 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_VERSION ARGs, 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.yml entry (hipblas/amd64, prebuilt gRPC base, tag-suffix matches index.yaml URIs exactly)
  • scripts/lib/backend-filter.mjs: path filter above the generic llama-cpp suffix + shared-input rule (make test-ci-scripts: 41 pass, 0 fail)
  • bump_deps.yaml registration; pin kept in ROCMFP4_VERSION?= form the bot greps for
  • backend/index.yaml meta + latest/development image entries (YAML validated)
  • Makefile wiring incl. test-extra-backend-rocmfp4 loading a real ROCmFP4 GGUF (fork-only decodable, bonsai pattern; artifact is 13.75 GiB — workstation test, noted in the target comment)
  • Importer preference swap + table-driven test (passes in isolation; the suite's live-HF specs for unrelated importers flake identically with and without this change)
  • Docs: docs/content/features/backends.md

Happy 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

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 localai-org-maint-bot left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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>
@walcz-de
walcz-de force-pushed the feat/rocmfp4-backend branch from 4696c51 to 3a0c54c Compare August 21, 2026 06:51
@walcz-de

Copy link
Copy Markdown
Contributor Author

Thanks for the careful review — all three addressed:

  • bump_deps.yaml: the new item is re-indented under matrix.include alongside its neighbours; the ds4 item below it is back at its normal level. Workflow YAML validates again.
  • backend/cpp/rocmfp4/patches/README.md: rewritten for rocmfp4 and the walcz-de/llama.cpp-ROCmFP4 fork — including why the directory is expected to stay empty (the fork is rebased onto the same LLAMA_VERSION the shared gRPC server is written against, unlike the Bonsai/Prism situation it was templated from).
  • DCO: the failure was a name mismatch between the commit author (stefanwalcz) and the sign-off (Stefan Walcz), same address on both sides. The branch's author fields are normalized so author and sign-off now match exactly on every commit; the attestation itself is unchanged.

@localai-org-maint-bot localai-org-maint-bot left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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>
@walcz-de

Copy link
Copy Markdown
Contributor Author

Done — apply-patches.sh header and success message rewritten for rocmfp4 / walcz-de/llama.cpp-ROCmFP4, including why patches/ is expected to stay empty here. I also swept the backend directory for further copied identity: the two remaining bonsai mentions in the Makefile are deliberate comparisons to the pattern this wrapper follows ("like bonsai, the additions are weight types decoded inside libllama"), not leftovers — happy to reword those too if you'd prefer.

@localai-org-maint-bot localai-org-maint-bot left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants