From b1f0fa8faa613d982c297f59f0dc5a0782aa9c3a Mon Sep 17 00:00:00 2001
From: singlecatlmx <2797705140@qq.com>
Date: Sat, 3 Oct 2026 20:27:43 +0800
Subject: [PATCH 1/4] Add ROCm/HIP backend selection for VLA models

On Linux gfx1151 with ROCm 7.14.60850 and llama.cpp b11223, CPU and HIP Release builds pass 8/8 CTest cases. SmolVLA and pi0.5 CPU BF16 actions remain byte-identical to upstream main@8eec763; HIP actions are finite and byte-identical across five fresh runs each. HIP-vs-CPU BF16 max abs is 1.992e-3 for SmolVLA and 8.792e-4 for pi0.5.

An invalid HIP device ordinal falls back to CPU with identical action bytes. CMake rejects HIP with CUDA/Vulkan or dynamic backend loading; HIP builds exclude vla.cpp CUDA objects and NVIDIA runtime links. Synthetic-input HIP P50 is 395.0 ms for SmolVLA and 387.8 ms for pi0.5 under the recorded high clock setting. Other models and task success need validation on this upstream revision.
---
 CMakeLists.txt          |  7 +++++--
 src/backend.h           | 31 ++++++++++++++++++++++---------
 src/cuda/vla_cuda_ops.h |  2 +-
 3 files changed, 28 insertions(+), 12 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index ce078c1..481c699 100644
--- a/CMakeLists.txt
+++ b/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()
@@ -285,6 +285,9 @@ if(GGML_CUDA)
     )
     target_link_libraries(vla_cuda_ops PUBLIC CUDA::cublas CUDA::cudart ggml-cuda)
     target_link_libraries(vla_core PRIVATE vla_cuda_ops)
+elseif(GGML_HIP)
+    # ggml-hip exports GGML_USE_CUDA publicly; keep vla.cpp's own CUDA code out.
+    target_compile_definitions(vla_core PUBLIC VLA_USE_HIP)
 endif()
 
 if(GGML_SYCL AND NOT GGML_CUDA)
