Add AMD ROCm/HIP backend for SmolVLA and π0.5 on gfx1151

CHANGELOG.md

@@ -6,6 +6,11 @@ Notable changes to vla.cpp. Format loosely follows [Keep a Changelog](https://ke
 
 ### Added
 
+- **ROCm/HIP backend for VLA inference.** `-DGGML_HIP=ON` selects ggml's HIP
+  backend without compiling vla.cpp's CUDA-only kernels. Linux `gfx1151` with
+  ROCm 7.14 and llama.cpp `b11223` has fixed-input and synthetic benchmark
+  coverage for SmolVLA and π0.5. See `docs/backend/rocm.md` for the measured
+  numerical differences, build commands and current limits.
 - **FoldQuant W8A8 / W4A4 inference** for GR00T N1.5 / N1.6 / N1.7 and π0.5. A
   FoldQuant GGUF carries the language backbone and the action
   module as INT8 or INT4 codes with per-row scales in a block-Hadamard,

CMakeLists.txt

@@ -19,7 +19,7 @@ elseif(UNIX)
 endif()
 
 set(_vla_accel "")
-foreach(_flag GGML_CUDA GGML_SYCL GGML_METAL GGML_OPENVINO GGML_HEXAGON GGML_OPENCL)
+foreach(_flag GGML_CUDA GGML_HIP GGML_VULKAN GGML_SYCL GGML_METAL GGML_OPENVINO GGML_HEXAGON GGML_OPENCL)
     if(${_flag})
         list(APPEND _vla_accel ${_flag})
     endif()
@@ -39,7 +39,7 @@ if(_vla_accel_n GREATER 0 AND GGML_BACKEND_DL)
         "GGML_BACKEND_DL=ON is not supported with ${_vla_accel}: the backend is "
         "built as a module and src/backend.h links its init directly.")
 endif()
-foreach(_flag GGML_VULKAN GGML_HIP GGML_MUSA GGML_CANN GGML_WEBGPU)
+foreach(_flag GGML_VULKAN GGML_MUSA GGML_CANN GGML_WEBGPU)
     if(${_flag})
         message(WARNING "${_flag}=ON reaches vlm-server through llama only; src/backend.h has no ${_flag} path, so the VLA archs run on CPU.")
     endif()
@@ -345,6 +345,9 @@ if(GGML_CUDA)
     )
     target_link_libraries(vla_cuda_ops PUBLIC CUDA::cublas CUDA::cudart ggml-cuda vla_fq_kernels)
     target_link_libraries(vla_core PRIVATE vla_cuda_ops)
+elseif(GGML_HIP)
+    # ggml-hip also exports GGML_USE_CUDA; distinguish HIP from CUDA-only code.
+    target_compile_definitions(vla_core PUBLIC GGML_USE_HIP)
 endif()
 
 if(GGML_SYCL AND NOT GGML_CUDA)

README.md

