diff --git a/CHANGELOG.md b/CHANGELOG.md index 07b6029..1c3da0e 100644 --- a/CHANGELOG.md +++ b/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, diff --git a/CMakeLists.txt b/CMakeLists.txt index 2a01c88..4c5f2d7 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() @@ -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) diff --git a/README.md b/README.md index 2f73c8e..6b08ad3 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. --- @@ -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 | diff --git a/docs/ARCHITECTURE.md b/docs/ARCHITECTURE.md index 7f7939b..e7fed74 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..81f4f50 --- /dev/null +++ b/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=` 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. diff --git a/src/backend.h b/src/backend.h index c786426..a4caf1f 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 @@ -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 diff --git a/src/cuda/vla_cuda_ops.h b/src/cuda/vla_cuda_ops.h index b16c3b3..6d8e4a2 100644 --- a/src/cuda/vla_cuda_ops.h +++ b/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