diff --git a/src/backend.h b/src/backend.h
index 7f61a95..87b0ba0 100644
--- a/src/backend.h
+++ b/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
@@ -244,8 +244,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)
@@ -269,7 +269,20 @@ inline int backend_device_index() {
 inline Backend backend_init(const char * tag, int n_threads) {
     Backend b;
 
-#ifdef GGML_USE_CUDA
+#if defined(VLA_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_CUDA)
     {
         const int dev = backend_device_index();
         b.handle = ggml_backend_cuda_init(dev);
diff --git a/src/cuda/vla_cuda_ops.h b/src/cuda/vla_cuda_ops.h
index b6085b6..6f15968 100644
--- a/src/cuda/vla_cuda_ops.h
+++ b/src/cuda/vla_cuda_ops.h
@@ -21,7 +21,7 @@
 
 namespace vla {
 
-#ifdef GGML_USE_CUDA
+#if defined(GGML_USE_CUDA) && !defined(VLA_USE_HIP)
 void cuda_register_bf16_ops();
 #else
 inline void cuda_register_bf16_ops() {}

From 9803295622c44628eae0e167ffa2403cfdf1918b Mon Sep 17 00:00:00 2001
From: singlecatlmx <2797705140@qq.com>
Date: Sat, 3 Oct 2026 21:28:32 +0800
Subject: [PATCH 2/4] Document validated gfx1151 ROCm/HIP coverage

Limit the ROCm support matrix to SmolVLA and pi0.5 on Linux gfx1151, ROCm 7.14.60850 and llama.cpp b11223. CPU/HIP Release CTest passes 8/8; CPU actions match upstream main@8eec763 byte for byte, and five HIP runs per model are finite and identical. HIP-vs-CPU BF16 max abs is 1.992e-3 and 8.792e-4 respectively.

Document GGUF SHA-256 values, build/runtime paths, fallback checks and three-round synthetic latency using stock vla-bench P50, P90 and Vision mean. Other architectures, server/client execution and robot-task success remain unverified on this revision.
---
 CHANGELOG.md         |   8 +++
 README.md            |  59 +++++++++++------
 docs/ARCHITECTURE.md |   4 +-
 docs/backend/rocm.md | 146 +++++++++++++++++++++++++++++++++++++++++++
 4 files changed, 196 insertions(+), 21 deletions(-)
 create mode 100644 docs/backend/rocm.md

diff --git a/CHANGELOG.md b/CHANGELOG.md
index 5164360..cf86d39 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -4,6 +4,14 @@ Notable changes to vla.cpp. Format loosely follows [Keep a Changelog](https://ke
 
 ## [Unreleased]
 
+### 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.
+
 ## [0.4.0] - 2026-09-30
 
 ### Added
diff --git a/README.md b/README.md
index 9ea731d..0b00e70 100644
--- a/README.md
+++ b/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.
 
 ---
 
@@ -227,7 +248,7 @@ Wiring, recording, training and queue sizing are in the
 | [docs/MODELS.md](docs/MODELS.md) | Converting safetensors checkpoints to GGUF, quantizing to Q8_0/Q4_0 |
 | [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 |
diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md
index 4ce8426..46a2012 100644
--- a/docs/ARCHITECTURE.md
+++ b/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
 
diff --git a/docs/backend/rocm.md b/docs/backend/rocm.md
new file mode 100644
index 0000000..f987e9c
--- /dev/null
+++ b/docs/backend/rocm.md
@@ -0,0 +1,146 @@
+# 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 Linux 6.17.0-40 on a 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@8eec763` 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 eight
+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 shown in the before/after manifests), and
+CPU governor `performance`. The GPU process probe saw no running process
+before or after the rounds. 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 | 397.5 | 210.2 |
+| π0.5 LIBERO | 2 × 224 | 128 | 387.8 | 403.8 | 46.8 |
+
+The π0.5 rounds had P50 values 386.3, 390.4 and 387.8 ms; the worst P90 was
+403.8 ms. These results use upstream v0.4.0's model code 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.
+- 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.

From 30fb53711671e3d852d429d9eb30f48367c6888e Mon Sep 17 00:00:00 2001
From: singlecatlmx <2797705140@qq.com>
Date: Sat, 3 Oct 2026 20:44:07 +0800
Subject: [PATCH 3/4] Use GGML_USE_HIP for VLA HIP backend guards

ggml-hip also exports GGML_USE_CUDA. Define GGML_USE_HIP for vla_core and use it to select the HIP rung and exclude VLA CUDA-only BF16 code.

Verified on gfx1151, ROCm 7.14.60850 and llama.cpp b11223: CPU and HIP Release rebuilds, 8/8 CTest each. SmolVLA and pi0.5 BF16 fixed-input action bytes match their pre-change CPU/HIP baselines; the HIP banner remains ROCm/HIP. Numeric output did not move. The prior CPU/HIP max-abs differences were 1.992e-3 and 8.792e-4. Performance was not remeasured for this macro-only change; the A9 3x20 engine baseline remains the reference.
---
 CMakeLists.txt          |  4 ++--
 src/backend.h           | 20 ++++++++++----------
 src/cuda/vla_cuda_ops.h |  2 +-
 3 files changed, 13 insertions(+), 13 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 481c699..5ab1dc9 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -286,8 +286,8 @@ if(GGML_CUDA)
     target_link_libraries(vla_cuda_ops PUBLIC CUDA::cublas CUDA::cudart ggml-cuda)
     target_link_libraries(vla_core PRIVATE vla_cuda_ops)
 elseif(GGML_HIP)
-    # ggml-hip exports GGML_USE_CUDA publicly; keep vla.cpp's own CUDA code out.
-    target_compile_definitions(vla_core PUBLIC VLA_USE_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)
diff --git a/src/backend.h b/src/backend.h
index 87b0ba0..7d0eb0c 100644
--- a/src/backend.h
+++ b/src/backend.h
@@ -269,28 +269,28 @@ inline int backend_device_index() {
 inline Backend backend_init(const char * tag, int n_threads) {
     Backend b;
 
-#if defined(VLA_USE_HIP)
+#if defined(GGML_USE_CUDA) && !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);
+            b.is_cuda = true;
+            std::printf("%s: backend = CUDA (device %d)\n", tag, dev);
         } else {
-            std::fprintf(stderr, "%s: ggml_backend_cuda_init failed on ROCm; falling back to CPU\n", tag);
+            std::fprintf(stderr, "%s: ggml_backend_cuda_init failed; falling back to CPU\n", tag);
         }
     }
-#elif defined(GGML_USE_CUDA)
+#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) {
-            b.is_cuda = true;
-            std::printf("%s: backend = CUDA (device %d)\n", tag, dev);
+            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; falling back to CPU\n", tag);
+            std::fprintf(stderr, "%s: ggml_backend_cuda_init failed on ROCm; falling back to CPU\n", tag);
         }
     }
 #elif defined(GGML_USE_SYCL)
diff --git a/src/cuda/vla_cuda_ops.h b/src/cuda/vla_cuda_ops.h
index 6f15968..4db2393 100644
--- a/src/cuda/vla_cuda_ops.h
+++ b/src/cuda/vla_cuda_ops.h
@@ -21,7 +21,7 @@
 
 namespace vla {
 
-#if defined(GGML_USE_CUDA) && !defined(VLA_USE_HIP)
+#if defined(GGML_USE_CUDA) && !defined(GGML_USE_HIP)
 void cuda_register_bf16_ops();
 #else
 inline void cuda_register_bf16_ops() {}

From 002a41c24fbf4c35538b4175e754bf730c51d94a Mon Sep 17 00:00:00 2001
From: singlecatlmx <2797705140@qq.com>
Date: Wed, 7 Oct 2026 18:15:10 +0800
Subject: [PATCH 4/4] Refresh ROCm validation after upstream merge

Document CPU/HIP tests and two-model fixed-input numerics on main 1adf078. Both Release builds pass 10/10 CTest on gfx1151 with ROCm 7.14.60850 and llama.cpp b11223. CPU actions match upstream main byte for byte; HIP actions are finite and repeat across five processes. HIP-vs-CPU max abs errors remain 1.992e-3 for SmolVLA and 8.792e-4 for pi05. Fresh 3x20 synthetic benchmarks report median P50 395.0/389.1 ms and worst P90 396.9/403.1 ms. Mark FoldQuant quantized GGUFs as unmeasured on HIP.
---
 docs/backend/rocm.md | 31 ++++++++++++++++---------------
 1 file changed, 16 insertions(+), 15 deletions(-)

diff --git a/docs/backend/rocm.md b/docs/backend/rocm.md
index f987e9c..81f4f50 100644
--- a/docs/backend/rocm.md
+++ b/docs/backend/rocm.md
@@ -5,10 +5,10 @@
 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 Linux 6.17.0-40 on a 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.
+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
 
@@ -59,9 +59,9 @@ CUDA or Vulkan, and rejects `GGML_BACKEND_DL=ON` with HIP.
 `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@8eec763` for both rows below. Each HIP row returned the complete
+`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 eight
+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 |
@@ -97,22 +97,21 @@ portion of each output. Use five fresh HIP processes to check repeatability.
 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 shown in the before/after manifests), and
-CPU governor `performance`. The GPU process probe saw no running process
-before or after the rounds. The stock `vla-bench` reports P50, P90 and mean
+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 | 397.5 | 210.2 |
-| π0.5 LIBERO | 2 × 224 | 128 | 387.8 | 403.8 | 46.8 |
+| 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 386.3, 390.4 and 387.8 ms; the worst P90 was
-403.8 ms. These results use upstream v0.4.0's model code and llama.cpp
-`b11223`. Do not compare them as a speedup against an older revision without
-a paired benchmark.
+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:
@@ -138,6 +137,8 @@ manifest when presenting new results.
   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.