@@ -12,8 +12,9 @@ A C++ inference engine for **Vision-Language-Action (VLA) models**, built on [`l
 It runs the open VLA policies - SmolVLA, π0, BitVLA, Evo-1, GR00T N1.5/1.6/1.7 and more -
 under one runtime, each packaged as a single self-contained GGUF that needs no Python or
 PyTorch at inference time. The binaries drive robots on **CPU**, **Apple Silicon**, **CUDA** -
-from consumer GPUs down to Jetson-class boards - **Intel GPUs and NPUs** via
-SYCL and OpenVINO, or **Qualcomm Snapdragon CPUs, Adreno GPUs and Hexagon NPUs**
+from consumer GPUs down to Jetson-class boards - **AMD Radeon 8060S** via ROCm/HIP,
+**Intel GPUs and NPUs** via SYCL and OpenVINO, or **Qualcomm Snapdragon CPUs,
+Adreno GPUs and Hexagon NPUs**
 via OpenCL and the Hexagon backend.
 
 [**Learn vla.cpp**](https://fai-modelopt-tech.github.io/learn-vla-cpp/) walks through the engine design and how each policy is implemented on ggml.
@@ -57,6 +58,8 @@ tarball). Details, and the macOS and Windows notes, are in
 - CMake ≥ 3.22
 - A C++17 compiler (GCC 11+ or Clang 14+)
 - CUDA 12.x or 13.x (optional - required only for CUDA GPU builds)
+- ROCm with HIP, hipBLAS and rocBLAS (optional - see
+  [the tested AMD configuration](docs/backend/rocm.md))
 - Intel oneAPI 2025.x + GPU compute runtime (optional - only for Intel GPU
   builds, see [docs/backend/sycl.md](docs/backend/sycl.md))
 - OpenVINO 2026.x runtime (optional - only for Intel CPU/GPU/NPU builds via
@@ -97,6 +100,19 @@ cmake -B build \
 cmake --build build -j$(nproc)
 ```
 
+For ROCm/HIP, use a separate build directory and the target reported by
+`rocminfo` (`gfx1151` in the tested example):
+
+```bash
+HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" \
+cmake -B build-rocm -G Ninja -DGGML_HIP=ON -DGPU_TARGETS=gfx1151 \
+    -DCMAKE_BUILD_TYPE=Release
+cmake --build build-rocm -j$(nproc)
+```
+
+The [ROCm backend notes](docs/backend/rocm.md) cover the runtime library path,
+device selection, validation and limits.
+
 If CMake cannot find CUDA, point the environment at it explicitly:
 
 ```bash
@@ -114,7 +130,7 @@ tree. `pip install ./bindings/python` builds the Python bindings, see
 [bindings/python/README.md](bindings/python/README.md).
 
 Check [docs/backend](docs/backend) for compiling `vla.cpp` on other platforms.
-WSL2, Apple Silicon, and Intel GPU are all tested.
+ROCm on Linux `gfx1151`, WSL2, Apple Silicon, and Intel GPU are all tested.
 To build and run in containers instead, see [docs/DOCKER.md](docs/DOCKER.md).
 
 ---
@@ -148,21 +164,26 @@ variables are in [docs/USAGE.md](docs/USAGE.md).
 Models (rows) against platforms (columns). Legend: `Y` =
 supported (released and benchmarked), `~` = in progress, `-` = planned.
 
-| Model | CPU (x86-64 / ARM) | CUDA | [SYCL (Intel)](docs/backend/sycl.md) | [Metal](docs/backend/metal.md) | [OpenVINO](docs/backend/ov.md) | [Hexagon](docs/backend/hexagon-windows.md) |
-|---|:--:|:--:|:--:|:--:|:--:|:--:|
-| [SmolVLA](https://hf.co/vrfai/smolvla-libero-gguf)             | Y | Y | Y | Y | Y | Y | 
-| [π0](https://hf.co/vrfai/pi0-libero-finetuned-v044-gguf)       | Y | Y | - | Y | Y | ~ | 
-| [π0.5](https://hf.co/vrfai/pi05-libero-gguf)                   | Y | Y | - | Y | Y | Y | 
-| [GR00T N1.5](https://hf.co/vrfai/gr00tn1d5-libero-object-gguf) | Y | Y | - | Y | Y | Y | 
-| [GR00T N1.6](https://hf.co/vrfai/gr00tn1d6-libero-gguf)        | Y | Y | - | Y | Y | Y | 
-| [GR00T N1.7](https://hf.co/vrfai/gr00tn1d7-libero-gguf)        | Y | Y | - | Y | Y | Y | 
-| [BitVLA](https://hf.co/vrfai/bitvla-libero-gguf)               | Y | Y | - | ~ | - | - | 
-| [Evo-1](https://hf.co/vrfai/evo1-libero-gguf)                  | Y | Y | Y | Y | Y | Y | 
-| [VLA-Adapter](https://hf.co/vrfai/vla-adapter-libero-gguf)     | Y | Y | ~ | Y | Y | Y | 
-| [OpenVLA-OFT](https://hf.co/vrfai/openvla-oft-libero-gguf)     | Y | Y | - | Y | Y | - | 
-| [VLA-JEPA](https://hf.co/vrfai/vla-jepa-libero)                | Y | Y | - | Y | Y | ~ | 
-| [Octo-Small](https://hf.co/vrfai/octo-small-libero-gguf)       | Y | Y | Y | Y | - | Y | 
-| [TurboVLA](https://hf.co/vrfai/turbovla-libero-gguf)           | Y | Y | Y | Y | Y | Y | 
+| Model | CPU (x86-64 / ARM) | CUDA | [ROCm (AMD)](docs/backend/rocm.md) | [SYCL (Intel)](docs/backend/sycl.md) | [Metal](docs/backend/metal.md) | [OpenVINO](docs/backend/ov.md) | [Hexagon](docs/backend/hexagon-windows.md) |
+|---|:--:|:--:|:--:|:--:|:--:|:--:|:--:|
+| [SmolVLA](https://hf.co/vrfai/smolvla-libero-gguf)             | Y | Y | Y | Y | Y | Y | Y |
+| [π0](https://hf.co/vrfai/pi0-libero-finetuned-v044-gguf)       | Y | Y | ~ | - | Y | Y | ~ |
+| [π0.5](https://hf.co/vrfai/pi05-libero-gguf)                   | Y | Y | Y | - | Y | Y | Y |
+| [GR00T N1.5](https://hf.co/vrfai/gr00tn1d5-libero-object-gguf) | Y | Y | ~ | - | Y | Y | Y |
+| [GR00T N1.6](https://hf.co/vrfai/gr00tn1d6-libero-gguf)        | Y | Y | ~ | - | Y | Y | Y |
+| [GR00T N1.7](https://hf.co/vrfai/gr00tn1d7-libero-gguf)        | Y | Y | ~ | - | Y | Y | Y |
+| [BitVLA](https://hf.co/vrfai/bitvla-libero-gguf)               | Y | Y | - | - | ~ | - | - |
+| [Evo-1](https://hf.co/vrfai/evo1-libero-gguf)                  | Y | Y | ~ | Y | Y | Y | Y |
+| [VLA-Adapter](https://hf.co/vrfai/vla-adapter-libero-gguf)     | Y | Y | ~ | ~ | Y | Y | Y |
+| [OpenVLA-OFT](https://hf.co/vrfai/openvla-oft-libero-gguf)     | Y | Y | - | - | Y | Y | - |
+| [VLA-JEPA](https://hf.co/vrfai/vla-jepa-libero)                | Y | Y | - | - | Y | Y | ~ |
+| [Octo-Small](https://hf.co/vrfai/octo-small-libero-gguf)       | Y | Y | - | Y | Y | - | Y |
+| [TurboVLA](https://hf.co/vrfai/turbovla-libero-gguf)           | Y | Y | ~ | Y | Y | Y | Y |
+
+ROCm `Y` currently means fixed-input inference and synthetic-input benchmark
+coverage on Linux `gfx1151` with ROCm 7.14 and llama.cpp `b11223`. The `~`
+rows await validation on this upstream revision; older results do not establish
+their current status. These engine checks do not establish robot-task success.
 
 ---
 
@@ -228,7 +249,7 @@ Wiring, recording, training and queue sizing are in the
 | [docs/QUANTIZATION.md](docs/QUANTIZATION.md) | FoldQuant W8A8 / W4A4: the GGUF contract, the integer arithmetic, converting a FoldQuantVLA quantized model |
 | [docs/DOCKER.md](docs/DOCKER.md) | Building and running the eval in containers |
 | [docs/ARCHITECTURE.md](docs/ARCHITECTURE.md) | Engine design: layers, the prediction path, backends, adding an architecture |
-| [docs/backend/](docs/backend) | Per-backend build and run notes: [SYCL](docs/backend/sycl.md), [OpenVINO](docs/backend/ov.md), [Metal](docs/backend/metal.md), [Hexagon](docs/backend/hexagon.md), [Hexagon on Windows](docs/backend/hexagon-windows.md), [WSL2](docs/backend/wsl.md) |
+| [docs/backend/](docs/backend) | Per-backend build and run notes: [ROCm](docs/backend/rocm.md), [SYCL](docs/backend/sycl.md), [OpenVINO](docs/backend/ov.md), [Metal](docs/backend/metal.md), [Hexagon](docs/backend/hexagon.md), [Hexagon on Windows](docs/backend/hexagon-windows.md), [WSL2](docs/backend/wsl.md) |
 | [docs/benchmark/](docs/benchmark) | Per-device latency and memory for every model, and the fastest flags per device |
 | [docs/KNOWN_ISSUES.md](docs/KNOWN_ISSUES.md) | Known issues and their resolutions |
 | [docs/ADOPTION.md](docs/ADOPTION.md) | C ABI, Python bindings, release packaging, and what is left |

docs/ARCHITECTURE.md

@@ -2,8 +2,8 @@
 
 vla.cpp runs Vision-Language-Action (VLA) policies on the ggml/llama.cpp runtime.
 Every model is a self-contained GGUF that the engine loads, detects, and drives on
-CPU, CUDA, Metal, SYCL, OpenVINO, OpenCL or Hexagon. This page is the map; the
-source is the detail.
+CPU, CUDA, ROCm/HIP, Metal, SYCL, OpenVINO, OpenCL or Hexagon. This page is the
+map; the source is the detail.
 
 ## Layers
 

docs/backend/rocm.md

@@ -0,0 +1,147 @@
+# ROCm/HIP backend
+
+`vla.cpp` can run a VLA model through the HIP backend in the pinned
+`llama.cpp`/ggml dependency. This page records the validated configuration and
+the current model coverage. ROCm is selected at configure time, in a separate
+build directory; it is not selected automatically on an AMD machine.
+
+The results below are from upstream `main@1adf078` merged with this HIP
+backend on Linux 6.17.0-40, Ryzen AI Max+ 395 / Radeon 8060S (`gfx1151`),
+ROCm 7.14.60850, and the repository's llama.cpp tag `b11223` (ggml commit
+`4da6337`). Other GPUs and ROCm versions have not been measured here.
+
+## Build and select the device
+
+Install HIP, hipBLAS, rocBLAS, Ninja, and the host dependencies in the main
+README. Confirm the GPU target with `rocminfo`; substitute that target for
+`gfx1151` if using another GPU.
+
+```bash
+hipconfig --version
+rocminfo | grep -m1 gfx
+HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" \
+cmake -B build-rocm -G Ninja -DCMAKE_BUILD_TYPE=Release \
+    -DGGML_HIP=ON -DGPU_TARGETS=gfx1151 \
+    -DVLA_BUILD_TESTS=ON -DVLA_BUILD_SERVER=OFF -DVLA_SPM=OFF
+cmake --build build-rocm -j$(nproc)
+```
+
+The flags above match this validation. `VLA_BUILD_SERVER=OFF` avoids the
+protobuf and ZeroMQ build dependencies; it does not disable `vla-cli`,
+`vla-bench`, or the test harness. The server and simulator path was not
+retested on this revision.
+
+If ROCm is installed outside the system library search path, export its
+library directory before running CTest or any binary. On the test system,
+CTest initially failed to load `libhipblas.so.3` until this path was set.
+
+```bash
+export VLA_ROCM_ROOT="$(hipconfig -R)"
+export LD_LIBRARY_PATH="$VLA_ROCM_ROOT/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}"
+ctest --test-dir build-rocm --output-on-failure
+```
+
+`VLA_DEVICE=<n>` chooses the HIP device ordinal (default 0). A successful
+model load prints `backend = ROCm/HIP (device 0: AMD Radeon 8060S Graphics)`.
+If initialization fails, the core logs the reason and uses CPU. An invalid
+ordinal (`VLA_DEVICE=999`) was checked: it fell back to CPU and returned the
+same fixed-input action bytes as the CPU build.
+
+ggml-hip still exposes CUDA-named entry points and may print `CUDA graph`
+messages. The backend banner identifies the actual device. vla.cpp's own
+CUDA-only BF16 and BitVLA kernels are excluded from the HIP build; the tested
+`libvla_core.so` linked `libggml-hip`, hipBLAS, rocBLAS and the HIP runtime,
+with no CUDA runtime or cuBLAS dependency. CMake rejects HIP combined with
+CUDA or Vulkan, and rejects `GGML_BACKEND_DL=ON` with HIP.
+
+## Fixed-input validation
+
+`vla_predict_check` fixes images, language tokens, state and diffusion noise.
+Both the CPU reference and HIP run used BF16 resident weights.
+The CPU build of this candidate produced byte-identical actions to upstream
+`main@1adf078` for both rows below. Each HIP row returned the complete
+50 × 32 action array, with finite values and identical action bytes across
+five independent processes. CPU and HIP Release builds each passed all ten
+first-party CTest cases.
+
+| Published GGUF | Views × image size | CPU reference | HIP vs CPU max abs | RMS | Cosine |
+|---|---:|---|---:|---:|---:|
+| SmolVLA LIBERO | 2 × 512 | BF16 | 1.992e-3 | 2.421e-4 | 0.9999991 |
+| π0.5 LIBERO | 2 × 224 | BF16 | 8.792e-4 | 1.169e-4 | 0.9999998 |
+
+The GGUFs came from `vrfai/smolvla-libero-gguf` (local file
+`smolvla-libero.gguf`, SHA-256
+`6fb2d475c98b4c2cef3e27c4eff4e67b483740cbf983fff320a3b8a5e5f74fe8`)
+and `vrfai/pi05-libero-gguf` (local file `pi05-libero.gguf`, SHA-256
+`f9b585b12cbe56bc6b531d40cc45dd1f9036478d5b9f200c1b32aaf323ef6471`).
+The comparisons are against the same GGUF and weight dtype on CPU; they are
+engine checks, not LIBERO success-rate measurements.
+
+To repeat the fixed-input checks, build CPU with the same `b11223` pin and
+`-DVLA_BUILD_TESTS=ON`, then run the corresponding commands on both builds:
+
+```bash
+export SMOLVLA_GGUF=/path/to/smolvla-libero.gguf
+export PI05_GGUF=/path/to/pi05-libero.gguf
+VLA_IMG_SIZE=512 ./build-rocm/tests/vla_predict_check "$SMOLVLA_GGUF" '' 2 --weight-dtype bf16 > smolvla-hip.txt
+VLA_IMG_SIZE=224 ./build-rocm/tests/vla_predict_check "$PI05_GGUF" '' 2 --weight-dtype bf16 > pi05-hip.txt
+VLA_IMG_SIZE=512 ./build-cpu/tests/vla_predict_check "$SMOLVLA_GGUF" '' 2 --weight-dtype bf16 > smolvla-cpu.txt
+VLA_IMG_SIZE=224 ./build-cpu/tests/vla_predict_check "$PI05_GGUF" '' 2 --weight-dtype bf16 > pi05-cpu.txt
+```
+
+The tools print the `action_len=` line followed by action values; compare that
+portion of each output. Use five fresh HIP processes to check repeatability.
+
+## Synthetic-input engine latency
+
+These are in-process `predict()` timings with fixed synthetic inputs. They
+exclude transport and simulator work. Each row used three fresh processes,
+each with three warmups and 20 timed calls, BF16 resident weights, GPU
+performance level `high` (2.9 GHz observed before and after the rounds), and
+CPU governor `performance`. The stock `vla-bench` reports P50, P90 and mean
+Vision time for each round. The table uses the median of the three round P50
+values, the worst round P90, and the median of the three round Vision means.
+All values are milliseconds.
+
+| Model | Views × size | Tokens | P50 | Worst P90 | Vision mean |
+|---|---:|---:|---:|---:|---:|
+| SmolVLA LIBERO | 2 × 512 | 48 | 395.0 | 396.9 | 210.3 |
+| π0.5 LIBERO | 2 × 224 | 128 | 389.1 | 403.1 | 47.1 |
+
+The π0.5 rounds had P50 values 388.1, 389.1 and 389.3 ms; the worst P90 was
+403.1 ms. These results use upstream `main@1adf078` and llama.cpp `b11223`.
+Do not compare them as a speedup against an older revision without a paired
+benchmark.
+
+Run each command three times in fresh processes, changing the log path for
+each round:
+
+```bash
+./build-rocm/vla-bench --ckpt "$SMOLVLA_GGUF" --images 2 --size 512 \
+    --tokens 48 --weight-dtype bf16 --warmup 3 --reps 20 \
+    > smolvla-round1.log 2>&1
+./build-rocm/vla-bench --ckpt "$PI05_GGUF" --images 2 --size 224 \
+    --tokens 128 --weight-dtype bf16 --warmup 3 --reps 20 \
+    > pi05-round1.log 2>&1
+```
+
+Within a round, `vla-bench` calculates P50 and P90 from sorted samples with
+linear interpolation at `p × (n - 1)`. Keep the logs and an environment
+manifest when presenting new results.
+
+## Current limits
+
+- The README's ROCm `Y` entries cover only the two configurations measured
+  above on Linux `gfx1151`. Other models tested on the earlier `b10729` pin
+  need fresh numerical and benchmark checks on `b11223` before being marked
+  supported. π0's HIP numerical path also remains under investigation.
+- BitVLA's published int2 GGUF requires vla.cpp's CUDA-only kernels. That
+  path is not available in a HIP build.
+- FoldQuant quantized GGUFs added on the newer mainline were not measured on
+  HIP. This PR does not add native HIP FoldQuant integer kernels.
+- The server/client route, task success rate, and long-duration soak have not
+  been rerun on this upstream revision. Synthetic tokens and images do not
+  establish policy equivalence in a robot task.
+- The VLA core uses one HIP backend for a full graph and has no per-op CPU
+  scheduler fallback. An unsupported HIP op fails prediction instead of
+  silently running part of the graph on CPU.

src/backend.h

@@ -21,12 +21,12 @@
  * instead; the ladder lives here once.
  *
  * Exactly one accelerator is compiled in, picked by the CMake flag that was
- * used (`GGML_CUDA` / `GGML_SYCL` / `GGML_METAL` / `GGML_OPENVINO` /
- * `GGML_HEXAGON` / `GGML_OPENCL`). The core drives a single backend through
- * `gallocr` rather than a scheduler. The first four run every op the archs
- * build, so there is no per-op CPU fallback for them. Hexagon and OpenCL do
- * not, and they are wrapped by @ref fallback_backend_new, which runs the ops
- * they reject on the CPU.
+ * used (`GGML_CUDA` / `GGML_HIP` / `GGML_SYCL` / `GGML_METAL` /
+ * `GGML_OPENVINO` / `GGML_HEXAGON` / `GGML_OPENCL`). The core drives a single
+ * backend through `gallocr` rather than a scheduler. CUDA, HIP, SYCL, Metal and
+ * OpenVINO have no per-op CPU fallback: an unsupported op fails prediction.
+ * Hexagon and OpenCL are wrapped by @ref fallback_backend_new, which runs the
+ * ops they reject on the CPU.
  */
 
 #pragma once
@@ -242,8 +242,8 @@ constexpr bool default_flash_attn() {
 #endif
 }
 
-/// GPU ordinal for CUDA and SYCL; `VLA_DEVICE` overrides. Junk is rejected, not
-/// silently read as device 0.
+/// GPU ordinal for CUDA, HIP and SYCL; `VLA_DEVICE` overrides. Junk is rejected,
+/// not silently read as device 0.
 inline int backend_device_index() {
     const char * e = std::getenv("VLA_DEVICE");
     if (!e || !*e)
@@ -267,7 +267,7 @@ inline int backend_device_index() {
 inline Backend backend_init(const char * tag, int n_threads) {
     Backend b;
 
-#ifdef GGML_USE_CUDA
+#if defined(GGML_USE_CUDA) && !defined(GGML_USE_HIP)
     {
         const int dev = backend_device_index();
         b.handle = ggml_backend_cuda_init(dev);
@@ -278,6 +278,19 @@ inline Backend backend_init(const char * tag, int n_threads) {
             std::fprintf(stderr, "%s: ggml_backend_cuda_init failed; falling back to CPU\n", tag);
         }
     }
+#elif defined(GGML_USE_HIP)
+    {
+        // ggml-hip compiles the ggml-cuda sources and keeps their entry points.
+        const int dev = backend_device_index();
+        b.handle = ggml_backend_cuda_init(dev);
+        if (b.handle) {
+            char desc[256] = { 0 };
+            ggml_backend_cuda_get_device_description(dev, desc, sizeof(desc));
+            std::printf("%s: backend = ROCm/HIP (device %d: %s)\n", tag, dev, desc);
+        } else {
+            std::fprintf(stderr, "%s: ggml_backend_cuda_init failed on ROCm; falling back to CPU\n", tag);
+        }
+    }
 #elif defined(GGML_USE_SYCL)
     {
         // ggml-sycl's VMM pool hands out virtual-memory-backed pointers that

src/cuda/vla_cuda_ops.h

@@ -24,7 +24,7 @@
 
 namespace vla {
 
-#ifdef GGML_USE_CUDA
+#if defined(GGML_USE_CUDA) && !defined(GGML_USE_HIP)
 void cuda_register_bf16_ops();
 void cuda_register_foldquant_ops();
 #else