From 5fd8b732489ac5716d99f8f8815710e4366b7bfc Mon Sep 17 00:00:00 2001 From: Stefan Walcz Date: Thu, 20 Aug 2026 13:40:24 +0200 Subject: [PATCH 1/8] feat(rocmfp4): add the ROCmFP4 backend (AMD RDNA3.5) 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 --- .docker/rocmfp4-compile.sh | 39 +++++++ Makefile | 8 +- backend/Dockerfile.rocmfp4 | 160 ++++++++++++++++++++++++++ backend/cpp/rocmfp4/Makefile | 96 ++++++++++++++++ backend/cpp/rocmfp4/apply-patches.sh | 48 ++++++++ backend/cpp/rocmfp4/package.sh | 39 +++++++ backend/cpp/rocmfp4/patches/README.md | 19 +++ backend/cpp/rocmfp4/run.sh | 77 +++++++++++++ 8 files changed, 484 insertions(+), 2 deletions(-) create mode 100755 .docker/rocmfp4-compile.sh create mode 100644 backend/Dockerfile.rocmfp4 create mode 100644 backend/cpp/rocmfp4/Makefile create mode 100755 backend/cpp/rocmfp4/apply-patches.sh create mode 100755 backend/cpp/rocmfp4/package.sh create mode 100644 backend/cpp/rocmfp4/patches/README.md create mode 100755 backend/cpp/rocmfp4/run.sh diff --git a/.docker/rocmfp4-compile.sh b/.docker/rocmfp4-compile.sh new file mode 100755 index 000000000000..31a3de8ebea0 --- /dev/null +++ b/.docker/rocmfp4-compile.sh @@ -0,0 +1,39 @@ +#!/usr/bin/env bash +# Shared compile logic for backend/Dockerfile.rocmfp4. +# Sourced (via bind mount) from both builder-fromsource and builder-prebuilt stages. + +set -euxo pipefail + +export CCACHE_DIR=/root/.ccache +ccache --max-size=5G || true +ccache -z || true + +export CMAKE_ARGS="${CMAKE_ARGS:-} -DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache -DCMAKE_CUDA_COMPILER_LAUNCHER=ccache" + +if [[ -n "${CUDA_DOCKER_ARCH:-}" ]]; then + CUDA_ARCH_ESC="${CUDA_DOCKER_ARCH//;/\\;}" + export CMAKE_ARGS="${CMAKE_ARGS} -DCMAKE_CUDA_ARCHITECTURES=${CUDA_ARCH_ESC}" + echo "CMAKE_ARGS(env) = ${CMAKE_ARGS}" + rm -rf /LocalAI/backend/cpp/rocmfp4-*-build +fi + +cd /LocalAI/backend/cpp/rocmfp4 + +if [ -z "${BUILD_TYPE:-}" ]; then + # Pure CPU image: one ggml CPU_ALL_VARIANTS build replaces the per-microarch binaries. + # arm64: the armv9.2 SME variants need gcc-14 (gcc-13 rejects +sme). + if [ "${TARGETARCH}" = "arm64" ]; then + apt-get update -qq && apt-get install -y -qq gcc-14 g++-14 + export CC=gcc-14 CXX=g++-14 + fi + make rocmfp4-cpu-all +else + # GPU build (cublas/hipblas/sycl/vulkan/...): single fallback CPU build, the accelerator + # does the compute. Keeps the GPU compile from also building the CPU variant matrix and + # avoids the gcc-14 apt step on GPU base images such as nvidia l4t. + make rocmfp4-fallback +fi +make rocmfp4-grpc +make rocmfp4-rpc-server + +ccache -s || true diff --git a/Makefile b/Makefile index c68cd55223e0..c9068021d8a9 100644 --- a/Makefile +++ b/Makefile @@ -1,5 +1,5 @@ # Disable parallel execution for backend builds -.NOTPARALLEL: backends/diffusers backends/llama-cpp backends/turboquant backends/bonsai backends/outetts backends/piper backends/stablediffusion-ggml backends/trellis2cpp backends/trellis2cpp-darwin backends/whisper backends/crispasr backends/parakeet-cpp backends/moss-transcribe-cpp backends/nemo-speech-cpp backends/faster-whisper backends/silero-vad backends/local-store backends/valkey-store backends/cloud-proxy backends/huggingface backends/rfdetr backends/rfdetr-cpp backends/insightface backends/speaker-recognition backends/kitten-tts backends/kokoro backends/chatterbox backends/llama-cpp-darwin backends/neutts build-darwin-python-backend build-darwin-go-backend backends/mlx backends/diffuser-darwin backends/mlx-vlm backends/mlx-audio backends/mlx-distributed backends/stablediffusion-ggml-darwin backends/vllm backends/vllm-omni backends/longcat-video backends/sglang backends/moonshine backends/pocket-tts backends/qwen-tts backends/faster-qwen3-tts backends/qwen-asr backends/nemo backends/voxcpm backends/whisperx backends/ace-step backends/acestep-cpp backends/fish-speech backends/voxtral backends/opus backends/trl backends/llama-cpp-quantization backends/kokoros backends/sam3-cpp backends/qwen3-tts-cpp backends/moss-tts-cpp backends/magpie-tts-cpp backends/vllm-cpp backends/omnivoice-cpp backends/vibevoice-cpp backends/localvqe backends/tinygrad backends/sherpa-onnx backends/ds4 backends/ds4-darwin backends/liquid-audio backends/supertonic backends/depth-anything-cpp backends/privacy-filter backends/privacy-filter-darwin backends/audio-cpp backends/audio-cpp-darwin +.NOTPARALLEL: backends/rocmfp4 backends/diffusers backends/llama-cpp backends/turboquant backends/bonsai backends/outetts backends/piper backends/stablediffusion-ggml backends/trellis2cpp backends/trellis2cpp-darwin backends/whisper backends/crispasr backends/parakeet-cpp backends/moss-transcribe-cpp backends/nemo-speech-cpp backends/faster-whisper backends/silero-vad backends/local-store backends/valkey-store backends/cloud-proxy backends/huggingface backends/rfdetr backends/rfdetr-cpp backends/insightface backends/speaker-recognition backends/kitten-tts backends/kokoro backends/chatterbox backends/llama-cpp-darwin backends/neutts build-darwin-python-backend build-darwin-go-backend backends/mlx backends/diffuser-darwin backends/mlx-vlm backends/mlx-audio backends/mlx-distributed backends/stablediffusion-ggml-darwin backends/vllm backends/vllm-omni backends/longcat-video backends/sglang backends/moonshine backends/pocket-tts backends/qwen-tts backends/faster-qwen3-tts backends/qwen-asr backends/nemo backends/voxcpm backends/whisperx backends/ace-step backends/acestep-cpp backends/fish-speech backends/voxtral backends/opus backends/trl backends/llama-cpp-quantization backends/kokoros backends/sam3-cpp backends/qwen3-tts-cpp backends/moss-tts-cpp backends/magpie-tts-cpp backends/vllm-cpp backends/omnivoice-cpp backends/vibevoice-cpp backends/localvqe backends/tinygrad backends/sherpa-onnx backends/ds4 backends/ds4-darwin backends/liquid-audio backends/supertonic backends/depth-anything-cpp backends/privacy-filter backends/privacy-filter-darwin backends/audio-cpp backends/audio-cpp-darwin GOCMD=go GOTEST=$(GOCMD) test @@ -1284,6 +1284,9 @@ BACKEND_TURBOQUANT = turboquant|turboquant|.|false|false # weight-quant kernels the Bonsai / Ternary-Bonsai models ship in. Reuses # backend/cpp/llama-cpp grpc-server sources via a thin wrapper Makefile. BACKEND_BONSAI = bonsai|bonsai|.|false|false +# rocmfp4 is a llama.cpp fork carrying the ROCmFP4 / ROCmFPx weight formats +# (ggml types 100-107) for AMD RDNA3.5 APUs. Stock llama.cpp rejects those types. +BACKEND_ROCMFP4 = rocmfp4|rocmfp4|.|false|false # ds4 is antirez/ds4, a DeepSeek V4 Flash-specific inference engine. # Single-model; hardware-only validation lives at tests/e2e-backends/ # (BACKEND_BINARY mode); see docs/superpowers/plans/2026-05-11-ds4-backend.md. @@ -1399,6 +1402,7 @@ $(eval $(call generate-docker-build-target,$(BACKEND_LLAMA_CPP))) $(eval $(call generate-docker-build-target,$(BACKEND_IK_LLAMA_CPP))) $(eval $(call generate-docker-build-target,$(BACKEND_TURBOQUANT))) $(eval $(call generate-docker-build-target,$(BACKEND_BONSAI))) +$(eval $(call generate-docker-build-target,$(BACKEND_ROCMFP4))) $(eval $(call generate-docker-build-target,$(BACKEND_DS4))) $(eval $(call generate-docker-build-target,$(BACKEND_PRIVACY_FILTER))) $(eval $(call generate-docker-build-target,$(BACKEND_AUDIO_CPP))) @@ -1471,7 +1475,7 @@ $(eval $(call generate-docker-build-target,$(BACKEND_SUPERTONIC))) docker-save-%: backend-images docker save local-ai-backend:$* -o backend-images/$*.tar -docker-build-backends: docker-build-llama-cpp docker-build-ik-llama-cpp docker-build-turboquant docker-build-bonsai docker-build-ds4 docker-build-rerankers docker-build-vllm docker-build-vllm-omni docker-build-longcat-video docker-build-sglang docker-build-transformers docker-build-outetts docker-build-diffusers docker-build-kokoro docker-build-faster-whisper docker-build-crispasr docker-build-coqui docker-build-chatterbox docker-build-vibevoice docker-build-liquid-audio docker-build-moonshine docker-build-pocket-tts docker-build-qwen-tts docker-build-fish-speech docker-build-faster-qwen3-tts docker-build-qwen-asr docker-build-nemo docker-build-voxcpm docker-build-whisperx docker-build-ace-step docker-build-acestep-cpp docker-build-voxtral docker-build-mlx-distributed docker-build-trl docker-build-llama-cpp-quantization docker-build-tinygrad docker-build-kokoros docker-build-sam3-cpp docker-build-rfdetr-cpp docker-build-qwen3-tts-cpp docker-build-moss-tts-cpp docker-build-magpie-tts-cpp docker-build-vllm-cpp docker-build-omnivoice-cpp docker-build-vibevoice-cpp docker-build-localvqe docker-build-insightface docker-build-speaker-recognition docker-build-sherpa-onnx docker-build-cloud-proxy docker-build-supertonic docker-build-depth-anything-cpp docker-build-moss-transcribe-cpp docker-build-nemo-speech-cpp docker-build-privacy-filter docker-build-trellis2cpp docker-build-valkey-store docker-build-audio-cpp +docker-build-backends: docker-build-rocmfp4 docker-build-llama-cpp docker-build-ik-llama-cpp docker-build-turboquant docker-build-bonsai docker-build-ds4 docker-build-rerankers docker-build-vllm docker-build-vllm-omni docker-build-longcat-video docker-build-sglang docker-build-transformers docker-build-outetts docker-build-diffusers docker-build-kokoro docker-build-faster-whisper docker-build-crispasr docker-build-coqui docker-build-chatterbox docker-build-vibevoice docker-build-liquid-audio docker-build-moonshine docker-build-pocket-tts docker-build-qwen-tts docker-build-fish-speech docker-build-faster-qwen3-tts docker-build-qwen-asr docker-build-nemo docker-build-voxcpm docker-build-whisperx docker-build-ace-step docker-build-acestep-cpp docker-build-voxtral docker-build-mlx-distributed docker-build-trl docker-build-llama-cpp-quantization docker-build-tinygrad docker-build-kokoros docker-build-sam3-cpp docker-build-rfdetr-cpp docker-build-qwen3-tts-cpp docker-build-moss-tts-cpp docker-build-magpie-tts-cpp docker-build-vllm-cpp docker-build-omnivoice-cpp docker-build-vibevoice-cpp docker-build-localvqe docker-build-insightface docker-build-speaker-recognition docker-build-sherpa-onnx docker-build-cloud-proxy docker-build-supertonic docker-build-depth-anything-cpp docker-build-moss-transcribe-cpp docker-build-nemo-speech-cpp docker-build-privacy-filter docker-build-trellis2cpp docker-build-valkey-store docker-build-audio-cpp ######################################################## ### Mock Backend for E2E Tests diff --git a/backend/Dockerfile.rocmfp4 b/backend/Dockerfile.rocmfp4 new file mode 100644 index 000000000000..2819d7bf5602 --- /dev/null +++ b/backend/Dockerfile.rocmfp4 @@ -0,0 +1,160 @@ +ARG BASE_IMAGE=ubuntu:24.04 +# BUILDER_BASE_IMAGE defaults to BASE_IMAGE so the Dockerfile parses even +# when no prebuilt base is supplied. The builder-prebuilt stage is only +# entered when BUILDER_TARGET=builder-prebuilt, so a "wrong" fallback +# content here is harmless — BuildKit prunes the unreferenced builder. +ARG BUILDER_BASE_IMAGE=${BASE_IMAGE} +# BUILDER_TARGET selects which builder stage the final scratch image copies +# package output from. Declared at global scope (before any FROM) so it's +# usable in `FROM ${BUILDER_TARGET}` below. Default keeps local +# `make backends/rocmfp4` on the from-source path. +ARG BUILDER_TARGET=builder-fromsource +ARG APT_MIRROR="" +ARG APT_PORTS_MIRROR="" + + +# ============================================================================ +# Stage: builder-fromsource — self-contained build path. +# Runs .docker/install-base-deps.sh (apt deps + cmake + protoc + gRPC + +# conditional CUDA/ROCm/Vulkan), copies /opt/grpc to /usr/local, then +# compiles the variant. Used when BUILDER_TARGET=builder-fromsource (the +# default; local `make backends/rocmfp4`). +# +# The install script is the same one that backend/Dockerfile.base-grpc-builder +# runs, so the result is bit-equivalent to the prebuilt-base path +# (builder-prebuilt below). +# ============================================================================ +FROM ${BASE_IMAGE} AS builder-fromsource +ARG BUILD_TYPE +ARG CUDA_MAJOR_VERSION +ARG CUDA_MINOR_VERSION +ARG CMAKE_FROM_SOURCE=false +# CUDA Toolkit 13.x compatibility: CMake 3.31.9+ fixes toolchain detection/arch table issues +ARG CMAKE_VERSION=3.31.10 +ARG GRPC_VERSION=v1.65.0 +ARG GRPC_MAKEFLAGS="-j4 -Otarget" +ARG SKIP_DRIVERS=false +ARG TARGETARCH +ARG TARGETVARIANT +ARG GO_VERSION=1.25.4 +ARG UBUNTU_VERSION=2404 +ARG APT_MIRROR +ARG APT_PORTS_MIRROR +ARG AMDGPU_TARGETS="" +ARG BACKEND=rerankers +# CUDA target archs, e.g. --build-arg CUDA_DOCKER_ARCH='75;86;89;120' +ARG CUDA_DOCKER_ARCH +ARG CMAKE_ARGS + +ENV BUILD_TYPE=${BUILD_TYPE} \ + CUDA_MAJOR_VERSION=${CUDA_MAJOR_VERSION} \ + CUDA_MINOR_VERSION=${CUDA_MINOR_VERSION} \ + CMAKE_FROM_SOURCE=${CMAKE_FROM_SOURCE} \ + CMAKE_VERSION=${CMAKE_VERSION} \ + GRPC_VERSION=${GRPC_VERSION} \ + GRPC_MAKEFLAGS=${GRPC_MAKEFLAGS} \ + SKIP_DRIVERS=${SKIP_DRIVERS} \ + TARGETARCH=${TARGETARCH} \ + UBUNTU_VERSION=${UBUNTU_VERSION} \ + APT_MIRROR=${APT_MIRROR} \ + APT_PORTS_MIRROR=${APT_PORTS_MIRROR} \ + AMDGPU_TARGETS=${AMDGPU_TARGETS} \ + CUDA_DOCKER_ARCH=${CUDA_DOCKER_ARCH} \ + CMAKE_ARGS=${CMAKE_ARGS} \ + DEBIAN_FRONTEND=noninteractive + +# CUDA on PATH (no-op when CUDA isn't installed) +ENV PATH=/usr/local/cuda/bin:${PATH} +# HipBLAS / ROCm on PATH (no-op when ROCm isn't installed) +ENV PATH=/opt/rocm/bin:${PATH} + +WORKDIR /build + +# Install everything via the shared script — the same one that +# backend/Dockerfile.base-grpc-builder runs, so the prebuilt CI base and +# this from-source path are bit-equivalent. +RUN --mount=type=bind,source=.docker/install-base-deps.sh,target=/usr/local/sbin/install-base-deps \ + --mount=type=bind,source=.docker/apt-mirror.sh,target=/usr/local/sbin/apt-mirror \ + bash /usr/local/sbin/install-base-deps + +# Mirror builder-prebuilt: copy gRPC from /opt/grpc to /usr/local so +# CMake's find_package finds it at the canonical prefix the Makefile expects. +RUN cp -a /opt/grpc/. /usr/local/ + +COPY . /LocalAI + +# BuildKit cache mount for ccache. See Dockerfile.llama-cpp (commit 9228e5b4) +# for rationale. rocmfp4 is a llama.cpp fork that reuses +# backend/cpp/llama-cpp source via a thin wrapper Makefile, so MOST TUs +# are content-identical to the upstream llama-cpp build. Sharing a cache +# id with llama-cpp could give cross-fork hits — but for now keep them +# separate so a regression in one doesn't poison the other. Revisit +# sharing after measuring the actual hit rate. +# +# The compile body is shared with builder-prebuilt via .docker/rocmfp4-compile.sh. +RUN --mount=type=bind,source=.docker/rocmfp4-compile.sh,target=/usr/local/sbin/compile.sh \ + --mount=type=cache,target=/root/.ccache,id=rocmfp4-ccache-${TARGETARCH}-${BUILD_TYPE},sharing=locked \ + bash /usr/local/sbin/compile.sh + + +# Copy libraries using a script to handle architecture differences +RUN make -BC /LocalAI/backend/cpp/rocmfp4 package + + +# ============================================================================ +# Stage: builder-prebuilt — uses the pre-built base from +# quay.io/go-skynet/ci-cache:base-grpc-* (built by .github/workflows/base-images.yml). +# That image already has gRPC at /opt/grpc + apt deps + CUDA/ROCm/Vulkan +# pre-installed, so we just copy gRPC to /usr/local and compile. Used when +# BUILDER_TARGET=builder-prebuilt (CI when the matrix entry sets +# builder-base-image). +# ============================================================================ +FROM ${BUILDER_BASE_IMAGE} AS builder-prebuilt + +ARG BUILD_TYPE +ENV BUILD_TYPE=${BUILD_TYPE} +ARG CUDA_DOCKER_ARCH +ENV CUDA_DOCKER_ARCH=${CUDA_DOCKER_ARCH} +ARG CMAKE_ARGS +ENV CMAKE_ARGS=${CMAKE_ARGS} +# AMDGPU_TARGETS must be forwarded into the env here too — backend/cpp/llama-cpp/Makefile +# (which the rocmfp4 Makefile reuses via a sibling build dir) errors out when the var +# is empty on a hipblas build, and the prebuilt path is what CI exercises most of the +# time. The builder-fromsource stage above already does this; mirror it here. +ARG AMDGPU_TARGETS +ENV AMDGPU_TARGETS=${AMDGPU_TARGETS} +ARG TARGETARCH +ARG TARGETVARIANT + +# The base-grpc-* image installs gRPC to /opt/grpc but doesn't copy it to +# /usr/local. Mirror what the from-source path does so the compile step +# can find gRPC at the canonical prefix the Makefile expects. +RUN cp -a /opt/grpc/. /usr/local/ + +COPY . /LocalAI + +RUN --mount=type=bind,source=.docker/rocmfp4-compile.sh,target=/usr/local/sbin/compile.sh \ + --mount=type=cache,target=/root/.ccache,id=rocmfp4-ccache-${TARGETARCH}-${BUILD_TYPE},sharing=locked \ + bash /usr/local/sbin/compile.sh + +RUN make -BC /LocalAI/backend/cpp/rocmfp4 package + + +# ============================================================================ +# Final stage — copies package output from one of the two builders. +# BUILDER_TARGET selects which one. BuildKit prunes the unreferenced builder. +# +# BuildKit doesn't support variable expansion in `COPY --from=` directly, +# so we resolve the ARG by aliasing the chosen builder to a fixed stage +# name via `FROM ${BUILDER_TARGET} AS builder` and then COPY --from=builder. +# BUILDER_TARGET itself is declared as a global ARG at the top of this +# file (required for use in FROM), so we just re-import it into this +# stage's scope before the FROM directive. +# ============================================================================ +FROM ${BUILDER_TARGET} AS builder + +FROM scratch + + +# Copy all available binaries (the build process only creates the appropriate ones for the target architecture) +COPY --from=builder /LocalAI/backend/cpp/rocmfp4/package/. ./ diff --git a/backend/cpp/rocmfp4/Makefile b/backend/cpp/rocmfp4/Makefile new file mode 100644 index 000000000000..987cccf8a1ce --- /dev/null +++ b/backend/cpp/rocmfp4/Makefile @@ -0,0 +1,96 @@ +# Pinned to the HEAD of the `rocmfp4` branch on https://github.com/walcz-de/llama.cpp-ROCmFP4. +# Auto-bumped nightly by .github/workflows/bump_deps.yaml. +ROCMFP4_VERSION?=0b8c6fbda8c18f7ae4bd30811996c24f04ddc0dc +LLAMA_REPO?=https://github.com/walcz-de/llama.cpp-ROCmFP4 + +CMAKE_ARGS?= +BUILD_TYPE?= +NATIVE?=false +ONEAPI_VARS?=/opt/intel/oneapi/setvars.sh +TARGET?=--target grpc-server +JOBS?=$(shell nproc 2>/dev/null || sysctl -n hw.ncpu 2>/dev/null || echo 1) +ARCH?=$(shell uname -m) + +CURRENT_MAKEFILE_DIR := $(dir $(abspath $(lastword $(MAKEFILE_LIST)))) +LLAMA_CPP_DIR := $(CURRENT_MAKEFILE_DIR)/../llama-cpp + +GREEN := \033[0;32m +RESET := \033[0m + +# rocmfp4 is a llama.cpp fork carrying the ROCmFP4 / ROCmFPx weight-quantization +# formats (ggml types 100-107) for AMD RDNA3.5 APUs. Stock llama.cpp rejects those +# tensor types, which is the whole reason this backend exists as a separate image +# rather than a flag on llama-cpp. +# +# Like bonsai, the additions are model *weight* types decoded inside libllama and +# therefore transparent to the shared gRPC server: no grpc-server.cpp allow-list +# patch is needed (unlike turboquant, whose KV-cache types do need one). So we reuse +# backend/cpp/llama-cpp's grpc-server.cpp / CMakeLists.txt / Makefile verbatim and +# only swap which repo and commit the fetch step pulls. +PATCHES_DIR := $(CURRENT_MAKEFILE_DIR)/patches + +define rocmfp4-build + rm -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build + cp -rf $(LLAMA_CPP_DIR) $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build + # Drop patches vendored for upstream llama.cpp: this fork carries them already + # (it is based on the same pin), so re-applying them would reject. + rm -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build/patches + $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build purge + bash $(LLAMA_CPP_DIR)/disable-score-task.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build/grpc-server.cpp + bash $(LLAMA_CPP_DIR)/disable-tts-task.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build/grpc-server.cpp + $(info $(GREEN)I rocmfp4 build info:$(1)$(RESET)) + LLAMA_REPO=$(LLAMA_REPO) LLAMA_VERSION=$(ROCMFP4_VERSION) \ + $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build llama.cpp + bash $(CURRENT_MAKEFILE_DIR)/apply-patches.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build/llama.cpp $(PATCHES_DIR) + CMAKE_ARGS="$(CMAKE_ARGS) $(2)" TARGET="$(3)" \ + LLAMA_REPO=$(LLAMA_REPO) LLAMA_VERSION=$(ROCMFP4_VERSION) \ + $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build grpc-server + cp -rfv $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build/grpc-server rocmfp4-$(1) +endef + +rocmfp4-avx2: + $(call rocmfp4-build,avx2,-DGGML_AVX=on -DGGML_AVX2=on -DGGML_AVX512=off -DGGML_FMA=on -DGGML_F16C=on,--target grpc-server) + +rocmfp4-avx512: + $(call rocmfp4-build,avx512,-DGGML_AVX=on -DGGML_AVX2=off -DGGML_AVX512=on -DGGML_FMA=on -DGGML_F16C=on,--target grpc-server) + +rocmfp4-avx: + $(call rocmfp4-build,avx,-DGGML_AVX=on -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server) + +rocmfp4-fallback: + $(call rocmfp4-build,fallback,-DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server) + +# Single-build CPU backend via ggml CPU_ALL_VARIANTS (mirrors llama-cpp-cpu-all). +rocmfp4-cpu-all: + rm -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build + cp -rf $(LLAMA_CPP_DIR) $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build + rm -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/patches + $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build purge + bash $(LLAMA_CPP_DIR)/disable-score-task.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/grpc-server.cpp + bash $(LLAMA_CPP_DIR)/disable-tts-task.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/grpc-server.cpp + $(info $(GREEN)I rocmfp4 build info:cpu-all-variants$(RESET)) + LLAMA_REPO=$(LLAMA_REPO) LLAMA_VERSION=$(ROCMFP4_VERSION) \ + $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build llama.cpp + bash $(CURRENT_MAKEFILE_DIR)/apply-patches.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/llama.cpp $(PATCHES_DIR) + SHARED_LIBS=ON EXTRA_CMAKE_ARGS="-DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON" TARGET="--target grpc-server --target ggml" \ + LLAMA_REPO=$(LLAMA_REPO) LLAMA_VERSION=$(ROCMFP4_VERSION) \ + $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build grpc-server + cp -rfv $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/grpc-server rocmfp4-cpu-all + rm -rf ggml-shared-libs && mkdir -p ggml-shared-libs + find $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/llama.cpp/build \( -name '*.so*' -o -name '*.dylib' \) -exec cp -av {} ggml-shared-libs/ \; + @echo "Collected ggml shared backends:" && ls -la ggml-shared-libs/ + +rocmfp4-grpc: + $(call rocmfp4-build,grpc,-DGGML_RPC=ON -DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server --target rpc-server) + +rocmfp4-rpc-server: rocmfp4-grpc + cp -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-grpc-build/llama.cpp/build/bin/rpc-server rocmfp4-rpc-server + +package: + bash package.sh + +purge: + rm -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-*-build + rm -rf rocmfp4-* package + +clean: purge diff --git a/backend/cpp/rocmfp4/apply-patches.sh b/backend/cpp/rocmfp4/apply-patches.sh new file mode 100755 index 000000000000..7ac850581262 --- /dev/null +++ b/backend/cpp/rocmfp4/apply-patches.sh @@ -0,0 +1,48 @@ +#!/bin/bash +# Apply the bonsai patch series to a cloned PrismML llama.cpp (prism branch) checkout. +# +# The prism fork branched from upstream llama.cpp before a number of API changes that the +# shared backend/cpp/llama-cpp/grpc-server.cpp depends on. We carry those upstream commits +# as patch files under backend/cpp/bonsai/patches/ and apply them here so the reused +# grpc-server source compiles against the fork unmodified. +# +# Drop the corresponding patch from patches/ whenever the fork catches up with upstream — +# the build will fail fast if a patch stops applying, which is the signal to retire it. + +set -euo pipefail + +if [[ $# -ne 2 ]]; then + echo "usage: $0 " >&2 + exit 2 +fi + +SRC_DIR=$1 +PATCHES_DIR=$2 + +if [[ ! -d "$SRC_DIR" ]]; then + echo "source dir does not exist: $SRC_DIR" >&2 + exit 2 +fi + +if [[ ! -d "$PATCHES_DIR" ]]; then + echo "no patches dir at $PATCHES_DIR, nothing to apply" + exit 0 +fi + +shopt -s nullglob +patches=("$PATCHES_DIR"/*.patch) +shopt -u nullglob + +if [[ ${#patches[@]} -eq 0 ]]; then + echo "no .patch files in $PATCHES_DIR, nothing to apply" + exit 0 +fi + +cd "$SRC_DIR" + +for patch in "${patches[@]}"; do + echo "==> applying $patch" + git apply --verbose "$patch" +done + +echo "all bonsai patches applied successfully" diff --git a/backend/cpp/rocmfp4/package.sh b/backend/cpp/rocmfp4/package.sh new file mode 100755 index 000000000000..691c1abd801a --- /dev/null +++ b/backend/cpp/rocmfp4/package.sh @@ -0,0 +1,39 @@ +#!/bin/bash + +# Script to copy the appropriate libraries based on architecture +# This script is used in the final stage of the Dockerfile + +set -e + +CURDIR=$(dirname "$(realpath $0)") +REPO_ROOT="${CURDIR}/../../.." + +# Create lib directory +mkdir -p $CURDIR/package/lib + +cp -avrf $CURDIR/rocmfp4-* $CURDIR/package/ +cp -rfv $CURDIR/run.sh $CURDIR/package/ + +# Bundle the ggml shared backends from the CPU_ALL_VARIANTS build into package/lib. ggml +# discovers the per-microarch libggml-cpu-*.so by scanning the executable directory, which +# (via the bundled lib/ld.so that run.sh launches through) resolves to lib/. See the +# matching comment in backend/cpp/llama-cpp/package.sh. No-op on the fallback/ROCm builds. +if [ -d "$CURDIR/ggml-shared-libs" ]; then + echo "Bundling ggml shared backends (CPU_ALL_VARIANTS)..." + cp -avf $CURDIR/ggml-shared-libs/*.so* $CURDIR/package/lib/ +fi + +# Detect architecture and copy appropriate libraries +source "$CURDIR/../../../scripts/build/package-system-libs.sh" "$CURDIR/package/lib" "" + +# Package GPU libraries based on BUILD_TYPE +GPU_LIB_SCRIPT="${REPO_ROOT}/scripts/build/package-gpu-libs.sh" +if [ -f "$GPU_LIB_SCRIPT" ]; then + echo "Packaging GPU libraries for BUILD_TYPE=${BUILD_TYPE:-cpu}..." + source "$GPU_LIB_SCRIPT" "$CURDIR/package/lib" + package_gpu_libs +fi + +echo "Packaging completed successfully" +ls -liah $CURDIR/package/ +ls -liah $CURDIR/package/lib/ diff --git a/backend/cpp/rocmfp4/patches/README.md b/backend/cpp/rocmfp4/patches/README.md new file mode 100644 index 000000000000..3bfb74183f98 --- /dev/null +++ b/backend/cpp/rocmfp4/patches/README.md @@ -0,0 +1,19 @@ +# bonsai fork skew patches + +The `bonsai` backend reuses `backend/cpp/llama-cpp/grpc-server.cpp` (written against +LocalAI's pinned *upstream* llama.cpp) but compiles it against the PrismML `prism` fork, +which branched from upstream some commits earlier. Any upstream API change that the shared +gRPC server depends on, but that the fork does not yet carry, is back-ported here as a +`*.patch` file and applied to the cloned fork checkout by `../apply-patches.sh`. + +CI treats both this directory and `backend/cpp/llama-cpp/` as Bonsai inputs, since +the wrapper copies and builds the shared llama.cpp backend sources. + +Rules: + +- One upstream commit (or minimal hunk) per patch, named `NNNN-short-description.patch`. +- Patches are applied with `git apply` from the fork's checkout root. +- `apply-patches.sh` fails fast if a patch stops applying cleanly — that is the signal the + fork has caught up (or diverged), so re-cut or drop the patch. +- Keep this set as small as possible; the long-term fix is the fork rebasing onto a newer + upstream (or Q1_0/Q2_0 landing in mainline llama.cpp, retiring this backend entirely). diff --git a/backend/cpp/rocmfp4/run.sh b/backend/cpp/rocmfp4/run.sh new file mode 100755 index 000000000000..0ccf3cae7137 --- /dev/null +++ b/backend/cpp/rocmfp4/run.sh @@ -0,0 +1,77 @@ +#!/bin/bash +set -ex + +# Get the absolute current dir where the script is located +CURDIR=$(dirname "$(realpath "$0")") + +cd / + +echo "CPU info:" +grep -e "model\sname" /proc/cpuinfo | head -1 +grep -e "flags" /proc/cpuinfo | head -1 + +BINARY=rocmfp4-fallback + +# x86/arm64 ship a single rocmfp4-cpu-all built with ggml CPU_ALL_VARIANTS: ggml's +# backend registry dlopens the best libggml-cpu-*.so for this host, so no shell-side +# probing. ROCm ships only rocmfp4-fallback, so fall back to it when cpu-all is absent. +if [ -e "$CURDIR"/rocmfp4-cpu-all ]; then + BINARY=rocmfp4-cpu-all +fi + +if [ -n "$LLAMACPP_GRPC_SERVERS" ]; then + if [ -e "$CURDIR"/rocmfp4-grpc ]; then + BINARY=rocmfp4-grpc + fi +fi + +# Extend ld library path with the dir where this script is located/lib +if [ "$(uname)" == "Darwin" ]; then + export DYLD_LIBRARY_PATH="$CURDIR"/lib:$DYLD_LIBRARY_PATH +else + export LD_LIBRARY_PATH="$CURDIR"/lib:$LD_LIBRARY_PATH + # Tell rocBLAS where to find TensileLibrary data (GPU kernel tuning files) + if [ -d "$CURDIR/lib/rocblas/library" ]; then + export ROCBLAS_TENSILE_LIBPATH="$CURDIR"/lib/rocblas/library + fi + # Same for hipBLASLt (rocblaslt): the bundled libhipblaslt.so resolves its + # TensileLibrary_lazy_gfx*.dat kernel data relative to itself, so point it at + # the bundled data or it falls back to slow generic kernels (issue #10660). + if [ -d "$CURDIR/lib/hipblaslt/library" ]; then + export HIPBLASLT_TENSILE_LIBPATH="$CURDIR"/lib/hipblaslt/library + fi + # Backends built for Intel GPUs carry a copy of the Intel graphics driver, + # and libze_loader is only there in those builds. Level Zero looks for a + # driver on its own, so point it at the copy that came with this backend: it + # was built against the same C library, while the machine's own driver may + # not have been, and loading that one can crash on start. + # + # Anything the user set is left alone, so a machine with a graphics card + # newer than the driver carried here can still be told to use its own. + # Nothing is said about OpenCL: no OpenCL driver is carried, so anything we + # set there would leave OpenCL worse off than the machine's own setup. + if [ -e "$CURDIR/lib/libze_loader.so.1" ]; then + if [ -e "$CURDIR/lib/libze_intel_gpu.so.1" ] && [ -z "${ZE_ENABLE_ALT_DRIVERS:-}" ]; then + export ZE_ENABLE_ALT_DRIVERS="$CURDIR"/lib/libze_intel_gpu.so.1 + fi + # Ask the driver how much graphics memory is free. Without this, the + # backend reads zero on an integrated graphics chip, because such a chip + # shares the system memory instead of having its own. + if [ -z "${ZES_ENABLE_SYSMAN:-}" ]; then + export ZES_ENABLE_SYSMAN=1 + fi + fi +fi + +# If there is a lib/ld.so, use it +if [ -f "$CURDIR"/lib/ld.so ]; then + echo "Using lib/ld.so" + echo "Using binary: $BINARY" + exec "$CURDIR"/lib/ld.so "$CURDIR"/$BINARY "$@" +fi + +echo "Using binary: $BINARY" +exec "$CURDIR"/$BINARY "$@" + +# We should never reach this point, however just in case we do, run fallback +exec "$CURDIR"/rocmfp4-fallback "$@" From 0bad0677825736e90c056eb89d01612b9b782635 Mon Sep 17 00:00:00 2001 From: Stefan Walcz Date: Thu, 20 Aug 2026 13:59:42 +0200 Subject: [PATCH 2/8] feat(rocmfp4): register the backend across catalog, CI and importer 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 --- .docker/rocmfp4-compile.sh | 6 ++++ .github/backend-matrix.yml | 14 +++++++++ .github/workflows/bump_deps.yaml | 4 +++ backend/Dockerfile.rocmfp4 | 12 ++++++++ backend/index.yaml | 47 +++++++++++++++++++++++++++++ core/gallery/importers/llama-cpp.go | 5 ++- docs/content/features/backends.md | 1 + scripts/lib/backend-filter.mjs | 8 ++++- 8 files changed, 95 insertions(+), 2 deletions(-) diff --git a/.docker/rocmfp4-compile.sh b/.docker/rocmfp4-compile.sh index 31a3de8ebea0..7547ab5f0960 100755 --- a/.docker/rocmfp4-compile.sh +++ b/.docker/rocmfp4-compile.sh @@ -4,6 +4,12 @@ set -euxo pipefail +# Docker ARG defaults arrive as empty strings, and an empty-but-defined variable +# still beats Make's ?= assignment. Drop them so the Makefile's pin wins unless a +# local build genuinely overrides it. +[ -n "${LLAMA_REPO:-}" ] || unset LLAMA_REPO || true +[ -n "${ROCMFP4_VERSION:-}" ] || unset ROCMFP4_VERSION || true + export CCACHE_DIR=/root/.ccache ccache --max-size=5G || true ccache -z || true diff --git a/.github/backend-matrix.yml b/.github/backend-matrix.yml index e7bada4b4e91..0052634a69d6 100644 --- a/.github/backend-matrix.yml +++ b/.github/backend-matrix.yml @@ -2243,6 +2243,20 @@ include: dockerfile: "./backend/Dockerfile.llama-cpp" context: "./" ubuntu-version: '2404' + - build-type: 'hipblas' + cuda-major-version: "" + cuda-minor-version: "" + platforms: 'linux/amd64' + tag-latest: 'auto' + tag-suffix: '-gpu-rocm-hipblas-rocmfp4' + builder-base-image: 'quay.io/go-skynet/ci-cache:base-grpc-rocm-amd64' + runs-on: 'ubuntu-latest' + base-image: "rocm/dev-ubuntu-24.04:7.2.1" + skip-drivers: 'false' + backend: "rocmfp4" + dockerfile: "./backend/Dockerfile.rocmfp4" + context: "./" + ubuntu-version: '2404' - build-type: 'hipblas' cuda-major-version: "" cuda-minor-version: "" diff --git a/.github/workflows/bump_deps.yaml b/.github/workflows/bump_deps.yaml index 069b094a04d4..38147d3eb82c 100644 --- a/.github/workflows/bump_deps.yaml +++ b/.github/workflows/bump_deps.yaml @@ -26,6 +26,10 @@ jobs: variable: "BONSAI_VERSION" branch: "prism" file: "backend/cpp/bonsai/Makefile" + - repository: "walcz-de/llama.cpp-ROCmFP4" + variable: "ROCMFP4_VERSION" + branch: "rocmfp4" + file: "backend/cpp/rocmfp4/Makefile" - repository: "antirez/ds4" variable: "DS4_VERSION" branch: "main" diff --git a/backend/Dockerfile.rocmfp4 b/backend/Dockerfile.rocmfp4 index 2819d7bf5602..7beb88e97685 100644 --- a/backend/Dockerfile.rocmfp4 +++ b/backend/Dockerfile.rocmfp4 @@ -42,6 +42,18 @@ ARG APT_MIRROR ARG APT_PORTS_MIRROR ARG AMDGPU_TARGETS="" ARG BACKEND=rerankers +# Local-build escape hatch: point the fetch at a git daemon on the host instead of +# GitHub. Defaults are empty, so CI keeps using the Makefile's pinned repo+sha. +ARG LLAMA_REPO= +ARG ROCMFP4_VERSION= +ENV LLAMA_REPO=${LLAMA_REPO} +ENV ROCMFP4_VERSION=${ROCMFP4_VERSION} +# Local-build escape hatch: point the fetch at a git daemon on the host instead of +# GitHub. Defaults are empty, so CI keeps using the Makefile's pinned repo+sha. +ARG LLAMA_REPO= +ARG ROCMFP4_VERSION= +ENV LLAMA_REPO=${LLAMA_REPO} +ENV ROCMFP4_VERSION=${ROCMFP4_VERSION} # CUDA target archs, e.g. --build-arg CUDA_DOCKER_ARCH='75;86;89;120' ARG CUDA_DOCKER_ARCH ARG CMAKE_ARGS diff --git a/backend/index.yaml b/backend/index.yaml index 3420d85496b7..e17431b7c901 100644 --- a/backend/index.yaml +++ b/backend/index.yaml @@ -74,6 +74,43 @@ nvidia-cuda-12: "cuda12-turboquant" nvidia-l4t-cuda-12: "nvidia-l4t-arm64-turboquant" nvidia-l4t-cuda-13: "cuda13-nvidia-l4t-arm64-turboquant" +- &rocmfp4 + name: "rocmfp4" + alias: "rocmfp4" + license: mit + description: | + Fork of llama.cpp carrying the ROCmFP4 / ROCmFPx 4-bit 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 these tensor types, so this backend exists to serve the + published ROCmFP4 weights rather than to be a faster llama.cpp. + + What the format buys on this hardware is MEMORY, not arithmetic: roughly 22% + smaller than Q4_K_M at perplexity parity, which can decide whether a model + fits in unified RAM. 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 gain is + bandwidth. Measured on a Radeon 8060S (gfx1151): 335 tok/s prefill and + 14.2 tok/s decode on a 27B model. + + llama-cpp stays the recommendation for everything else. + urls: + - https://github.com/walcz-de/llama.cpp-ROCmFP4 + - https://github.com/charlie12345/ROCmFPX + tags: + - text-to-text + - LLM + - GPU + - HIP + - AMD + - rocmfp4 + - fp4 + - strix-halo + capabilities: + amd: "rocm-rocmfp4" +- !!merge <<: *rocmfp4 + name: "rocmfp4-development" + capabilities: + amd: "rocm-rocmfp4-development" - &bonsai name: "bonsai" alias: "bonsai" @@ -2859,6 +2896,16 @@ uri: "quay.io/go-skynet/local-ai-backends:master-gpu-nvidia-cuda-13-bonsai" mirrors: - localai/localai-backends:master-gpu-nvidia-cuda-13-bonsai +- !!merge <<: *rocmfp4 + name: "rocm-rocmfp4" + uri: "quay.io/go-skynet/local-ai-backends:latest-gpu-rocm-hipblas-rocmfp4" + mirrors: + - localai/localai-backends:latest-gpu-rocm-hipblas-rocmfp4 +- !!merge <<: *rocmfp4 + name: "rocm-rocmfp4-development" + uri: "quay.io/go-skynet/local-ai-backends:master-gpu-rocm-hipblas-rocmfp4" + mirrors: + - localai/localai-backends:master-gpu-rocm-hipblas-rocmfp4 - !!merge <<: *bonsai name: "rocm-bonsai" uri: "quay.io/go-skynet/local-ai-backends:latest-gpu-rocm-hipblas-bonsai" diff --git a/core/gallery/importers/llama-cpp.go b/core/gallery/importers/llama-cpp.go index a1cbb6d1bc3c..bc70de74e790 100644 --- a/core/gallery/importers/llama-cpp.go +++ b/core/gallery/importers/llama-cpp.go @@ -37,6 +37,7 @@ func (i *LlamaCPPImporter) AdditionalBackends() []KnownBackendEntry { return []KnownBackendEntry{ {Name: "ik-llama-cpp", Modality: "text", Description: "GGUF drop-in replacement for llama-cpp with ik-quants"}, {Name: "turboquant", Modality: "text", Description: "GGUF drop-in replacement for llama-cpp with TurboQuant optimizations"}, + {Name: "rocmfp4", Modality: "text", Description: "GGUF drop-in replacement for llama-cpp reading ROCmFP4 4-bit weights on AMD RDNA3.5"}, {Name: "vllm-cpp", Modality: "text", Description: "vLLM-style continuous-batching engine (vllm.cpp) consuming GGUF, by the LocalAI team"}, } } @@ -136,7 +137,9 @@ func (i *LlamaCPPImporter) Import(details Details) (gallery.ModelConfig, error) backend := "llama-cpp" if b, ok := preferencesMap["backend"].(string); ok { switch b { - case "ik-llama-cpp", "turboquant", "vllm-cpp": + // Preference-only: ROCmFP4 GGUFs carry ordinary .gguf names, so there is no + // safe auto-detect signal. Without an explicit preference this stays llama-cpp. + case "ik-llama-cpp", "turboquant", "vllm-cpp", "rocmfp4": backend = b } } diff --git a/docs/content/features/backends.md b/docs/content/features/backends.md index 7d2ec66a4f4a..a18487780a73 100644 --- a/docs/content/features/backends.md +++ b/docs/content/features/backends.md @@ -174,6 +174,7 @@ For getting started, see the available backends in LocalAI here: https://github. LocalAI supports various types of backends: - **LLM Backends**: For running language models (e.g., llama.cpp, vLLM, vllm.cpp, SGLang, transformers, MLX) +- **AMD FP4 Backend**: `rocmfp4` — a llama.cpp fork reading the ROCmFP4 / ROCmFPx 4-bit weight formats (ggml types 100-107) on AMD RDNA3.5 APUs (Strix Point / Strix Halo / Gorgon). Stock llama.cpp rejects these tensor types. The format trades memory, not arithmetic: ~22% smaller than Q4_K_M at perplexity parity, with the matmul running as int8 dot products since RDNA3.5 has no FP4 matrix instruction. `llama-cpp` remains the recommendation for everything else. - **Speech-to-Text Backends**: For transcription, forced alignment and speaker diarization (e.g., whisper.cpp, parakeet.cpp, moss-transcribe.cpp, [NeMo-Speech.cpp]({{%relref "features/nemo-speech-cpp" %}}), faster-whisper, NeMo, [audio.cpp]({{%relref "features/audio-cpp" %}})) - **Text-to-Speech Backends**: For speech synthesis (e.g., piper, Kokoro, VibeVoice, Qwen3-TTS, [NeMo-Speech.cpp]({{%relref "features/nemo-speech-cpp" %}}), [audio.cpp]({{%relref "features/audio-cpp" %}})) - **Sound Generation Backends**: For music and audio generation (e.g., ACE-Step, [audio.cpp]({{%relref "features/audio-cpp" %}})) diff --git a/scripts/lib/backend-filter.mjs b/scripts/lib/backend-filter.mjs index 9312cbf6d0c0..6c6e976a7092 100644 --- a/scripts/lib/backend-filter.mjs +++ b/scripts/lib/backend-filter.mjs @@ -78,6 +78,12 @@ export function inferBackendPath(item) { // via a thin wrapper Makefile. Changes to either dir should retrigger it. return `backend/cpp/turboquant/`; } + if (item.dockerfile.endsWith("rocmfp4")) { + // rocmfp4 is a llama.cpp fork carrying the ROCmFP4 weight formats; like + // bonsai it reuses backend/cpp/llama-cpp sources via a wrapper Makefile. + // Must stay above the generic llama-cpp suffix test. + return `backend/cpp/rocmfp4/`; + } if (item.dockerfile.endsWith("bonsai")) { // bonsai is a llama.cpp fork that reuses backend/cpp/llama-cpp sources // via a thin wrapper Makefile. Changes to either dir should retrigger it. @@ -152,7 +158,7 @@ export function backendChanged(backend, pathPrefix, changedFiles) { // Fork backends reuse backend/cpp/llama-cpp sources via thin wrappers; // changes to either directory must retrigger their pipelines. - return (backend === "turboquant" || backend === "bonsai") && + return (backend === "turboquant" || backend === "bonsai" || backend === "rocmfp4") && changedFiles.some(file => file.startsWith("backend/cpp/llama-cpp/")); } From 35f3004638794bbcf9b461d42c1f0f9c3189e29e Mon Sep 17 00:00:00 2001 From: Stefan Walcz Date: Thu, 20 Aug 2026 16:14:40 +0200 Subject: [PATCH 3/8] fix(rocmfp4): use the current RPC target name (rpc-server -> ggml-rpc-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 --- backend/cpp/rocmfp4/Makefile | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/backend/cpp/rocmfp4/Makefile b/backend/cpp/rocmfp4/Makefile index 987cccf8a1ce..5f9db1ef9197 100644 --- a/backend/cpp/rocmfp4/Makefile +++ b/backend/cpp/rocmfp4/Makefile @@ -1,6 +1,6 @@ # Pinned to the HEAD of the `rocmfp4` branch on https://github.com/walcz-de/llama.cpp-ROCmFP4. # Auto-bumped nightly by .github/workflows/bump_deps.yaml. -ROCMFP4_VERSION?=0b8c6fbda8c18f7ae4bd30811996c24f04ddc0dc +ROCMFP4_VERSION?=0b6afd26bfa64f51c14794b4dfbe3da8946076d3 LLAMA_REPO?=https://github.com/walcz-de/llama.cpp-ROCmFP4 CMAKE_ARGS?= @@ -81,10 +81,13 @@ rocmfp4-cpu-all: @echo "Collected ggml shared backends:" && ls -la ggml-shared-libs/ rocmfp4-grpc: - $(call rocmfp4-build,grpc,-DGGML_RPC=ON -DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server --target rpc-server) + $(call rocmfp4-build,grpc,-DGGML_RPC=ON -DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server --target ggml-rpc-server) +# Upstream renamed the RPC binary rpc-server -> ggml-rpc-server. The bonsai wrapper +# this was modelled on still uses the old name because its fork predates the rename; +# ours is based on a current tree, so it must use the new one. rocmfp4-rpc-server: rocmfp4-grpc - cp -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-grpc-build/llama.cpp/build/bin/rpc-server rocmfp4-rpc-server + cp -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-grpc-build/llama.cpp/build/bin/ggml-rpc-server rocmfp4-rpc-server package: bash package.sh From 9e570226dbcdb5b644857211a75b51c0887896a7 Mon Sep 17 00:00:00 2001 From: Stefan Walcz Date: Thu, 20 Aug 2026 20:30:27 +0200 Subject: [PATCH 4/8] test(rocmfp4): importer preference swap + E2E target with a real FP4 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 --- Makefile | 10 ++++++++++ core/gallery/importers/llama-cpp_test.go | 17 +++++++++++++++++ 2 files changed, 27 insertions(+) diff --git a/Makefile b/Makefile index c9068021d8a9..ddde4eb25ab2 100644 --- a/Makefile +++ b/Makefile @@ -750,6 +750,16 @@ test-extra-backend-turboquant: docker-build-turboquant BACKEND_TEST_CACHE_TYPE_V=turbo3 \ $(MAKE) test-extra-backend +## rocmfp4: exercises the llama.cpp-fork backend with a real ROCmFP4 model — the +## published Strix Halo build of Qwen3.8-27B, whose weight quants (ggml types +## 100/103) are *only* decodable by this fork. Loading it is what makes the +## backend distinct from stock llama-cpp. Note the artifact is 13.75 GiB; this +## target is meant for a workstation with the weights cached, not for slim CI. +test-extra-backend-rocmfp4: docker-build-rocmfp4 + BACKEND_IMAGE=local-ai-backend:rocmfp4 \ + BACKEND_TEST_MODEL_URL=https://huggingface.co/kingjones777/Qwen3.8-27B-ROCmFP4-STRIX-MTP-GGUF/resolve/main/Qwen3.8-27B-Q4_0_ROCMFP4_STRIX.gguf \ + $(MAKE) test-extra-backend + ## bonsai: exercises the llama.cpp-fork backend with a real Q1_0 (1-bit) model — ## the PrismML Bonsai-8B GGUF, whose weight quant is *only* decodable by the fork's ## Q1_0 kernels. Loading it is what makes this backend distinct from stock llama-cpp; diff --git a/core/gallery/importers/llama-cpp_test.go b/core/gallery/importers/llama-cpp_test.go index 139b29884bb7..453035e6d05f 100644 --- a/core/gallery/importers/llama-cpp_test.go +++ b/core/gallery/importers/llama-cpp_test.go @@ -181,6 +181,23 @@ var _ = Describe("LlamaCPPImporter", func() { Expect(modelConfig.Files[0].Filename).To(Equal("my-model.gguf")) }) + It("swaps the emitted backend to rocmfp4 when preferred", func() { + preferences := json.RawMessage(`{"backend": "rocmfp4"}`) + details := Details{ + URI: "https://example.com/my-model.gguf", + Preferences: preferences, + } + + modelConfig, err := importer.Import(details) + + Expect(err).ToNot(HaveOccurred()) + Expect(modelConfig.ConfigFile).To(ContainSubstring("backend: rocmfp4"), fmt.Sprintf("Model config: %+v", modelConfig)) + Expect(modelConfig.ConfigFile).NotTo(ContainSubstring("backend: llama-cpp\n"), fmt.Sprintf("Model config: %+v", modelConfig)) + Expect(modelConfig.ConfigFile).To(ContainSubstring("model: my-model.gguf"), fmt.Sprintf("Model config: %+v", modelConfig)) + Expect(len(modelConfig.Files)).To(Equal(1)) + Expect(modelConfig.Files[0].Filename).To(Equal("my-model.gguf")) + }) + It("swaps the emitted backend to vllm-cpp when preferred, keeping engine-side templating", func() { preferences := json.RawMessage(`{"backend": "vllm-cpp"}`) details := Details{ From 3a0c54c89c732cc83d50f73834b5963147242813 Mon Sep 17 00:00:00 2001 From: Stefan Walcz Date: Fri, 21 Aug 2026 08:50:40 +0200 Subject: [PATCH 5/8] =?UTF-8?q?fix(rocmfp4):=20review=20cleanup=20?= =?UTF-8?q?=E2=80=94=20bump=5Fdeps=20indentation,=20rocmfp4=20patches=20RE?= =?UTF-8?q?ADME?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Addresses the maintainer review on #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 --- .github/workflows/bump_deps.yaml | 8 ++++---- backend/cpp/rocmfp4/patches/README.md | 28 ++++++++++++--------------- 2 files changed, 16 insertions(+), 20 deletions(-) diff --git a/.github/workflows/bump_deps.yaml b/.github/workflows/bump_deps.yaml index 38147d3eb82c..b9187732d7ae 100644 --- a/.github/workflows/bump_deps.yaml +++ b/.github/workflows/bump_deps.yaml @@ -26,10 +26,10 @@ jobs: variable: "BONSAI_VERSION" branch: "prism" file: "backend/cpp/bonsai/Makefile" - - repository: "walcz-de/llama.cpp-ROCmFP4" - variable: "ROCMFP4_VERSION" - branch: "rocmfp4" - file: "backend/cpp/rocmfp4/Makefile" + - repository: "walcz-de/llama.cpp-ROCmFP4" + variable: "ROCMFP4_VERSION" + branch: "rocmfp4" + file: "backend/cpp/rocmfp4/Makefile" - repository: "antirez/ds4" variable: "DS4_VERSION" branch: "main" diff --git a/backend/cpp/rocmfp4/patches/README.md b/backend/cpp/rocmfp4/patches/README.md index 3bfb74183f98..e56c93c7000e 100644 --- a/backend/cpp/rocmfp4/patches/README.md +++ b/backend/cpp/rocmfp4/patches/README.md @@ -1,19 +1,15 @@ -# bonsai fork skew patches +# rocmfp4 fork skew patches -The `bonsai` backend reuses `backend/cpp/llama-cpp/grpc-server.cpp` (written against -LocalAI's pinned *upstream* llama.cpp) but compiles it against the PrismML `prism` fork, -which branched from upstream some commits earlier. Any upstream API change that the shared -gRPC server depends on, but that the fork does not yet carry, is back-ported here as a -`*.patch` file and applied to the cloned fork checkout by `../apply-patches.sh`. +The `rocmfp4` backend reuses `backend/cpp/llama-cpp/grpc-server.cpp` (written against +LocalAI's pinned *upstream* llama.cpp) but compiles it against the +[walcz-de/llama.cpp-ROCmFP4](https://github.com/walcz-de/llama.cpp-ROCmFP4) fork, which +carries the ROCmFP4 / ROCmFPx tensor formats (ggml types 100-107) for AMD RDNA3.5 APUs. -CI treats both this directory and `backend/cpp/llama-cpp/` as Bonsai inputs, since -the wrapper copies and builds the shared llama.cpp backend sources. +The fork tracks the same upstream pin as `backend/cpp/llama-cpp` (it is re-based onto the +`LLAMA_VERSION` LocalAI pins), so in the common case **no patches are needed here** and this +directory stays empty. If the pins ever drift — upstream changes an API the shared gRPC +server depends on before the fork has rebased — the gap is back-ported here as a `*.patch` +file and applied to the cloned fork checkout by `../apply-patches.sh`. -Rules: - -- One upstream commit (or minimal hunk) per patch, named `NNNN-short-description.patch`. -- Patches are applied with `git apply` from the fork's checkout root. -- `apply-patches.sh` fails fast if a patch stops applying cleanly — that is the signal the - fork has caught up (or diverged), so re-cut or drop the patch. -- Keep this set as small as possible; the long-term fix is the fork rebasing onto a newer - upstream (or Q1_0/Q2_0 landing in mainline llama.cpp, retiring this backend entirely). +Patches apply with `git apply` from the fork checkout root. Name them +`NNNN-short-description.patch` so the apply order stays deterministic. From df039060626019bfdcf7fe714fe889be3e0734fa Mon Sep 17 00:00:00 2001 From: Stefan Walcz Date: Fri, 21 Aug 2026 10:17:43 +0200 Subject: [PATCH 6/8] fix(rocmfp4): rewrite apply-patches.sh comments and message for this 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 --- backend/cpp/rocmfp4/apply-patches.sh | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/backend/cpp/rocmfp4/apply-patches.sh b/backend/cpp/rocmfp4/apply-patches.sh index 7ac850581262..bdee6ba0627a 100755 --- a/backend/cpp/rocmfp4/apply-patches.sh +++ b/backend/cpp/rocmfp4/apply-patches.sh @@ -1,10 +1,12 @@ #!/bin/bash -# Apply the bonsai patch series to a cloned PrismML llama.cpp (prism branch) checkout. +# Apply the rocmfp4 patch series to a cloned walcz-de/llama.cpp-ROCmFP4 checkout. # -# The prism fork branched from upstream llama.cpp before a number of API changes that the -# shared backend/cpp/llama-cpp/grpc-server.cpp depends on. We carry those upstream commits -# as patch files under backend/cpp/bonsai/patches/ and apply them here so the reused -# grpc-server source compiles against the fork unmodified. +# The rocmfp4 fork is rebased onto the same LLAMA_VERSION LocalAI pins for the shared +# backend/cpp/llama-cpp/grpc-server.cpp, so in the common case patches/ is empty and this +# script is a no-op. If the pins ever drift — upstream changes an API the shared gRPC +# server depends on before the fork has rebased — the gap is carried as patch files under +# backend/cpp/rocmfp4/patches/ and applied here so the reused grpc-server source compiles +# against the fork unmodified. # # Drop the corresponding patch from patches/ whenever the fork catches up with upstream — # the build will fail fast if a patch stops applying, which is the signal to retire it. @@ -45,4 +47,4 @@ for patch in "${patches[@]}"; do git apply --verbose "$patch" done -echo "all bonsai patches applied successfully" +echo "all rocmfp4 patches applied successfully" From f34ebbb0523e7855fbb9a38c68cb91b32b9a5879 Mon Sep 17 00:00:00 2001 From: stefanwalcz Date: Fri, 4 Sep 2026 12:59:07 +0200 Subject: [PATCH 7/8] feat(rocmfp4): ship llama-quantize, drop the paths this backend never 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) Claude-Session: https://claude.ai/code/session_01PHcZ8CTpZEoK3qgroNnkB4 (cherry picked from commit 2286566a972814ccd42612e27fae28d286144c74) --- .docker/rocmfp4-compile.sh | 19 ++++----------- backend/Dockerfile.rocmfp4 | 6 ----- backend/cpp/rocmfp4/Makefile | 38 ++++++----------------------- backend/cpp/rocmfp4/run.sh | 29 ++-------------------- scripts/lib/backend-filter_test.mjs | 23 +++++++++++++++-- 5 files changed, 36 insertions(+), 79 deletions(-) diff --git a/.docker/rocmfp4-compile.sh b/.docker/rocmfp4-compile.sh index 7547ab5f0960..d4dc873a4ce1 100755 --- a/.docker/rocmfp4-compile.sh +++ b/.docker/rocmfp4-compile.sh @@ -25,21 +25,12 @@ fi cd /LocalAI/backend/cpp/rocmfp4 -if [ -z "${BUILD_TYPE:-}" ]; then - # Pure CPU image: one ggml CPU_ALL_VARIANTS build replaces the per-microarch binaries. - # arm64: the armv9.2 SME variants need gcc-14 (gcc-13 rejects +sme). - if [ "${TARGETARCH}" = "arm64" ]; then - apt-get update -qq && apt-get install -y -qq gcc-14 g++-14 - export CC=gcc-14 CXX=g++-14 - fi - make rocmfp4-cpu-all -else - # GPU build (cublas/hipblas/sycl/vulkan/...): single fallback CPU build, the accelerator - # does the compute. Keeps the GPU compile from also building the CPU variant matrix and - # avoids the gcc-14 apt step on GPU base images such as nvidia l4t. - make rocmfp4-fallback -fi +# rocmfp4 is ROCm-only: the single matrix entry is hipblas/amd64, so BUILD_TYPE is always +# set and the CPU-image branch this was templated from is unreachable. One fallback CPU +# build, the accelerator does the compute. +make rocmfp4-fallback make rocmfp4-grpc make rocmfp4-rpc-server +make rocmfp4-quantize ccache -s || true diff --git a/backend/Dockerfile.rocmfp4 b/backend/Dockerfile.rocmfp4 index 7beb88e97685..ee45db0697af 100644 --- a/backend/Dockerfile.rocmfp4 +++ b/backend/Dockerfile.rocmfp4 @@ -48,12 +48,6 @@ ARG LLAMA_REPO= ARG ROCMFP4_VERSION= ENV LLAMA_REPO=${LLAMA_REPO} ENV ROCMFP4_VERSION=${ROCMFP4_VERSION} -# Local-build escape hatch: point the fetch at a git daemon on the host instead of -# GitHub. Defaults are empty, so CI keeps using the Makefile's pinned repo+sha. -ARG LLAMA_REPO= -ARG ROCMFP4_VERSION= -ENV LLAMA_REPO=${LLAMA_REPO} -ENV ROCMFP4_VERSION=${ROCMFP4_VERSION} # CUDA target archs, e.g. --build-arg CUDA_DOCKER_ARCH='75;86;89;120' ARG CUDA_DOCKER_ARCH ARG CMAKE_ARGS diff --git a/backend/cpp/rocmfp4/Makefile b/backend/cpp/rocmfp4/Makefile index 5f9db1ef9197..a01593b63ede 100644 --- a/backend/cpp/rocmfp4/Makefile +++ b/backend/cpp/rocmfp4/Makefile @@ -48,40 +48,11 @@ define rocmfp4-build cp -rfv $(CURRENT_MAKEFILE_DIR)/../rocmfp4-$(1)-build/grpc-server rocmfp4-$(1) endef -rocmfp4-avx2: - $(call rocmfp4-build,avx2,-DGGML_AVX=on -DGGML_AVX2=on -DGGML_AVX512=off -DGGML_FMA=on -DGGML_F16C=on,--target grpc-server) - -rocmfp4-avx512: - $(call rocmfp4-build,avx512,-DGGML_AVX=on -DGGML_AVX2=off -DGGML_AVX512=on -DGGML_FMA=on -DGGML_F16C=on,--target grpc-server) - -rocmfp4-avx: - $(call rocmfp4-build,avx,-DGGML_AVX=on -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server) - rocmfp4-fallback: $(call rocmfp4-build,fallback,-DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server) -# Single-build CPU backend via ggml CPU_ALL_VARIANTS (mirrors llama-cpp-cpu-all). -rocmfp4-cpu-all: - rm -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build - cp -rf $(LLAMA_CPP_DIR) $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build - rm -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/patches - $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build purge - bash $(LLAMA_CPP_DIR)/disable-score-task.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/grpc-server.cpp - bash $(LLAMA_CPP_DIR)/disable-tts-task.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/grpc-server.cpp - $(info $(GREEN)I rocmfp4 build info:cpu-all-variants$(RESET)) - LLAMA_REPO=$(LLAMA_REPO) LLAMA_VERSION=$(ROCMFP4_VERSION) \ - $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build llama.cpp - bash $(CURRENT_MAKEFILE_DIR)/apply-patches.sh $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/llama.cpp $(PATCHES_DIR) - SHARED_LIBS=ON EXTRA_CMAKE_ARGS="-DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON" TARGET="--target grpc-server --target ggml" \ - LLAMA_REPO=$(LLAMA_REPO) LLAMA_VERSION=$(ROCMFP4_VERSION) \ - $(MAKE) -C $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build grpc-server - cp -rfv $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/grpc-server rocmfp4-cpu-all - rm -rf ggml-shared-libs && mkdir -p ggml-shared-libs - find $(CURRENT_MAKEFILE_DIR)/../rocmfp4-cpu-all-build/llama.cpp/build \( -name '*.so*' -o -name '*.dylib' \) -exec cp -av {} ggml-shared-libs/ \; - @echo "Collected ggml shared backends:" && ls -la ggml-shared-libs/ - rocmfp4-grpc: - $(call rocmfp4-build,grpc,-DGGML_RPC=ON -DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server --target ggml-rpc-server) + $(call rocmfp4-build,grpc,-DGGML_RPC=ON -DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server --target ggml-rpc-server --target llama-quantize) # Upstream renamed the RPC binary rpc-server -> ggml-rpc-server. The bonsai wrapper # this was modelled on still uses the old name because its fork predates the rename; @@ -89,6 +60,13 @@ rocmfp4-grpc: rocmfp4-rpc-server: rocmfp4-grpc cp -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-grpc-build/llama.cpp/build/bin/ggml-rpc-server rocmfp4-rpc-server +# The FP4 formats are only usable if people can produce weights, not just consume the +# single published model. The fork carries the quantization recipes and CLI arguments, +# so shipping llama-quantize next to the server turns "one model" into "any GGUF you +# already have". Built by the grpc target above, which needs no accelerator. +rocmfp4-quantize: rocmfp4-grpc + cp -rf $(CURRENT_MAKEFILE_DIR)/../rocmfp4-grpc-build/llama.cpp/build/bin/llama-quantize rocmfp4-quantize + package: bash package.sh diff --git a/backend/cpp/rocmfp4/run.sh b/backend/cpp/rocmfp4/run.sh index 0ccf3cae7137..5c7b2f75f2ff 100755 --- a/backend/cpp/rocmfp4/run.sh +++ b/backend/cpp/rocmfp4/run.sh @@ -12,12 +12,8 @@ grep -e "flags" /proc/cpuinfo | head -1 BINARY=rocmfp4-fallback -# x86/arm64 ship a single rocmfp4-cpu-all built with ggml CPU_ALL_VARIANTS: ggml's -# backend registry dlopens the best libggml-cpu-*.so for this host, so no shell-side -# probing. ROCm ships only rocmfp4-fallback, so fall back to it when cpu-all is absent. -if [ -e "$CURDIR"/rocmfp4-cpu-all ]; then - BINARY=rocmfp4-cpu-all -fi +# ROCm-only backend: the accelerator does the compute, so the CPU side is the single +# rocmfp4-fallback build. No cpu-all variant is produced, hence no probing here. if [ -n "$LLAMACPP_GRPC_SERVERS" ]; then if [ -e "$CURDIR"/rocmfp4-grpc ]; then @@ -40,27 +36,6 @@ else if [ -d "$CURDIR/lib/hipblaslt/library" ]; then export HIPBLASLT_TENSILE_LIBPATH="$CURDIR"/lib/hipblaslt/library fi - # Backends built for Intel GPUs carry a copy of the Intel graphics driver, - # and libze_loader is only there in those builds. Level Zero looks for a - # driver on its own, so point it at the copy that came with this backend: it - # was built against the same C library, while the machine's own driver may - # not have been, and loading that one can crash on start. - # - # Anything the user set is left alone, so a machine with a graphics card - # newer than the driver carried here can still be told to use its own. - # Nothing is said about OpenCL: no OpenCL driver is carried, so anything we - # set there would leave OpenCL worse off than the machine's own setup. - if [ -e "$CURDIR/lib/libze_loader.so.1" ]; then - if [ -e "$CURDIR/lib/libze_intel_gpu.so.1" ] && [ -z "${ZE_ENABLE_ALT_DRIVERS:-}" ]; then - export ZE_ENABLE_ALT_DRIVERS="$CURDIR"/lib/libze_intel_gpu.so.1 - fi - # Ask the driver how much graphics memory is free. Without this, the - # backend reads zero on an integrated graphics chip, because such a chip - # shares the system memory instead of having its own. - if [ -z "${ZES_ENABLE_SYSMAN:-}" ]; then - export ZES_ENABLE_SYSMAN=1 - fi - fi fi # If there is a lib/ld.so, use it diff --git a/scripts/lib/backend-filter_test.mjs b/scripts/lib/backend-filter_test.mjs index f419afd3a37a..5f589a5f3a43 100644 --- a/scripts/lib/backend-filter_test.mjs +++ b/scripts/lib/backend-filter_test.mjs @@ -65,6 +65,12 @@ const includes = [ "tag-suffix": "-turboquant", "base-image": "ubuntu:24.04", }, + { + backend: "rocmfp4", + dockerfile: "./backend/Dockerfile.rocmfp4", + "tag-suffix": "-rocmfp4", + "base-image": "ubuntu:24.04", + }, { backend: "audio-cpp", dockerfile: "./backend/Dockerfile.audio-cpp", @@ -213,10 +219,23 @@ test("tests for the packaging scripts do not rebuild anything", () => { assert.deepEqual(filteredDarwin, []); }); -test("turboquant still retriggers on llama-cpp source changes", () => { +test("llama.cpp forks still retrigger on llama-cpp source changes", () => { const { filtered } = run(["backend/cpp/llama-cpp/grpc-server.cpp"]); - assert.deepEqual(names(filtered), ["llama-cpp", "turboquant"]); + assert.deepEqual(names(filtered), ["llama-cpp", "rocmfp4", "turboquant"]); +}); + +// --------------------------------------------------------------------------- +// rocmfp4: its inferBackendPath case must sit ABOVE the generic llama-cpp +// suffix test. Without it "Dockerfile.rocmfp4" would not resolve to its own +// directory and a PR touching only this backend would get no CI job at all. +// --------------------------------------------------------------------------- + +test("rocmfp4 source changes rebuild rocmfp4 and nothing else", () => { + const { filtered, changedBackends } = run(["backend/cpp/rocmfp4/Makefile"]); + + assert.deepEqual(names(filtered), ["rocmfp4"]); + assert.ok(changedBackends.has("rocmfp4")); }); // --------------------------------------------------------------------------- From fd1944ccde7c59eb411b735b3e39b12503142a4e Mon Sep 17 00:00:00 2001 From: stefanwalcz Date: Fri, 4 Sep 2026 13:02:10 +0200 Subject: [PATCH 8/8] feat(rocmfp4): build the grpc target with CPU vector instructions on 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) Claude-Session: https://claude.ai/code/session_01PHcZ8CTpZEoK3qgroNnkB4 (cherry picked from commit b784a35b10630cd4cd3d7d2f4011663beb289a01) --- backend/cpp/rocmfp4/Makefile | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/backend/cpp/rocmfp4/Makefile b/backend/cpp/rocmfp4/Makefile index a01593b63ede..e3db34e81435 100644 --- a/backend/cpp/rocmfp4/Makefile +++ b/backend/cpp/rocmfp4/Makefile @@ -51,8 +51,13 @@ endef rocmfp4-fallback: $(call rocmfp4-build,fallback,-DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server) +# CPU features are ON here, unlike the portable fallback build. This backend only +# targets gfx1150-1153, which are APU integrated GPUs: the CPU sits on the same +# package and is always Zen 4/5, so AVX/AVX2/AVX512/FMA/F16C/BMI2 are guaranteed +# present and nothing portable is lost. It matters because this build also produces +# llama-quantize, and quantization is pure CPU work over every weight in the model. rocmfp4-grpc: - $(call rocmfp4-build,grpc,-DGGML_RPC=ON -DGGML_AVX=off -DGGML_AVX2=off -DGGML_AVX512=off -DGGML_FMA=off -DGGML_F16C=off -DGGML_BMI2=off,--target grpc-server --target ggml-rpc-server --target llama-quantize) + $(call rocmfp4-build,grpc,-DGGML_RPC=ON -DGGML_AVX=on -DGGML_AVX2=on -DGGML_AVX512=on -DGGML_FMA=on -DGGML_F16C=on -DGGML_BMI2=on,--target grpc-server --target ggml-rpc-server --target llama-quantize) # Upstream renamed the RPC binary rpc-server -> ggml-rpc-server. The bonsai wrapper # this was modelled on still uses the old name because its fork predates the rename;