diff --git a/CHANGELOG.md b/CHANGELOG.md index 8d2436c7..48f547e9 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -9,15 +9,35 @@ Notable changes to vla.cpp. Format loosely follows [Keep a Changelog](https://ke - **Hexagon NPU on the Qualcomm Dragonwing IQ-10 (Linux).** `scripts/build_linux_hexagon.sh` cross-builds vla.cpp with `GGML_HEXAGON=ON` in Qualcomm's toolchain container, `vla-server` included. On the IQ-10 (v81, - llama.cpp `b11223`) SmolVLA, π0.5, Evo-1 and GR00T N1.5 / N1.6 / N1.7 run - within 2.9e-3 of the CPU build on fixed inputs and serve through - `eval/client`. Concats the NPU refuses on buffers above 4 GB of DSP address - space and matmuls of a 2D weight against a batch, which it computes wrong - (GR00T N1.6 with two views), run on the CPU, and GR00T N1.7's LM leaves - flash attention (F16 scores on the HMX) unless `--flash-attn` is given. - GR00T N1.5 and N1.6 are above the bar on a served LIBERO observation, as the - CPU is at F16. See `docs/backend/hexagon.md` for the build, the skel install - and the numbers. + llama.cpp `b11223`) SmolVLA, π0, π0.5, Evo-1, GR00T N1.5 / N1.6 / N1.7, + VLA-JEPA, VLA-Adapter, OpenVLA-OFT, Octo-Small and TurboVLA run within 2.9e-3 + of the CPU build on fixed inputs and serve through `eval/client`. Concats the + NPU refuses on buffers above 4 GB of DSP address space, copies that read a + weight (also answered `NO_SUPPORT`; VLA-Adapter), matmuls of a 2D weight + against a batch, which it computes wrong, and repeats of a computed tensor, + which can read stale cache lines (TurboVLA), run on the CPU, and GR00T N1.7's + LM leaves flash attention (F16 scores on the HMX) unless `--flash-attn` is + given. Three matmuls that the HMX does not take move onto it: SigLIP's + 4304-wide MLP is zero-padded to 4320 at load, π0.5's SigLIP takes flash + attention, and Evo-1's LM attention takes F16 K/V. GR00T N1.5, N1.6, π0.5 and + Evo-1 run 27-40% faster, within the bar on fixed inputs. π0.5's VLM and action + expert take flash attention too, and the experts' KV concats run as row copies + that stay on the NPU (π0, π0.5, SmolVLA). VLA-Adapter's and OpenVLA-OFT's + DINOv2 + SigLIP towers take the same padding and flash attention, `linear()` + multiplies a batch of activations as one matrix of rows, and large transposing + copies run on the CPU: OpenVLA-OFT, VLA-Adapter and TurboVLA run 51-57% + faster, Evo-1 and SmolVLA 8-10%. VLA-JEPA's LM leaves flash attention unless + `--flash-attn` is given, which brings it within the bar on a served LIBERO + observation at 13% more latency. VLA-Adapter's LM and action head and Octo's + block transformer pad their softmax rows to a multiple of 32, and + VLA-Adapter's and OpenVLA-OFT's DINOv2 build GELU (erf) from ops the NPU runs: + VLA-Adapter 17% faster, Octo 4%. ACT runs within 6.1e-4 of the CPU F32 + reference (1.0e-3 served): pads that read a weight, also answered + `NO_SUPPORT`, run on the CPU, and its ResNet's im2col matrices are + zero-padded to a multiple of 32 so the convs' matmuls take the HMX (152.9 ms, + 170.8 ms on the CPU). GR00T N1.5 and N1.6 are above the bar on a served LIBERO + observation, as the CPU is at F16. See + `docs/backend/hexagon.md` for the build, the skel install and the numbers. - **ROCm/HIP backend for VLA inference.** `-DGGML_HIP=ON` selects ggml's HIP backend without compiling vla.cpp's CUDA-only kernels. On Linux `gfx1200` (Radeon RX 9060 XT) with ROCm 7.14 and llama.cpp `b11223`, every arch except diff --git a/README.md b/README.md index 0cacbd2e..bdd9e87b 100644 --- a/README.md +++ b/README.md @@ -50,7 +50,7 @@ cmake --build build -j$(nproc) For a GPU or NPU, follow the backend's page in [docs/backend/](docs/backend): [CUDA](docs/backend/cuda.md), [ROCm](docs/backend/rocm.md), [SYCL](docs/backend/sycl.md), [OpenVINO](docs/backend/ov.md), -[Metal](docs/backend/metal.md), [Hexagon](docs/backend/hexagon-windows.md). +[Metal](docs/backend/metal.md), [Hexagon](docs/backend/hexagon.md). For containers, see [docs/DOCKER.md](docs/DOCKER.md). --- @@ -86,10 +86,10 @@ Per-device latency and memory are in [docs/benchmark/](docs/benchmark). Use [ggml-compiler](https://github.com/khanhnd61-vr/ggml-compiler) and follow [the instruction](docs/ARCHITECTURE.md#adding-an-architecture) to add a new model. -| Model | CPU (x86-64 / ARM) | [CUDA (NVIDIA)](docs/backend/cuda.md) | [ROCm (AMD)](docs/backend/rocm.md) | [SYCL (Intel)](docs/backend/sycl.md) | [Metal (Apple)](docs/backend/metal.md) | [OpenVINO (Intel)](docs/backend/ov.md) | [Hexagon (Qualcomm)](docs/backend/hexagon-windows.md) | +| Model | CPU (x86-64 / ARM) | [CUDA (NVIDIA)](docs/backend/cuda.md) | [ROCm (AMD)](docs/backend/rocm.md) | [SYCL (Intel)](docs/backend/sycl.md) | [Metal (Apple)](docs/backend/metal.md) | [OpenVINO (Intel)](docs/backend/ov.md) | [Hexagon (Qualcomm)](docs/backend/hexagon.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 | Y | Y | ~ | +| [π0](https://hf.co/vrfai/pi0-libero-finetuned-v044-gguf) | Y | Y | Y | Y | Y | Y | Y | | [π0.5](https://hf.co/vrfai/pi05-libero-gguf) | Y | Y | Y | Y | Y | Y | Y | | [GR00T N1.5](https://hf.co/vrfai/gr00tn1d5-libero-object-gguf) | Y | Y | Y | Y | Y | Y | Y | | [GR00T N1.6](https://hf.co/vrfai/gr00tn1d6-libero-gguf) | Y | Y | Y | Y | Y | Y | Y | @@ -97,14 +97,20 @@ Use [ggml-compiler](https://github.com/khanhnd61-vr/ggml-compiler) and follow [t | [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 | Y | | [VLA-Adapter](https://hf.co/vrfai/vla-adapter-libero-gguf) | Y | Y | Y | Y | Y | Y | Y | -| [OpenVLA-OFT](https://hf.co/vrfai/openvla-oft-libero-gguf) | Y | Y | Y | Y | Y | Y | - | -| [VLA-JEPA](https://hf.co/vrfai/vla-jepa-libero) | Y | Y | Y | Y | Y | Y | ~ | +| [OpenVLA-OFT](https://hf.co/vrfai/openvla-oft-libero-gguf) | Y | Y | Y | Y | Y | Y | Y | +| [VLA-JEPA](https://hf.co/vrfai/vla-jepa-libero) | Y | Y | Y | Y | Y | Y | Y | | [Octo-Small](https://hf.co/vrfai/octo-small-libero-gguf) | Y | Y | Y | Y | Y | - | Y | | [TurboVLA](https://hf.co/vrfai/turbovla-libero-gguf) | Y | Y | Y | Y | Y | Y | Y | -| [ACT](https://hf.co/ravediamond/act-aloha-sim-transfer-cube-gguf) | Y | Y | Y | ~ | Y | ~ | ~ | +| [ACT](https://hf.co/ravediamond/act-aloha-sim-transfer-cube-gguf) | Y | Y | Y | ~ | Y | ~ | Y | Legend: `Y` = supported (released and benchmarked), `~` = in progress, `-` = planned. +The Hexagon column is the Dragonwing IQ-10 under Linux +([docs/backend/hexagon.md](docs/backend/hexagon.md)). On a Snapdragon X laptop +under Windows, π0 and VLA-JEPA land above the 2.9e-3 fidelity bar, OpenVLA-OFT +was not attempted (about 14 GB at F16, over its 8 GB budget) and ACT was not run +([docs/backend/hexagon-windows.md](docs/backend/hexagon-windows.md)). + --- ## Rollout on a real robot diff --git a/docs/backend/hexagon.md b/docs/backend/hexagon.md index 6df93b30..14a4abd8 100644 --- a/docs/backend/hexagon.md +++ b/docs/backend/hexagon.md @@ -5,9 +5,9 @@ to take once it ran. > **Status: running on Windows on Arm and on a Dragonwing IQ-10 under Linux.** > vla.cpp builds with `GGML_HEXAGON=ON` and runs eleven checkpoints on a -> Snapdragon X laptop's v73 NPU and six on an IQ-10's v81 NPU under Linux, each -> checked against a CPU reference. The laptop build, its measurements and the -> upstream bugs found along the way are in +> Snapdragon X laptop's v73 NPU and thirteen on an IQ-10's v81 NPU under Linux, +> each checked against a CPU reference. The laptop build, its measurements and +> the upstream bugs found along the way are in > [hexagon-windows.md](hexagon-windows.md); the Linux board is in [Linux > (IQ-10)](#linux-iq-10). What follows is the backend in general: the parts that > hold on any Snapdragon part, and the IQ-9 and multi-NPU notes from before @@ -94,7 +94,14 @@ changed all three conclusions: it. - **`GELU_ERF` is still missing.** It now runs on the CPU (next point) in every DINOv2 and SigLIP-so400m tower, and costs what that implies: Evo-1 spends more - on copies than the NPU saves. + on copies than the NPU saves. An HTP kernel was written and checked on the + IQ-10 (b11223): more accurate than the CPU's (relative error 1.5e-6 against + 5.1e-5), and Evo-1 7% faster. It was left out: GR00T N1.7 sits close to the + bar, and the kernel moved its fixed-input chunk from 2.0e-3 to 3.0e-3. + VLA-Adapter's and OpenVLA-OFT's DINOv2 MLP and projector build it from ops the + NPU runs instead (`vla::gelu_erf()`: x*sigmoid(p(x)), p an odd cubic in x + clamped to [-5, 5], within 2.5e-5 of the exact GELU): VLA-Adapter 740 to + 705 ms, OpenVLA-OFT 1771 to 1740 ms. - **Absent is no longer fatal.** `src/backend_fallback.cpp` wraps the NPU in a backend that runs whatever it rejects on the CPU, without touching an arch. That also covers ops refused for shape rather than type: softmax rows that are @@ -109,6 +116,7 @@ these, and they compute something else: | `ADD`/`MUL`/... | `src1` broadcast over one dim but not a higher one | Refused, runs on the CPU | | `IM2COL` | Wrong unless kernel = stride > 1 and no padding | Refused, runs on the CPU | | `CPY` F32→F16 | Reads back wrong on the host | Copies feeding a CPU op run on the CPU | +| `REPEAT` | Reads stale cache lines of a source written by DMA earlier in the same op batch | Refused unless the source is a weight, runs on the CPU | | fused `MUL_MAT+ADD`, `RMS_NORM+MUL` | ~10x the error of the unfused ops | `GGML_HEXAGON_OPFUSION=0` by default | Each one flips the gripper channel, or comes close, on some checkpoint. The @@ -117,21 +125,91 @@ was right: a missing op fails at load, and these only show in the actions. [hexagon-windows.md](hexagon-windows.md#hexagon-issues-found-and-worked-around) has the numbers and how each was bisected. They belong upstream. -One more op is claimed and then refused at run time. On buffers mapped above +Two more ops are claimed and then refused at run time. On buffers mapped above 4 GB of DSP address space, which the IQ-10 under Linux hands out, the NPU runs `CONCAT` only as the dim-1 contiguous DMA copy and answers any other concat with `NO_SUPPORT`, and the host aborts on the first predict. The wrapper sends those -concats to the CPU. +concats to the CPU. A `CPY` or a `PAD` that reads a weight gets the same answer +(VLA-Adapter casts its F16 action queries to F32 in the graph, ACT zero-pads its +stem kernel), and also runs on the CPU. + +The experts' KV concats (π0, π0.5, SmolVLA) were such concats: along the token +dim of contiguous [head dim, heads, tokens] tensors, once per layer per denoise +step. Appending along that dim appends whole rows, so on Hexagon +`concat_dim2()` in `src/backend.h` does them as a dim-1 concat of the +[head dim x heads, tokens] views, which writes the same bytes and stays on the +NPU: π0.5 1379 to 1133 ms, SmolVLA 810 to 638 ms, outputs unchanged. π0 takes +the same path. Flash attention on the HMX (v81) keeps scores and probabilities in F16. That is a precision limit rather than a wrong op, but in GR00T N1.7's LM it moves the action chunk 5.4e-3 from the CPU reference. N1.7's LM therefore takes the explicit path unless `--flash-attn` is given: 2.0e-3, at the same latency. +VLA-JEPA's LM does the same: with flash attention a served LIBERO observation +lands 6.7e-3 from the reference (the CPU with flash attention 5.5e-3), without +it 2.6e-3, at a cost of 683 to 810 ms with two views. + +The softmax takes rows of up to 32 elements or a multiple of 32. Where a +sequence can grow without changing its result, vla.cpp pads it to +`gemm_width()` on Hexagon rather than send every layer's softmax to the CPU: +VLA-Adapter's causal LM and Octo's block transformer get zero tokens appended, +masked out as keys, and each VLA-Adapter action-head block stacks its three +key/value sets as rows, padded and masked, instead of concatenating scores +along their columns. VLA-Adapter goes from 859 to 747 ms, Octo-Small from 126.5 +to 122.6 ms. OpenVLA-OFT's LM gains nothing that way: its 22 padding tokens add +as much 7B matmul as the softmax saves. `MUL_MAT` with a 2D weight against a batch of activations (weight `ne[2..3]` of 1, activations above 1) is claimed and comes out wrong. GR00T N1.6's SigLIP2 tower takes two camera views that way: its action chunk lands 0.16 from the CPU -reference. The wrapper runs those matmuls on the CPU: 1.9e-3, and 25% faster. +reference. The wrapper runs those matmuls on the CPU. On Hexagon `linear()` in +`src/layers/linear.h` multiplies a contiguous batch as one 2D matrix of rows +instead, the same product, so most of them never reach the rule: GR00T N1.6 +with two views lands 8.8e-4 from the reference, and the matmuls of TurboVLA's +DINOv3, which takes its two views as a batch, run on the NPU. + +A copy that transposes the innermost dim, such as the `CONT` of a permuted V +for explicit attention, is a strided gather on the NPU: 48 ms for OpenVLA-OFT's +[330,128,32], 32 ms for Evo-1's [1025,112,8], 18 ms for TurboVLA's +[512,256,4]. The wrapper runs those of 256K elements or more on the CPU, which +writes the same bytes; the smaller ones take tens of microseconds on the NPU, +less than the trip to the CPU. + +The matrix unit (HMX) takes a matmul only when its inner dim K and its output +dim N are multiples of 32, and with a batch only an F16 src0. Anything else runs +on the vector units, at about 1/20 of the speed. `GGML_HEXAGON_PROFILE=1` on the +IQ-10 showed such matmuls dominating, and vla.cpp now avoids them in +Hexagon builds: + +- **SigLIP-so400m's MLP is 4304 wide**: 440 ms of GR00T N1.5's 565 ms tower. + The loader zero-pads fc1, its bias and fc2 to 4320 + (`WeightLoader::pad_mlp`, `gemm_width()` in `src/backend.h`), which leaves + the output unchanged. VLA-Adapter's and OpenVLA-OFT's SigLIP + (`src/modules/dual_tower.h`) take the same padding. +- **π0.5's attention ran explicit scores** in SigLIP (head dim 72, F32 x F32), + the VLM and the action expert, because its flash attention was tied to SYCL's + view-taking kernel. The VLM's and the expert's single KV head also sent their + matmuls to the CPU (the batched `MUL_MAT` rule above). All three now take the + plain flash attention on Hexagon (`use_flash_attention()` in `pi05.cpp`). + VLA-Adapter's and OpenVLA-OFT's DINOv2 and SigLIP towers did the same and now + take flash attention on Hexagon too: their vision stage goes from 790 to + 311 ms. +- **Evo-1's LM attention multiplied F32 K and V against 14 query heads.** K and + V now go to F16 (`fa_kv()`); the scores still accumulate in F32. +- **ACT's ResNet convs are im2col and one matmul** with the patches as its + first operand. The stem's K of 147 and layer3's and layer4's 1200 and 300 + output positions (640x480) are not multiples of 32: 80 ms of the 101 ms spent + on the NPU. On Hexagon `conv_2d_padded()` in `act.cpp` zero-pads the patch + matrix and the stem's kernel to `gemm_width()` and drops the extra rows: + 228.5 to 155.6 ms (mean, 5 reps). Its im2col (padded and overlapping), bias + adds (a gapped broadcast) and max pool still run on the CPU; under the + benchmark's protocol ACT takes 152.9 ms on the NPU against 170.8 ms on the + CPU. + +What is left on the vector units: GR00T's DiT attention (head dim 48) and Evo-1's +DiT cross-attention (head dim 112, 1025 keys), whose shapes the HMX does not +take, and Evo-1's InternViT attention, which stays on the CPU because F16 K/V +there measured 3.9e-2. ## What changed in `vla.cpp` @@ -280,43 +358,62 @@ export LD_LIBRARY_PATH=$PWD/pkg/lib `vla_predict_check` at the Hexagon defaults, against the same binary with `VLA_DEVICE=cpu --weight-dtype f32 --flash-attn 0`. A Hexagon build turns flash attention on by default even on the CPU, so the reference asks for the explicit -path that other builds run: +path that other builds run. Views are the ones the LIBERO client sends, so they +differ from the benchmark's for VLA-JEPA, VLA-Adapter, OpenVLA-OFT and Octo: -| Model | max\|Δ\| | -|---|--:| -| SmolVLA | 7.2e-4 | -| π0.5 | 2.7e-4 | -| Evo-1 | 4.1e-4 | -| GR00T N1.5 | 1.7e-3 | -| GR00T N1.6 | 1.4e-3 | -| GR00T N1.7 | 2.0e-3 | - -GR00T N1.6 here takes one view; with two it is 1.9e-3. The other archs have -not been run on the board yet. +| Model | Views | max\|Δ\| | +|---|--:|--:| +| SmolVLA | 2 | 7.2e-4 | +| π0 | 2 | 1.2e-3 | +| π0.5 | 2 | 4.8e-4 | +| Evo-1 | 1 | 5.5e-4 | +| GR00T N1.5 | 1 | 2.4e-3 | +| GR00T N1.6 | 1 | 1.4e-3 | +| GR00T N1.7 | 1 | 2.0e-3 | +| VLA-JEPA | 2 | 2.2e-3 | +| VLA-Adapter | 2 | 5.4e-4 | +| OpenVLA-OFT | 2 | 3.2e-4 | +| Octo-Small | 1 | 5.3e-4 | +| TurboVLA | 2 | 3.7e-4 | +| ACT | 1 | 6.1e-4 | + +GR00T N1.6 with two views is 8.8e-4. At the benchmark's single view VLA-JEPA is +1.9e-3, VLA-Adapter 5.7e-4 and OpenVLA-OFT 2.1e-4. ACT takes its one 640x480 +view, the benchmark's input. ### Serving `vla-server` on the board, driven from the host by `eval/client/vla_cpp_client.py` and the adapters `run_sim_client_direct.py` -uses, on one fixed LIBERO observation (two views, the task prompt, a fixed -robot state) with fixed noise. The same request goes to a Hexagon server and to -a CPU server with `--weight-dtype f32 --flash-attn 0`. Every model returns a -finite chunk, the same chunk for a repeated request, an error reply to a -malformed or an invalid request, and keeps serving afterwards. +uses, on one fixed LIBERO observation (two views, the task prompt, a fixed robot +state), with fixed noise for the archs that draw it. The same request goes to a +Hexagon server and to a CPU server with `--weight-dtype f32 --flash-attn 0`. +Every model returns a finite chunk, the same chunk for a repeated request, an +error reply to a malformed or an invalid request, and keeps serving afterwards. +ACT has no `eval/client` preset; it gets the same checks from raw `vla.proto` +requests carrying one fixed 640x480 view and a fixed 14-dim state. | Model | max\|Δ\| vs CPU F32 | CPU F16 vs CPU F32 | |---|--:|--:| | SmolVLA | 1.3e-3 | | -| π0.5 | 4.3e-4 | | -| Evo-1 | 4.3e-4 | | -| GR00T N1.5 | 8.3e-3 | 4.3e-2 | +| π0.5 | 3.5e-4 | | +| Evo-1 | 1.1e-3 | | +| GR00T N1.5 | 1.1e-2 | 4.3e-2 | | GR00T N1.6 | 1.3e-2 | 1.1e-2 | | GR00T N1.7 | 2.5e-3 | | - -GR00T N1.5 and N1.6 are above the bar on this input, by about what the CPU -itself gives with F16 weights: they are sensitive to 16-bit precision rather -than miscomputed. Running them with F32 weights on the NPU has not been tried -yet. +| π0 | 5.8e-4 | | +| VLA-JEPA | 2.6e-3 | 2.1e-2 | +| VLA-Adapter | 5.4e-4 | | +| OpenVLA-OFT | 5.9e-4 | | +| Octo-Small | 2.4e-4 | | +| TurboVLA | 8.5e-4 | | +| ACT | 1.0e-3 | | + +GR00T N1.5 and N1.6 are above the bar on this input, and so is the CPU itself +with F16 weights: N1.5 lands closer to the reference than the CPU at F16, N1.6 +slightly further (1.3e-2 against 1.1e-2). They are sensitive to 16-bit precision +rather than miscomputed. Running them with F32 weights on the NPU has not been +tried yet. ## IQ-10, imagined diff --git a/docs/benchmark/README.md b/docs/benchmark/README.md index 7c381c8f..50c184ea 100644 --- a/docs/benchmark/README.md +++ b/docs/benchmark/README.md @@ -12,7 +12,7 @@ AGX Orin row ran at `e6653e8`, on the bit-exact CUDA path ran at `d0e5d00`, PR #35 merged with main `f464653` (same llama.cpp pin), with fixed-input checks against the CPU build in [backend/rocm.md](../backend/rocm.md). The Qualcomm Dragonwing IQ-10 row ran at -`e8dfa34` (PR #69, same llama.cpp pin). Each report +`7c76ea9`, its ACT cell at `0f3e010` (same llama.cpp pin). Each report lists the device, the build, each model's configuration, latency and peak memory at the defaults, and the fastest runtime flags for that device. @@ -21,33 +21,35 @@ every report has. `N/A` means the arch does not run on that backend, and `-` means the model does not fit or fails there; each report says which and why. `n/r` means not run. -| Device | Octo-Small | TurboVLA | VLA-JEPA | VLA-Adapter | GR00T N1.5 | GR00T N1.6 | GR00T N1.7 | BitVLA | SmolVLA | π0 | π0.5 | OpenVLA-OFT | Evo-1 | -|---|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:| -| [RTX 5090 (CUDA), author's numbers](rtx-5090.md) | 2.69 | 4.88 | 14.3 | 16.0 | 16.1 | 19.4 | 19.5 | 20.8 | 38.2 | 28.8 | 29.2 | 34.5 | 49.1 | -| [RTX 3090 (CUDA)](rtx-3090.md) | 4.3 | 9.0 | 27.3 | 34.2 | 36.6 | 36.3 | 35.3 | 26.9 | 61.0 | 97.5 | 97.8 | 134.2 | 108.6 | -| [RTX 3060 (CUDA)](rtx-3060.md) | 8.3 | 21.6 | 57.1 | 74.7 | 86.8 | 78.1 | 82.8 | 59.1 | 125.1 | 228.2 | 229.9 | - | 230.8 | -| [RTX 5070 Laptop (CUDA)](rtx-5070-laptop.md) | 4.7 | 14.1 | 38.0 | 56.3 | 73.0 | 65.5 | 63.3 | 54.4 | 89.3 | 177.1 | 180.5 | - | 184.6 | -| [Radeon RX 9060 XT (ROCm)](rx-9060-xt.md) | 14.0 | 33.8 | 65.9 | 70.1 | 79.5 | 94.6 | 97.1 | N/A | 207.7 | 217.9 | 220.6 | 182.2 | 233.6 | -| [Jetson AGX Orin (CUDA)](jetson-agx-orin.md) | 14.3 | 28.0 | 60.6 | 79.1 | 96.7 | 101.2 | 93.0 | 76.8 | 164.3 | 269.0 | 266.6 | 297.4 | 298.5 | -| [Jetson Orin NX 16 GB (CUDA)](jetson-orin-nx.md) | 27.4 | 79.7 | 176.7 | 270.7 | 260.8 | 252.1 | 251.4 | 269.1 | 411.2 | 739.6 | 742.9 | - | 893.2 | -| [Jetson Orin Nano Super (CUDA)](jetson-orin-nano.md) | 31.9 | 89.3 | 193.9 | 297.5 | 293.4 | 272.3 | 271.7 | 335.0 | 462.8 | 845.2 | 849.7 | - | 1011.5 | -| [Apple M4 (Metal)](apple-m4.md) | 22.6 | 65.2 | 248.7 | 315.7 | 401.5 | 358.9 | 350.8 | N/A | 357.6 | 1141.1 | 1153.8 | 1650.7 | 829.3 | -| [Intel Arc B390 iGPU (SYCL)](core-ultra-x7-358h.md) | 15.5 | 25.0 | 102.1 | 62.6 | 101.2 | 106.1 | 178.4 | N/A | 61.7 | 269.3 | 275.2 | 434.9 | 270.4 | -| [Intel Arc B390 iGPU (OpenVINO)](core-ultra-x7-358h.md) | N/A | 28.5 | 116.3 | 170.2 | 106.0 | 122.8 | 131.4 | N/A | 213.1 | 620.5 | 232.5 | 1262.5 | 566.0 | -| [Intel Arc A380 (SYCL)](arc-a380.md) | 60.5 | 103.5 | 355.9 | 408.9 | 472.0 | 646.0 | 615.3 | N/A | 734.5 | 921.9 | 945.6 | - | 1047.3 | -| [Snapdragon X Hexagon NPU](snapdragon-x-hexagon.md) | 239.7 | 2280.8 | 821.0 | 4072.1 | 2755.7 | 2940.0 | 1400.2 | N/A | 1239.3 | 5730.0 | 9321.7 | - | 7208.6 | -| [Qualcomm Dragonwing IQ-10 Hexagon NPU](qualcomm-iq10.md) | n/r | n/r | n/r | n/r | 1020.9 | 1110.3 | 747.5 | N/A | 810.3 | n/r | 2914.6 | n/r | 3519.4 | -| [Intel AI Boost NPU (OpenVINO)](core-ultra-x7-358h.md) | N/A | 23.5 | 90.7 | 126.7 | 134.0 | 161.1 | 144.9 | N/A | 161.3 | 438.3 | 430.2 | 405.8 | 256.3 | -| [Intel Core i7-14700F (CPU)](core-i7-14700f.md) | 39.8 | 195.1 | 825.5 | 1140.6 | 1913.8 | 1627.9 | 1453.4 | N/A | 1689.2 | 6455.7 | 6068.1 | 9869.0 | 4159.6 | -| [Intel Core i9-14900HX (CPU)](core-i9-14900hx.md) | 67.0 | 314.7 | 1227.2 | 1659.3 | 1935.4 | 1688.9 | 1505.5 | N/A | 1822.6 | 6182.1 | 5735.0 | 9126.4 | 4118.2 | -| [Intel Core Ultra X7 358H (CPU)](core-ultra-x7-358h.md) | 48.0 | 199.0 | 1258.1 | 1645.4 | 1979.1 | 1697.3 | 1532.5 | N/A | 1758.0 | 6083.9 | 6120.3 | 9894.3 | 4025.8 | -| [Intel Core Ultra X7 358H (OpenVINO CPU)](core-ultra-x7-358h.md) | N/A | 238.4 | 1407.2 | 1831.2 | 2456.6 | 2300.4 | 2076.5 | N/A | 1899.5 | 7537.1 | 7522.7 | 11651.3 | 4626.5 | -| [Intel Core i5-12400F (CPU)](core-i5-12400f.md) | 78.1 | 392.7 | 1668.5 | 2165.7 | 2681.4 | 2216.8 | 1973.6 | N/A | 2288.4 | 8953.5 | 8637.9 | - | 5435.0 | -| [AMD Ryzen 5 5500 (CPU)](ryzen-5-5500.md) | 96.6 | 475.7 | 2133.0 | 2781.7 | 3447.0 | 2826.0 | 2518.8 | N/A | 3124.3 | 11321.3 | 11225.6 | - | 7347.2 | -| [Snapdragon X Oryon (CPU)](snapdragon-x-hexagon.md) | 109.5 | 548.9 | 1656.5 | 2013.1 | 2736.6 | 2226.0 | 1988.1 | N/A | 2080.9 | 9224.5 | 9140.6 | - | 5518.9 | +| Device | Octo-Small | ACT | TurboVLA | VLA-JEPA | VLA-Adapter | GR00T N1.5 | GR00T N1.6 | GR00T N1.7 | BitVLA | SmolVLA | π0 | π0.5 | OpenVLA-OFT | Evo-1 | +|---|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:|--:| +| [RTX 5090 (CUDA), author's numbers](rtx-5090.md) | 2.69 | n/r | 4.88 | 14.3 | 16.0 | 16.1 | 19.4 | 19.5 | 20.8 | 38.2 | 28.8 | 29.2 | 34.5 | 49.1 | +| [RTX 3090 (CUDA)](rtx-3090.md) | 4.3 | n/r | 9.0 | 27.3 | 34.2 | 36.6 | 36.3 | 35.3 | 26.9 | 61.0 | 97.5 | 97.8 | 134.2 | 108.6 | +| [RTX 3060 (CUDA)](rtx-3060.md) | 8.3 | n/r | 21.6 | 57.1 | 74.7 | 86.8 | 78.1 | 82.8 | 59.1 | 125.1 | 228.2 | 229.9 | - | 230.8 | +| [RTX 5070 Laptop (CUDA)](rtx-5070-laptop.md) | 4.7 | n/r | 14.1 | 38.0 | 56.3 | 73.0 | 65.5 | 63.3 | 54.4 | 89.3 | 177.1 | 180.5 | - | 184.6 | +| [Radeon RX 9060 XT (ROCm)](rx-9060-xt.md) | 14.0 | n/r | 33.8 | 65.9 | 70.1 | 79.5 | 94.6 | 97.1 | N/A | 207.7 | 217.9 | 220.6 | 182.2 | 233.6 | +| [Jetson AGX Orin (CUDA)](jetson-agx-orin.md) | 14.3 | n/r | 28.0 | 60.6 | 79.1 | 96.7 | 101.2 | 93.0 | 76.8 | 164.3 | 269.0 | 266.6 | 297.4 | 298.5 | +| [Jetson Orin NX 16 GB (CUDA)](jetson-orin-nx.md) | 27.4 | n/r | 79.7 | 176.7 | 270.7 | 260.8 | 252.1 | 251.4 | 269.1 | 411.2 | 739.6 | 742.9 | - | 893.2 | +| [Jetson Orin Nano Super (CUDA)](jetson-orin-nano.md) | 31.9 | n/r | 89.3 | 193.9 | 297.5 | 293.4 | 272.3 | 271.7 | 335.0 | 462.8 | 845.2 | 849.7 | - | 1011.5 | +| [Apple M4 (Metal)](apple-m4.md) | 22.6 | n/r | 65.2 | 248.7 | 315.7 | 401.5 | 358.9 | 350.8 | N/A | 357.6 | 1141.1 | 1153.8 | 1650.7 | 829.3 | +| [Intel Arc B390 iGPU (SYCL)](core-ultra-x7-358h.md) | 15.5 | n/r | 25.0 | 102.1 | 62.6 | 101.2 | 106.1 | 178.4 | N/A | 61.7 | 269.3 | 275.2 | 434.9 | 270.4 | +| [Intel Arc B390 iGPU (OpenVINO)](core-ultra-x7-358h.md) | N/A | n/r | 28.5 | 116.3 | 170.2 | 106.0 | 122.8 | 131.4 | N/A | 213.1 | 620.5 | 232.5 | 1262.5 | 566.0 | +| [Intel Arc A380 (SYCL)](arc-a380.md) | 60.5 | n/r | 103.5 | 355.9 | 408.9 | 472.0 | 646.0 | 615.3 | N/A | 734.5 | 921.9 | 945.6 | - | 1047.3 | +| [Snapdragon X Hexagon NPU](snapdragon-x-hexagon.md) | 239.7 | n/r | 2280.8 | 821.0 | 4072.1 | 2755.7 | 2940.0 | 1400.2 | N/A | 1239.3 | 5730.0 | 9321.7 | - | 7208.6 | +| [Qualcomm Dragonwing IQ-10 Hexagon NPU](qualcomm-iq10.md) | 117.6 | 148.5 | 188.1 | 505.4 | 703.2 | 591.6 | 696.5 | 759.5 | N/A | 583.6 | 1114.0 | 1113.0 | 1720.3 | 2333.6 | +| [Intel AI Boost NPU (OpenVINO)](core-ultra-x7-358h.md) | N/A | n/r | 23.5 | 90.7 | 126.7 | 134.0 | 161.1 | 144.9 | N/A | 161.3 | 438.3 | 430.2 | 405.8 | 256.3 | +| [Intel Core i7-14700F (CPU)](core-i7-14700f.md) | 39.8 | n/r | 195.1 | 825.5 | 1140.6 | 1913.8 | 1627.9 | 1453.4 | N/A | 1689.2 | 6455.7 | 6068.1 | 9869.0 | 4159.6 | +| [Intel Core i9-14900HX (CPU)](core-i9-14900hx.md) | 67.0 | n/r | 314.7 | 1227.2 | 1659.3 | 1935.4 | 1688.9 | 1505.5 | N/A | 1822.6 | 6182.1 | 5735.0 | 9126.4 | 4118.2 | +| [Intel Core Ultra X7 358H (CPU)](core-ultra-x7-358h.md) | 48.0 | n/r | 199.0 | 1258.1 | 1645.4 | 1979.1 | 1697.3 | 1532.5 | N/A | 1758.0 | 6083.9 | 6120.3 | 9894.3 | 4025.8 | +| [Intel Core Ultra X7 358H (OpenVINO CPU)](core-ultra-x7-358h.md) | N/A | n/r | 238.4 | 1407.2 | 1831.2 | 2456.6 | 2300.4 | 2076.5 | N/A | 1899.5 | 7537.1 | 7522.7 | 11651.3 | 4626.5 | +| [Intel Core i5-12400F (CPU)](core-i5-12400f.md) | 78.1 | n/r | 392.7 | 1668.5 | 2165.7 | 2681.4 | 2216.8 | 1973.6 | N/A | 2288.4 | 8953.5 | 8637.9 | - | 5435.0 | +| [AMD Ryzen 5 5500 (CPU)](ryzen-5-5500.md) | 96.6 | n/r | 475.7 | 2133.0 | 2781.7 | 3447.0 | 2826.0 | 2518.8 | N/A | 3124.3 | 11321.3 | 11225.6 | - | 7347.2 | +| [Snapdragon X Oryon (CPU)](snapdragon-x-hexagon.md) | 109.5 | n/r | 548.9 | 1656.5 | 2013.1 | 2736.6 | 2226.0 | 1988.1 | N/A | 2080.9 | 9224.5 | 9140.6 | - | 5518.9 | Notes on the table: +- ACT runs its ALOHA checkpoint on one 640x480 view. Only the IQ-10 report + measures it; the other rows are `n/r`. - The RTX 5090 row comes from the PR #32 author and was not re-measured. Its numbers are a minimum over 5 alternating rounds, not the protocol the other reports use (3 warmups + 20 timed reps, best of 3 processes). diff --git a/docs/benchmark/qualcomm-iq10.md b/docs/benchmark/qualcomm-iq10.md index b511afd0..39d1ef14 100644 --- a/docs/benchmark/qualcomm-iq10.md +++ b/docs/benchmark/qualcomm-iq10.md @@ -15,9 +15,9 @@ through vla.cpp's fallback wrapper. | Power | CPU governor `performance` | | OS | Ubuntu 24.04.4 LTS, kernel 6.6.110 | | Toolchain | `ghcr.io/snapdragon-toolchain/arm64-linux:v0.7`: clang 21.1.8, Hexagon SDK 6.6.0.0 (tools 19.0.07) | -| Commit | `e8dfa34` (PR #69, on top of #68) | +| Commit | `7c76ea9`; ACT at `0f3e010` | | llama.cpp | `b11223` | -| Date | 2026-11-02 11:02 to 11:48 (UTC+07) | +| Date | 2026-11-05 14:26 to 14:57 (UTC+07); ACT 17:00 to 17:01 | | Build | `scripts/build_linux_hexagon.sh` (servers included); the HTP skel processed for the board and installed in `/vendor/dsp/cdsp0..3` | | Runtime flags | defaults; the fastest flags per model are in the second table | | Method | 3 warmups + 20 timed reps per process, 3 processes. The process with the lowest mean is reported, and memory is the peak over all three. | @@ -31,8 +31,9 @@ buffers. Nothing else ran on the board. ## Model configuration Each model runs at its native view count and input size, with weights as -shipped. Checkpoints are the `vrfai/*` GGUFs; where a repo ships several, we use -the `libero_object` variant. Settings not listed are the defaults. +shipped. Checkpoints are the `vrfai/*` GGUFs, and ACT's is +`ravediamond/act-aloha-sim-transfer-cube-gguf`; where a repo ships several, we +use the `libero_object` variant. Settings not listed are the defaults. | Model | Checkpoint | Views | Input | Lang tokens | Action chunk | Action dim | Action head | Steps | |---|---|--:|--:|--:|--:|--:|---|--:| @@ -49,6 +50,7 @@ the `libero_object` variant. Settings not listed are the defaults. | OpenVLA-OFT | `openvla-oft-libero-gguf` | 1 | 224 | 16 | 8 | 7 | L1 regression head | 1 | | Octo-Small | `octo-small-libero-gguf` | 2 | 256 + 128 | 16 | 4 | 7 | DDPM diffusion | 20 | | TurboVLA | `turbovla-libero-gguf` | 2 | 256 | 16 | 12 | 7 | single-pass decoder | 1 | +| ACT | `act-aloha-sim-transfer-cube-gguf` | 1 | 640 x 480 | - | 100 | 14 | transformer decoder | 1 | - **Lang tokens** is the synthetic prompt length. VLA-JEPA adds its 32 `` tokens with `--extra-token 151697 --extra-count 32`. @@ -57,6 +59,8 @@ the `libero_object` variant. Settings not listed are the defaults. `VLA_GR00T_EMBODIMENT=new_embodiment` for N1.5, `libero_panda` for N1.6 and `libero_sim` for N1.7. **Octo** needs `VLA_OCTO_UNNORM_DATASET=libero_object`. - **Octo** takes a 256 px primary view and a 128 px wrist view. +- **ACT** is an ALOHA checkpoint (sim transfer cube): no language, a 14-dim + bimanual state and action, run with `--images 1 --size 640 --height 480`. ## Latency and memory @@ -65,12 +69,19 @@ vision stage separately. | Model | Views | Input | min ms | mean ms | p50 ms | p90 ms | vision ms | Peak RSS MiB | System RAM MiB | |---|--:|--:|--:|--:|--:|--:|--:|--:|--:| -| GR00T N1.7 | 1 | 256 | 747.5 | 754.7 | 753.0 | 762.7 | 196.1 | 191 | 5083 | -| SmolVLA | 2 | 512 | 810.3 | 818.6 | 820.5 | 822.6 | 283.0 | 146 | 1187 | -| GR00T N1.5 | 1 | 224 | 1020.9 | 1024.2 | 1024.2 | 1027.0 | 563.2 | 162 | 3650 | -| GR00T N1.6 | 1 | 224 | 1110.3 | 1114.5 | 1113.6 | 1118.5 | 559.8 | 147 | 4687 | -| π0.5 | 2 | 224 | 2914.6 | 2942.6 | 2934.4 | 2967.2 | 1429.6 | 285 | 5517 | -| Evo-1 | 1 | 448 | 3519.4 | 3572.4 | 3576.7 | 3588.2 | 854.5 | 319 | 1579 | +| Octo-Small | 2 | 256 + 128 | 117.6 | 119.4 | 118.4 | 120.7 | 21.2 | 253 | 723 | +| ACT | 1 | 640 x 480 | 148.5 | 152.9 | 152.5 | 157.3 | - | 194 | 467 | +| TurboVLA | 2 | 256 | 188.1 | 190.8 | 188.6 | 194.9 | - | 249 | 1075 | +| VLA-JEPA | 1 | 256 | 505.4 | 512.6 | 510.1 | 522.2 | 196.4 | 184 | 3969 | +| SmolVLA | 2 | 512 | 583.6 | 584.4 | 584.4 | 585.1 | 229.9 | 145 | 1196 | +| GR00T N1.5 | 1 | 224 | 591.6 | 594.9 | 593.4 | 597.1 | 138.9 | 162 | 3639 | +| GR00T N1.6 | 1 | 224 | 696.5 | 699.3 | 699.2 | 702.2 | 145.1 | 147 | 4692 | +| VLA-Adapter | 1 | 224 | 703.2 | 708.4 | 708.2 | 715.3 | 258.1 | 872 | 3566 | +| GR00T N1.7 | 1 | 256 | 759.5 | 766.3 | 766.0 | 773.4 | 197.6 | 191 | 5113 | +| π0 | 2 | 224 | 1114.0 | 1116.6 | 1116.4 | 1118.2 | 277.1 | 284 | 5537 | +| π0.5 | 2 | 224 | 1113.0 | 1117.3 | 1116.6 | 1120.7 | 277.9 | 285 | 5826 | +| OpenVLA-OFT | 1 | 224 | 1720.3 | 1725.5 | 1724.7 | 1731.0 | 263.4 | 844 | 15451 | +| Evo-1 | 1 | 448 | 2333.6 | 2348.3 | 2349.9 | 2360.1 | 820.6 | 319 | 1757 | ### Fastest configuration @@ -90,27 +101,44 @@ replaces the defaults only if its mean is more than 3% lower, both in the screen and in the full re-run. Flash attention and the matmul precision can move the action chunk. Success -rates measured at the defaults therefore do not carry over to these rows. +rates measured at the defaults therefore do not carry over to these rows. For +VLA-JEPA and GR00T N1.7, `--flash-attn` puts their LM back on flash attention, +which the defaults leave off because it is outside the fidelity bar (VLA-JEPA +6.7e-3 on a served LIBERO observation, N1.7 5.4e-3 on fixed inputs). | Model | Fastest flags | min ms | mean ms | p50 ms | p90 ms | vision ms | Peak RSS MiB | System RAM MiB | vs defaults | |---|---|--:|--:|--:|--:|--:|--:|--:|--:| -| GR00T N1.7 | *(defaults)* | 747.5 | 754.7 | 753.0 | 762.7 | 196.1 | 191 | 5083 | - | -| SmolVLA | *(defaults)* | 810.3 | 818.6 | 820.5 | 822.6 | 283.0 | 146 | 1187 | - | -| GR00T N1.5 | *(defaults)* | 1020.9 | 1024.2 | 1024.2 | 1027.0 | 563.2 | 162 | 3650 | - | -| GR00T N1.6 | *(defaults)* | 1110.3 | 1114.5 | 1113.6 | 1118.5 | 559.8 | 147 | 4687 | - | -| π0.5 | *(defaults)* | 2914.6 | 2942.6 | 2934.4 | 2967.2 | 1429.6 | 285 | 5517 | - | -| Evo-1 | *(defaults)* | 3519.4 | 3572.4 | 3576.7 | 3588.2 | 854.5 | 319 | 1579 | - | +| Octo-Small | *(defaults)* | 117.6 | 119.4 | 118.4 | 120.7 | 21.2 | 253 | 723 | - | +| ACT | *(defaults)* | 148.5 | 152.9 | 152.5 | 157.3 | - | 194 | 467 | - | +| TurboVLA | *(defaults)* | 188.1 | 190.8 | 188.6 | 194.9 | - | 249 | 1075 | - | +| VLA-JEPA | `--flash-attn --mm-prec default` | 442.1 | 452.4 | 450.2 | 460.6 | 198.3 | 184 | 3985 | -11.7% | +| SmolVLA | *(defaults)* | 583.6 | 584.4 | 584.4 | 585.1 | 229.9 | 145 | 1196 | - | +| GR00T N1.5 | *(defaults)* | 591.6 | 594.9 | 593.4 | 597.1 | 138.9 | 162 | 3639 | - | +| GR00T N1.6 | *(defaults)* | 696.5 | 699.3 | 699.2 | 702.2 | 145.1 | 147 | 4692 | - | +| VLA-Adapter | *(defaults)* | 703.2 | 708.4 | 708.2 | 715.3 | 258.1 | 872 | 3566 | - | +| GR00T N1.7 | `--flash-attn --mm-prec default` | 727.8 | 734.1 | 733.2 | 742.4 | 196.2 | 191 | 5081 | -4.2% | +| π0 | *(defaults)* | 1114.0 | 1116.6 | 1116.4 | 1118.2 | 277.1 | 284 | 5537 | - | +| π0.5 | *(defaults)* | 1113.0 | 1117.3 | 1116.6 | 1120.7 | 277.9 | 285 | 5826 | - | +| OpenVLA-OFT | *(defaults)* | 1720.3 | 1725.5 | 1724.7 | 1731.0 | 263.4 | 844 | 15451 | - | +| Evo-1 | *(defaults)* | 2333.6 | 2348.3 | 2349.9 | 2360.1 | 820.6 | 319 | 1757 | - |
Screen means (ms) per flag set | Model | defaults | `--flash-attn` | `--mm-prec default` | `--flash-attn --mm-prec default` | `--flash-attn 0` | |---|--:|--:|--:|--:|--:| -| GR00T N1.7 | 768.7 | 748.5 | 761.3 | 744.5 | 1086.6 | -| SmolVLA | 823.1 | 818.9 | 819.5 | 818.5 | 3018.8 | -| GR00T N1.5 | 1026.1 | 1023.0 | 1025.3 | 1026.1 | 1501.2 | -| GR00T N1.6 | 1114.4 | 1115.1 | 1114.4 | 1113.3 | 1290.3 | -| π0.5 | 2939.0 | 2956.9 | 2959.0 | 2962.7 | 2959.6 | -| Evo-1 | 3583.7 | 3593.9 | 3591.8 | 3582.6 | 5828.5 | +| Octo-Small | 121.8 | 120.4 | 121.8 | 123.2 | 123.8 | +| ACT | 143.4 | 153.4 | 157.5 | 155.8 | 174.8 | +| TurboVLA | 196.5 | 189.3 | 192.3 | 190.9 | 257.4 | +| VLA-JEPA | 514.8 | 454.3 | 512.0 | 452.9 | 1662.4 | +| SmolVLA | 584.6 | 585.0 | 585.2 | 584.6 | 2147.2 | +| GR00T N1.5 | 594.7 | 592.6 | 600.1 | 601.0 | 877.2 | +| GR00T N1.6 | 709.6 | 714.5 | 711.1 | 713.1 | 875.3 | +| VLA-Adapter | 712.4 | 711.1 | 713.7 | 712.8 | 989.2 | +| GR00T N1.7 | 765.5 | 739.0 | 758.3 | 733.1 | 865.3 | +| π0 | 1131.6 | 1138.1 | 1118.0 | 1128.8 | 2051.6 | +| π0.5 | 1125.8 | 1109.1 | 1134.2 | 1134.6 | 2036.0 | +| OpenVLA-OFT | 1737.9 | 1722.4 | 1715.7 | 1718.5 | 2005.8 | +| Evo-1 | 2345.2 | 2352.4 | 2355.4 | 2361.4 | 3877.1 |
@@ -118,8 +146,6 @@ rates measured at the defaults therefore do not carry over to these rows. - **BitVLA**: N/A off CUDA. The only published GGUF is int2-packed, and only CUDA builds load it. -- **π0, VLA-JEPA, VLA-Adapter, OpenVLA-OFT, Octo-Small, TurboVLA**: not run in - this round; their GGUFs were not on the board. ## Notes @@ -128,12 +154,12 @@ The HTP skel has to be processed for this board and installed in [backend/hexagon.md](../backend/hexagon.md#linux-iq-10). GR00T N1.6 runs one view here, as the configuration above fixes. With two views -its SigLIP2 tower hits the batched `MUL_MAT` the wrapper sends to the CPU; see +one batched `MUL_MAT` of its SigLIP2 tower still runs on the CPU; see [backend/hexagon.md](../backend/hexagon.md#op-coverage-measured). These rows time the defaults; they do not vouch for accuracy. On a served LIBERO -observation, GR00T N1.5 and N1.6 land above the 2.9e-3 fidelity bar, about as -far off as the CPU at F16; see +observation, GR00T N1.5 and N1.6 land above the 2.9e-3 fidelity bar: N1.5 +closer to the reference than the CPU at F16, N1.6 slightly further; see [backend/hexagon.md](../backend/hexagon.md#serving). ## Reproducing diff --git a/src/backend.h b/src/backend.h index bd574c0d..60a51755 100644 --- a/src/backend.h +++ b/src/backend.h @@ -191,6 +191,44 @@ inline ggml_tensor * fa_kv([[maybe_unused]] ggml_context * ctx, ggml_tensor * t) #endif } +/** + * @brief `ggml_concat(a, b, 2)` of [d0, d1, n] tensors, as the compiled-in + * backend runs it. + * + * Hexagon runs a concat on the NPU only as a dim-1 copy of whole rows; any other + * goes to the CPU fallback (src/backend_fallback.cpp). For contiguous inputs a + * dim-2 concat appends whole d0*d1 rows, so it is done as a dim-1 concat of the + * [d0*d1, n] views, which writes the same bytes. Everywhere else, and for + * non-contiguous inputs, this is ggml_concat(a, b, 2). + */ +inline ggml_tensor * concat_dim2(ggml_context * ctx, ggml_tensor * a, ggml_tensor * b) { +#ifdef GGML_USE_HEXAGON + if (ggml_is_contiguous(a) && ggml_is_contiguous(b) && a->ne[3] == 1 && b->ne[3] == 1 && + a->ne[0] == b->ne[0] && a->ne[1] == b->ne[1]) { + ggml_tensor * rows = ggml_concat(ctx, ggml_reshape_2d(ctx, a, a->ne[0] * a->ne[1], a->ne[2]), + ggml_reshape_2d(ctx, b, b->ne[0] * b->ne[1], b->ne[2]), 1); + return ggml_reshape_3d(ctx, rows, a->ne[0], a->ne[1], a->ne[2] + b->ne[2]); + } +#endif + return ggml_concat(ctx, a, b, 2); +} + +/** + * @brief A GEMM's inner width as the compiled-in backend's matmul takes it. + * + * Hexagon's matrix unit (HMX) takes an F16 weight only when both of its dims + * are multiples of 32; anything else runs on the vector units. SigLIP's + * 4304-wide MLP spends 440 ms there per tower on the IQ-10, against about + * 0.5 ms a matmul on the HMX. Everywhere else this returns @p n. + */ +inline int64_t gemm_width(int64_t n) { +#ifdef GGML_USE_HEXAGON + return (n + 31) / 32 * 32; +#else + return n; +#endif +} + /** * @brief `ggml_gelu` (the tanh form), in ops the compiled-in backend gets right. * @@ -211,6 +249,26 @@ inline ggml_tensor * gelu(ggml_context * ctx, ggml_tensor * x) { #endif } +/** + * @brief `ggml_gelu_erf`, in ops the compiled-in backend runs. + * + * ggml-hexagon has no GELU_ERF, so every MLP that uses it round-trips its + * activations through the CPU. There it is x*sigmoid(p(x)), p an odd cubic in + * x clamped to [-5, 5] fitted to logit(Phi(x)): within 2.5e-5 of the exact + * GELU for every x. Elsewhere this is plain ggml_gelu_erf. + */ +inline ggml_tensor * gelu_erf(ggml_context * ctx, ggml_tensor * x) { +#ifdef GGML_USE_HEXAGON + ggml_tensor * xc = ggml_clamp(ctx, x, -5.0f, 5.0f); + ggml_tensor * x2 = ggml_mul(ctx, xc, xc); + ggml_tensor * p = ggml_mul(ctx, ggml_scale_bias(ctx, x2, -7.0303601711e-4f, 7.4011308986e-2f), x2); + p = ggml_mul(ctx, ggml_scale_bias(ctx, p, 1.0f, 1.5950157475f), xc); + return ggml_mul(ctx, x, ggml_sigmoid(ctx, p)); +#else + return ggml_gelu_erf(ctx, x); +#endif +} + inline ggml_tensor * geglu(ggml_context * ctx, ggml_tensor * gate, ggml_tensor * up) { #if !defined(GGML_USE_METAL) && !defined(GGML_USE_OPENVINO) && \ !defined(GGML_USE_HEXAGON) && !defined(GGML_USE_OPENCL) diff --git a/src/backend_fallback.cpp b/src/backend_fallback.cpp index 7e0d4035..a4eaed17 100644 --- a/src/backend_fallback.cpp +++ b/src/backend_fallback.cpp @@ -351,6 +351,39 @@ bool accel_runs(FallbackCtx * fc, const ggml_tensor * t, const char ** why) { *why = "miscomputes a 2D weight against a batch in"; return false; } + if (fc->hexagon && t->op == GGML_OP_CPY && is_weight(t->src[0])) { + // Claimed, but a copy that reads a weight is answered with NO_SUPPORT + // and the host aborts. VLA-Adapter casts its F16 action queries + // [896,64] to F32 in the graph; with the cast on the CPU it is within + // 4.2e-4 of the CPU F32 reference. + *why = "cannot run (NO_SUPPORT) a weight-reading"; + return false; + } + if (fc->hexagon && t->op == GGML_OP_PAD && is_weight(t->src[0])) { + // The same for a pad that reads a weight: ACT zero-pads its stem kernel + // [147,64] to K = 160 so the conv's matmul takes the HMX, and the host + // aborts with the pad on the NPU. + *why = "cannot run (NO_SUPPORT) a weight-reading"; + return false; + } + if (fc->hexagon && t->op == GGML_OP_REPEAT && !is_weight(t->src[0])) { + // The NPU's REPEAT reads its source through the data cache, which is + // invalidated only between op batches, while the op that wrote the + // source used DMA, so within one batch it can read stale lines. + // TurboVLA repeats its ViT prefix [768,5] over two views: max|delta| + // 0.19 and different on every run, 3.1e-4 with the repeat on the CPU. + *why = "can read stale data in a"; + return false; + } + if (fc->hexagon && t->op == GGML_OP_CONT && t->src[0]->nb[0] != ggml_type_size(t->src[0]->type) && + ggml_nelements(t) >= 256 * 1024) { + // Runs, but a large copy that transposes the innermost dim is a strided + // gather on the NPU: 18 ms for TurboVLA's fusion values [512,256,4], 32 ms + // for Evo-1's [1025,112,8], 48 ms for OpenVLA-OFT's [330,128,32]. Small + // ones take tens of microseconds there, less than a trip to the CPU. + *why = "is slow on a large transposing"; + return false; + } return ggml_backend_supports_op(fc->accel, t); } diff --git a/src/layers/linear.h b/src/layers/linear.h index 0de035a3..bd2fe191 100644 --- a/src/layers/linear.h +++ b/src/layers/linear.h @@ -22,6 +22,17 @@ namespace vla { inline ggml_tensor * linear(ggml_context * C, ggml_tensor * W, ggml_tensor * b, ggml_tensor * x) { +#ifdef GGML_USE_HEXAGON + // Hexagon computes a 2D weight against a batch of activations wrong, and the + // fallback wrapper runs that on the CPU; the same rows as one 2D matmul stay + // on the NPU. + if (W->ne[2] * W->ne[3] == 1 && x->ne[2] * x->ne[3] > 1 && ggml_is_contiguous(x)) { + ggml_tensor * y = ggml_mul_mat(C, W, ggml_reshape_2d(C, x, x->ne[0], x->ne[1] * x->ne[2] * x->ne[3])); + if (b) + y = ggml_add(C, y, b); + return ggml_reshape_4d(C, y, W->ne[1], x->ne[1], x->ne[2], x->ne[3]); + } +#endif ggml_tensor * y = ggml_mul_mat(C, W, x); return b ? ggml_add(C, y, b) : y; } diff --git a/src/loader.cpp b/src/loader.cpp index 54826859..7730d6b5 100644 --- a/src/loader.cpp +++ b/src/loader.cpp @@ -20,6 +20,7 @@ #include #include #include +#include namespace vla { @@ -39,11 +40,39 @@ void WeightLoader::as_float(const std::string & name, int64_t K, int64_t N, Floa as_float_[name] = AsFloat{K, N, std::move(make)}; } +void WeightLoader::pad(const std::string & name, int64_t ne0, int64_t ne1) { + pad_[name] = {ne0, ne1}; +} + +void WeightLoader::pad_mlp(const std::string & fc1, const std::string & fc2) { + const ggml_tensor * w1 = g_.meta((fc1 + ".weight").c_str()); + const ggml_tensor * w2 = g_.meta((fc2 + ".weight").c_str()); + if (!w1 || !w2 || ggml_is_quantized(gemm_) || ggml_is_quantized(w1->type) || ggml_is_quantized(w2->type)) + return; + const int64_t ff = w1->ne[1]; + const int64_t fp = gemm_width(ff); + if (fp == ff || w2->ne[0] != ff) + return; + pad(fc1 + ".weight", w1->ne[0], fp); + pad(fc1 + ".bias", fp, 1); + pad(fc2 + ".weight", fp, w2->ne[1]); +} + bool WeightLoader::read_resident(const std::string & name, ggml_type type, bool gemma_norm, std::vector & out) { const auto it = as_float_.find(name); if (it == as_float_.end()) { out = g_.read_convert(name.c_str(), type, gemma_norm); - return !out.empty(); + const auto p = pad_.find(name); + if (p == pad_.end() || out.empty()) + return !out.empty(); + const ggml_tensor * src = g_.meta(name.c_str()); + const size_t rs = ggml_row_size(type, src->ne[0]); + const size_t rp = ggml_row_size(type, p->second.first); + std::vector padded(rp * (size_t) p->second.second, 0); + for (int64_t r = 0; r < src->ne[1]; ++r) + std::memcpy(padded.data() + r * rp, out.data() + r * rs, rs); + out.swap(padded); + return true; } const AsFloat & a = it->second; std::vector f; @@ -97,7 +126,23 @@ ggml_tensor * WeightLoader::declare(ggml_type want, bool required, bool gemma_no return nullptr; } - ggml_tensor * t = ggml_new_tensor(ctx_, g_.resident_type(src, want), ggml_n_dims(src), src->ne); + const ggml_type rt = g_.resident_type(src, want); + int64_t ne[GGML_MAX_DIMS]; + std::memcpy(ne, src->ne, sizeof(ne)); + int n_dims = ggml_n_dims(src); + if (const auto p = pad_.find(name); p != pad_.end()) { + const bool float_type = rt == GGML_TYPE_F32 || rt == GGML_TYPE_F16 || rt == GGML_TYPE_BF16; + if (!float_type || n_dims > 2 || p->second.first < ne[0] || p->second.second < ne[1]) { + std::fprintf(stderr, "vla(%s): cannot pad %s to [%lld, %lld]\n", arch_, name, + (long long) p->second.first, (long long) p->second.second); + ok_ = false; + return nullptr; + } + ne[0] = p->second.first; + ne[1] = p->second.second; + n_dims = 2; + } + ggml_tensor * t = ggml_new_tensor(ctx_, rt, n_dims, ne); if (!t) { std::fprintf(stderr, "vla(%s): ggml_new_tensor failed for %s (weight context too small?)\n", arch_, name); ok_ = false; diff --git a/src/loader.h b/src/loader.h index 2f71d099..cec75760 100644 --- a/src/loader.h +++ b/src/loader.h @@ -82,6 +82,16 @@ class WeightLoader { using FloatMaker = std::function &)>; void as_float(const std::string & name, int64_t K, int64_t N, FloatMaker make); + // Declare 2D (or 1D) float tensor `name` as [ne0, ne1], at least its shape in + // the file, with the extra elements zero. Call before declaring it. + void pad(const std::string & name, int64_t ne0, int64_t ne1); + + // Widens the MLP fc1 -> GELU -> fc2 (`fc1`/`fc2` without ".weight") to + // gemm_width() with zeros: the extra fc1 rows and bias are 0, GELU(0) is 0 and + // the extra fc2 columns are 0, so the output is the same. Weights quantized in + // the file or by the loader stay as they are. Call before declaring them. + void pad_mlp(const std::string & fc1, const std::string & fc2); + // For helpers that inspect the file before declaring (foldquant.cpp). gguf_reader & reader() { return g_; @@ -111,6 +121,7 @@ class WeightLoader { FloatMaker make; }; std::map as_float_; + std::map> pad_; const char * arch_; gguf_reader & g_; diff --git a/src/models/act.cpp b/src/models/act.cpp index d05ec7d4..56586def 100644 --- a/src/models/act.cpp +++ b/src/models/act.cpp @@ -176,12 +176,34 @@ struct ActModelArch : public ModelArchBase { namespace { +// conv_2d_f32 with the patch matrix padded with zeros to gemm_width() in both +// dims, for one view. The HMX takes a matmul only when those are multiples of +// 32; the stem's K of 147 and layer3's and layer4's 1200 and 300 rows are not. +// Returns nullptr where nothing would be padded. +ggml_tensor * conv_2d_padded(ggml_context * C, ggml_tensor * w, ggml_tensor * x, int stride, int pad) { + const int64_t ow = (x->ne[0] + 2*pad - w->ne[0]) / stride + 1, oh = (x->ne[1] + 2*pad - w->ne[1]) / stride + 1; + const int64_t K = w->ne[0]*w->ne[1]*w->ne[2], R = ow*oh, Kp = gemm_width(K), Rp = gemm_width(R); + if (x->ne[3] != 1 || (Kp == K && Rp == R)) + return nullptr; + ggml_tensor * col = ggml_im2col(C, w, x, stride, stride, pad, pad, 1, 1, true, GGML_TYPE_F32); + ggml_tensor * c2 = ggml_pad(C, ggml_reshape_2d(C, col, K, R), (int) (Kp - K), (int) (Rp - R), 0, 0); + ggml_tensor * w2 = ggml_reshape_2d(C, w, K, w->ne[3]); + if (Kp > K) + w2 = ggml_pad(C, w2, (int) (Kp - K), 0, 0, 0); + ggml_tensor * y = ggml_mul_mat(C, c2, w2); + if (Rp > R) + y = ggml_cont(C, ggml_view_2d(C, y, R, y->ne[1], y->nb[1], 0)); + return ggml_reshape_4d(C, y, col->ne[1], col->ne[2], w->ne[3], 1); +} + // [w, h, c, view] to [ow, oh, oc, view]. The kernel's type picks the path; see // ActModelArch::conv_direct. ggml_tensor * conv2d(ggml_context * C, const ConvW & cw, ggml_tensor * x, int stride, int pad) { - ggml_tensor * y = cw.w->type == GGML_TYPE_F16 - ? ggml_conv_2d_direct(C, cw.w, x, stride, stride, pad, pad, 1, 1) - : conv_2d_f32(C, cw.w, x, stride, pad); + ggml_tensor * y = nullptr; + if (cw.w->type == GGML_TYPE_F16) + y = ggml_conv_2d_direct(C, cw.w, x, stride, stride, pad, pad, 1, 1); + else if (!(y = conv_2d_padded(C, cw.w, x, stride, pad))) + y = conv_2d_f32(C, cw.w, x, stride, pad); return ggml_add(C, y, ggml_reshape_4d(C, cw.b, 1, 1, cw.b->ne[0], 1)); } diff --git a/src/models/evo1.cpp b/src/models/evo1.cpp index 547e0024..ce31ffab 100644 --- a/src/models/evo1.cpp +++ b/src/models/evo1.cpp @@ -188,8 +188,12 @@ ggml_tensor * build_qwen2_layer(ggml_context * C, const Evo1ModelArch & m, const att = ggml_reshape_2d(C, fa, hq, seq); } else { ggml_tensor * Q = ggml_cont(C, ggml_permute(C, q_rope, 0, 2, 1, 3)); - ggml_tensor * K = ggml_cont(C, ggml_permute(C, k_rope, 0, 2, 1, 3)); - ggml_tensor * V = ggml_cont(C, ggml_permute(C, ggml_reshape_3d(C, vp, hd, n_kv, seq), 1, 2, 0, 3)); + // fa_kv: on Hexagon K and V go to F16, so both matmuls run on the HMX, which + // takes no batched F32 src0 (1 s of vector-unit time per chunk on the IQ-10). + // The scores still accumulate in F32 and the softmax stays F32: 4.3e-4 from + // the CPU F32 reference there. Other backends keep F32 K/V. + ggml_tensor * K = fa_kv(C, ggml_cont(C, ggml_permute(C, k_rope, 0, 2, 1, 3))); + ggml_tensor * V = fa_kv(C, ggml_cont(C, ggml_permute(C, ggml_reshape_3d(C, vp, hd, n_kv, seq), 1, 2, 0, 3))); ggml_tensor * kq = ggml_mul_mat(C, K, Q); ggml_prec_set_acc(kq, GGML_PREC_F32); ggml_tensor * aw = ggml_soft_max_ext(C, kq, mask, scale, 0.0f); ggml_tensor * kqv = ggml_mul_mat(C, V, aw); diff --git a/src/models/octo.cpp b/src/models/octo.cpp index da5c0205..4774c722 100644 --- a/src/models/octo.cpp +++ b/src/models/octo.cpp @@ -1083,6 +1083,10 @@ bool run_transformer_graph(OctoRuntime& rt, const int n_readout = (int) layout.readout_seq_idx.size(); if (input.size() != (size_t) kHidden*seq) return false; + // Zero tokens appended to a multiple of 32 where gemm_width() pads (Hexagon), + // masked out as keys: the softmax rows are then a width the NPU takes. The + // readouts are gathered by index, so the extra rows never reach them. + const int pseq = (int) gemm_width(seq); OctoRuntime::BtKey key; for (const OctoSeqRun& r : layout.runs) @@ -1109,12 +1113,12 @@ bool run_transformer_graph(OctoRuntime& rt, if (!out_w || !out_b) return nullptr; - ggml_tensor * x = ggml_new_tensor_2d(C, GGML_TYPE_F32, kHidden, seq); + ggml_tensor * x = ggml_new_tensor_2d(C, GGML_TYPE_F32, kHidden, pseq); ggml_set_name(x, "octo.block_transformer.input"); ggml_set_input(x); io.input = x; - ggml_tensor * mask = ggml_new_tensor_2d(C, GGML_TYPE_F32, seq, seq); + ggml_tensor * mask = ggml_new_tensor_2d(C, GGML_TYPE_F32, pseq, pseq); ggml_set_name(mask, "octo.block_transformer.additive_mask"); ggml_set_input(mask); ggml_set_output(mask); @@ -1127,7 +1131,7 @@ bool run_transformer_graph(OctoRuntime& rt, io.readout_idx = readout_idx; for (int i=0; i<12; ++i) { - if (rt.fused_ops && seq <= 4096) { + if (rt.fused_ops && pseq == seq && seq <= 4096) { // The same block on the fused (bit-identical) kernels: LayerNorm, // bias + head split + V^T in one pass, exact attention with the // mask, and the bias + residual and bias + GELU epilogues. @@ -1144,11 +1148,11 @@ bool run_transformer_graph(OctoRuntime& rt, } ggml_tensor * n1 = layer_norm(C, x, blk_w[i][0], blk_w[i][1], ln_eps); ggml_tensor * qkv = linear(C, blk_w[i][2], blk_w[i][3], n1); - ggml_tensor * Q = ggml_cont(C, ggml_permute(C, head_view(C, qkv, head_dim, heads, seq, kHidden, 3, 0), 0, 2, 1, 3)); - ggml_tensor * K = ggml_cont(C, ggml_permute(C, head_view(C, qkv, head_dim, heads, seq, kHidden, 3, 1), 0, 2, 1, 3)); - ggml_tensor * V = ggml_cont(C, ggml_permute(C, head_view(C, qkv, head_dim, heads, seq, kHidden, 3, 2), 1, 2, 0, 3)); + ggml_tensor * Q = ggml_cont(C, ggml_permute(C, head_view(C, qkv, head_dim, heads, pseq, kHidden, 3, 0), 0, 2, 1, 3)); + ggml_tensor * K = ggml_cont(C, ggml_permute(C, head_view(C, qkv, head_dim, heads, pseq, kHidden, 3, 1), 0, 2, 1, 3)); + ggml_tensor * V = ggml_cont(C, ggml_permute(C, head_view(C, qkv, head_dim, heads, pseq, kHidden, 3, 2), 1, 2, 0, 3)); - ggml_tensor * merged = attention(C, Q, K, V, mask, attn_scale, kHidden, seq); + ggml_tensor * merged = attention(C, Q, K, V, mask, attn_scale, kHidden, pseq); ggml_tensor * residual = ggml_add(C, x, linear(C, blk_w[i][4], blk_w[i][5], merged)); ggml_tensor * n2 = layer_norm(C, residual, blk_w[i][6], blk_w[i][7], ln_eps); @@ -1177,10 +1181,21 @@ bool run_transformer_graph(OctoRuntime& rt, if (fresh) { std::vector mask; build_transformer_mask(layout, mask); + if (pseq > seq) { + std::vector padded((size_t) pseq*pseq, 0.0f); + for (int q=0; q= seq ? -FLT_MAX : q < seq ? mask[(size_t) q*seq+k] : 0.0f; + mask.swap(padded); + } ggml_backend_tensor_set(io.mask, mask.data(), 0, ggml_nbytes(io.mask)); ggml_backend_tensor_set(io.readout_idx, layout.readout_seq_idx.data(), 0, ggml_nbytes(io.readout_idx)); } - ggml_backend_tensor_set(io.input, input.data(), 0, ggml_nbytes(io.input)); + ggml_backend_tensor_set(io.input, input.data(), 0, input.size()*sizeof(float)); + if (pseq > seq) { + const std::vector zeros((size_t) kHidden*(pseq-seq), 0.0f); + ggml_backend_tensor_set(io.input, zeros.data(), input.size()*sizeof(float), zeros.size()*sizeof(float)); + } if (!octo_compute(rt, rt.transformer.graph(), "block transformer")) return false; readout_action.resize((size_t) kHidden*n_readout); diff --git a/src/models/pi05.cpp b/src/models/pi05.cpp index 1f4faa31..6eeece50 100644 --- a/src/models/pi05.cpp +++ b/src/models/pi05.cpp @@ -146,6 +146,14 @@ namespace { // straight off the [hd, H, S] views. Otherwise the unfused graph is used. bool use_fused_attention() { return vla::flash_attn_enabled() && vla::fa_takes_views(); } +// Where flash attention is on by default (Hexagon), the three sites also take the +// plain flash attention when the kernel does not take views. Their explicit +// scores are matmuls the NPU does not run: SigLIP's at head dim 72, and the VLM's +// and the expert's, whose single KV head the fallback wrapper sends to the CPU. +bool use_flash_attention() { + return use_fused_attention() || (vla::default_flash_attn() && vla::flash_attn_enabled()); +} + ggml_tensor * build_adarms( ggml_context * ctx, ggml_tensor * x, ggml_tensor * mod, int64_t h, float eps, ggml_tensor ** gate_out) { @@ -185,7 +193,7 @@ ggml_tensor * build_expert_layer( ggml_tensor * gate_ffn = ggml_view_1d(ctx, mod_ffn, h, (size_t) 2*h * sizeof(float)); ggml_tensor * xq_attn = fq_act_mod(ctx, &w.act_mod_qkv, w.fq_qkv, x_in, fqx->scale1_attn, fqx->shift_attn); ggml_tensor * h1 = gemma_attn(ctx, w, nullptr, positions, cfg, seq, cached_K, cached_V, mask, nullptr, nullptr, - GGML_TYPE_F32, use_fused_attention(), nullptr, x_in, nullptr, fqx->K16, fqx->V16, + GGML_TYPE_F32, use_flash_attention(), nullptr, x_in, nullptr, fqx->K16, fqx->V16, xq_attn, gate_attn); ggml_tensor * xq_ffn = fq_act_mod(ctx, &w.act_mod_gu, w.fq_gu, h1, fqx->scale1_ffn, fqx->shift_ffn); return gemma_mlp(ctx, w, nullptr, GGML_TYPE_F32, nullptr, h1, xq_ffn, gate_ffn); @@ -207,13 +215,13 @@ ggml_tensor * build_expert_layer( // out * gate + x reading the GEMM output in whatever type it has. ggml_tensor * xn = fused_adarms(ctx, x_in, mod_attn, at, *fa); ggml_tensor * o = gemma_attn(ctx, w, xn, positions, cfg, seq, cached_K, cached_V, mask, - nullptr, nullptr, at, use_fused_attention(), nullptr, nullptr, kv); + nullptr, nullptr, at, use_flash_attention(), nullptr, nullptr, kv); ggml_tensor * h1 = fused_gateres(ctx, x_in, o, mod_attn, *fg); ggml_tensor * xm = fused_adarms(ctx, h1, mod_ffn, at, *fa); return fused_gateres(ctx, h1, gemma_mlp(ctx, w, xm, at), mod_ffn, *fg); } - if (!mask && !kv && gemma_fused(w, at, use_fused_attention(), seq, seq + (cached_K ? cached_K->ne[2] : 0))) { + if (!mask && !kv && gemma_fused(w, at, use_flash_attention(), seq, seq + (cached_K ? cached_K->ne[2] : 0))) { auto part = [&](ggml_tensor * mod, int j) { return ggml_view_1d(ctx, mod, h, (size_t) j*h*sizeof(float)); }; @@ -237,7 +245,7 @@ ggml_tensor * build_expert_layer( // mm_act does not convert its input: a BF16 GEMM must be handed BF16 rows, // but adaRMS produces F32, whose bits would be misread as BF16. ggml_tensor * o_out = gemma_attn(ctx, w, as_type(ctx, x_norm, at), positions, cfg, seq, cached_K, cached_V, mask, - nullptr, nullptr, at, use_fused_attention(), x_norm, nullptr, kv, + nullptr, nullptr, at, use_flash_attention(), x_norm, nullptr, kv, fqx ? fqx->K16 : nullptr, fqx ? fqx->V16 : nullptr); o_out = as_type(ctx, o_out, GGML_TYPE_F32); @@ -544,7 +552,7 @@ std::unique_ptr pi05_create(const std::string& mmproj_path, std::printf("vla(pi05): expert weights INT4 (W4A8), %s\n", aex_w4_group > 0 ? ("group " + std::to_string(aex_w4_group)).c_str() : "per channel"); } - m->vis.vit.enc.cfg.flash_attn = use_fused_attention(); + m->vis.vit.enc.cfg.flash_attn = use_flash_attention(); // Fused SigLIP blocks (modules/encoder.h, EncCfg::fused): SYCL, BF16, fused // attention. A FoldQuant GGUF takes them too: its integer sites are the LM's // and the expert's, and its vision tower is BF16 like any other checkpoint's. @@ -837,7 +845,7 @@ std::vector Pi05ModelArch::predict(const Inputs& in) { pl.rms_spec); else h = gemma_layer(C, pl.blk[i], h, t_prefix_pos, cfg, n_prefix_g, - nullptr, nullptr, t_prefix_mask, &cK[i], &cV[i], act_type, use_fused_attention(), + nullptr, nullptr, t_prefix_mask, &cK[i], &cV[i], act_type, use_flash_attention(), kv_inplace ? &kv[i] : nullptr); } (void) h; diff --git a/src/models/smolvla.cpp b/src/models/smolvla.cpp index d298ec11..02c11414 100644 --- a/src/models/smolvla.cpp +++ b/src/models/smolvla.cpp @@ -732,8 +732,8 @@ ggml_tensor * build_expert_self_attn_layer( ggml_tensor * k_h = ggml_reshape_3d(ctx, mm_w(ctx, w.Wk, x_norm), cfg.head_dim, cfg.n_kv_heads, cfg.n_suffix); ggml_tensor * v_h = ggml_reshape_3d(ctx, mm_w(ctx, w.Wv, x_norm), cfg.head_dim, cfg.n_kv_heads, cfg.n_suffix); - ggml_tensor * K_full = ggml_concat(ctx, cached_K, rope(ctx, rope_spec(cfg), k_h, positions_full), 2); - ggml_tensor * V_full = ggml_concat(ctx, cached_V, v_h, 2); + ggml_tensor * K_full = concat_dim2(ctx, cached_K, rope(ctx, rope_spec(cfg), k_h, positions_full)); + ggml_tensor * V_full = concat_dim2(ctx, cached_V, v_h); return attn_mlp(ctx, w, x_in, rope(ctx, rope_spec(cfg), q_h, positions_full), K_full, V_full, mask_full, cfg, cfg.n_suffix); } diff --git a/src/models/vla_adapter.cpp b/src/models/vla_adapter.cpp index 234b500f..25c65a13 100644 --- a/src/models/vla_adapter.cpp +++ b/src/models/vla_adapter.cpp @@ -34,12 +34,14 @@ #include "layers/norm.h" #include "layers/rope.h" +#include #include #include #include #include #include #include +#include #include namespace vla { @@ -77,6 +79,7 @@ struct VlaAdapterModelArch : public ModelArchBase { ggml_tensor *t_ids=nullptr,*t_proj=nullptr,*t_pos=nullptr,*t_mask=nullptr; ggml_tensor *t_state=nullptr,*t_x0=nullptr,*norm_actions=nullptr; ggml_tensor *cT=nullptr,*sT=nullptr,*cA=nullptr,*sA=nullptr,*cK=nullptr,*sK=nullptr; + ggml_tensor *cJ=nullptr,*sJ=nullptr,*hmask=nullptr; bool consts=false; }; graph_cache main_graph; @@ -259,6 +262,10 @@ std::vector VlaAdapterModelArch::predict(const Inputs& in) { const int64_t NUM_PROMPT_TOKENS = NPROMPT-1; const int64_t NPATCH = NP * n_views; const int64_t SEQ = 1+NPATCH+(NPROMPT-1)+num_tokens+1; + // The LM runs on LSEQ tokens, zeros appended. The mask is causal, so no real + // token sees them, and on Hexagon the softmax rows become a multiple of 32, + // which is what the NPU takes. + const int64_t LSEQ = gemm_width(SEQ); const auto ti=clock::now(); // LM + action head graph depends only on the sequence layout. const MainKey mkey{ SEQ, n_views, NPROMPT }; @@ -276,12 +283,14 @@ std::vector VlaAdapterModelArch::predict(const Inputs& in) { ggml_tensor*erest=ggml_cont(C,ggml_view_2d(C,emb2,HC,NPROMPT+num_tokens,emb2->nb[1],emb2->nb[1])); ggml_tensor*t_proj=ggml_new_tensor_2d(C,GGML_TYPE_F32,HC,NPATCH); ggml_set_input(t_proj); ggml_tensor*mm_seq=ggml_concat(C,ggml_concat(C,e0,t_proj,1),erest,1); + if(LSEQ>SEQ) + mm_seq=ggml_pad(C,mm_seq,0,(int)(LSEQ-SEQ),0,0); - ggml_tensor*t_pos=ggml_new_tensor_1d(C,GGML_TYPE_I32,SEQ); ggml_set_input(t_pos); ggml_set_output(t_pos); - ggml_tensor*t_mask=ggml_new_tensor_2d(C,GGML_TYPE_F32,SEQ,SEQ); ggml_set_input(t_mask); ggml_set_output(t_mask); + ggml_tensor*t_pos=ggml_new_tensor_1d(C,GGML_TYPE_I32,LSEQ); ggml_set_input(t_pos); ggml_set_output(t_pos); + ggml_tensor*t_mask=ggml_new_tensor_2d(C,GGML_TYPE_F32,LSEQ,LSEQ); ggml_set_input(t_mask); ggml_set_output(t_mask); const float lsc=1.0f/std::sqrt((float)lm_head_dim); std::vector lout(lm_layers); ggml_tensor*x=mm_seq; - const bool fused_lm=fused_ops && cublas_rows(SEQ) && SEQ<=4096; + const bool fused_lm=fused_ops && LSEQ==SEQ && cublas_rows(SEQ) && SEQ<=4096; if(fused){ const FusedLmDims d{lm_head_dim,n_q,n_kv,lm_rope_base}; ggml_tensor*kbuf,*vbuf; fused_lm_kv(C,d,SEQ,&kbuf,&vbuf); @@ -305,19 +314,22 @@ std::vector VlaAdapterModelArch::predict(const Inputs& in) { } ggml_tensor*hn=ggml_mul(C,ggml_rms_norm(C,x,lm_rms_eps),l.attn_norm); ggml_tensor*qp=ggml_add(C,ggml_mul_mat(C,l.Wq,hn),l.bq),*kp=ggml_add(C,ggml_mul_mat(C,l.Wk,hn),l.bk),*vp=ggml_add(C,ggml_mul_mat(C,l.Wv,hn),l.bv); - ggml_tensor*qh=ggml_reshape_3d(C,qp,lm_head_dim,n_q,SEQ),*kh=ggml_reshape_3d(C,kp,lm_head_dim,n_kv,SEQ),*vh=ggml_reshape_3d(C,vp,lm_head_dim,n_kv,SEQ); + ggml_tensor*qh=ggml_reshape_3d(C,qp,lm_head_dim,n_q,LSEQ),*kh=ggml_reshape_3d(C,kp,lm_head_dim,n_kv,LSEQ),*vh=ggml_reshape_3d(C,vp,lm_head_dim,n_kv,LSEQ); ggml_tensor*qr=ggml_rope_ext(C,qh,t_pos,nullptr,(int)lm_head_dim,GGML_ROPE_TYPE_NEOX,0,lm_rope_base,1.0f,0.0f,1.0f,32.0f,1.0f); ggml_tensor*kr=ggml_rope_ext(C,kh,t_pos,nullptr,(int)lm_head_dim,GGML_ROPE_TYPE_NEOX,0,lm_rope_base,1.0f,0.0f,1.0f,32.0f,1.0f); ggml_tensor*Q=ggml_cont(C,ggml_permute(C,qr,0,2,1,3)),*K=ggml_cont(C,ggml_permute(C,kr,0,2,1,3)),*V=ggml_cont(C,ggml_permute(C,vh,1,2,0,3)); ggml_tensor*kq=ggml_mul_mat(C,K,Q); ggml_prec_set_acc(kq,GGML_PREC_F32); ggml_tensor*aw=ggml_soft_max_ext(C,kq,t_mask,lsc,0.0f); ggml_tensor*kqv=ggml_mul_mat(C,V,aw); - ggml_tensor*att=ggml_reshape_2d(C,ggml_cont(C,ggml_permute(C,kqv,0,2,1,3)),HC,SEQ); + ggml_tensor*att=ggml_reshape_2d(C,ggml_cont(C,ggml_permute(C,kqv,0,2,1,3)),HC,LSEQ); x=ggml_add(C,x,ggml_mul_mat(C,l.Wo,att)); ggml_tensor*hn2=ggml_mul(C,ggml_rms_norm(C,x,lm_rms_eps),l.ffn_norm); ggml_tensor*gt=ggml_silu(C,ggml_mul_mat(C,l.Wg,hn2)),*ut=ggml_mul_mat(C,l.Wu,hn2); x=ggml_add(C,x,ggml_mul_mat(C,l.Wd,ggml_mul(C,gt,ut))); lout[i]=x; } + if(LSEQ>SEQ) + for(auto&o:lout) + o=ggml_view_2d(C,o,HC,SEQ,o->nb[1],0); ggml_tensor*final_norm=ggml_mul(C,ggml_rms_norm(C,lout[lm_layers-1],lm_rms_eps),lm_out_norm); std::vector cond(head_blocks); @@ -333,9 +345,21 @@ std::vector VlaAdapterModelArch::predict(const Inputs& in) { ggml_tensor*ss=ggml_new_tensor_2d(C,GGML_TYPE_F32,HD,Lh); ggml_set_input(ss); ggml_set_output(ss); return {cc,ss}; }; + // Each head block attends its own tokens, the action queries and the vision + // tokens. Where gemm_width() pads (Hexagon), their keys and values are + // stacked as rows and padded to a multiple of 32, masked out: one softmax + // the NPU takes, and no column concats. + const int64_t LJ=chunk+num_tokens+1+NPATCH, LH=gemm_width(LJ); + const bool joint=LH>LJ; auto [cT,sT]=cs_tensor(chunk); - auto [cA,sA]=cs_tensor(num_tokens+1); - auto [cK,sK]=cs_tensor(NPATCH); + ggml_tensor *cA=nullptr,*sA=nullptr,*cK=nullptr,*sK=nullptr,*cJ=nullptr,*sJ=nullptr,*hmask=nullptr; + if(joint){ + std::tie(cJ,sJ)=cs_tensor(LH); + hmask=ggml_new_tensor_2d(C,GGML_TYPE_F32,LH,chunk); ggml_set_input(hmask); ggml_set_output(hmask); + }else{ + std::tie(cA,sA)=cs_tensor(num_tokens+1); + std::tie(cK,sK)=cs_tensor(NPATCH); + } ggml_tensor*t_x0=ggml_new_tensor_2d(C,GGML_TYPE_F32,action_dim*HC,chunk); ggml_set_input(t_x0); ggml_set_output(t_x0); ggml_tensor*hx=ggml_relu(C,linear(C,h_fc1w,h_fc1b,layer_norm(C,t_x0,h_ln1w,h_ln1b,head_ln_eps))); @@ -348,18 +372,31 @@ std::vector VlaAdapterModelArch::predict(const Inputs& in) { auto qk=[&](ggml_tensor*p,ggml_tensor*cs,ggml_tensor*sn){ return rope_pairwise(C,to_heads(C,p,HD,NH,p->ne[1]),cs,sn,HD); }; auto tov=[&](ggml_tensor*p){ return to_heads_v(C,p,HD,NH,p->ne[1]); }; ggml_tensor*q=qk(linear(C,w.Wq,w.bq,hx),cT,sT); - ggml_tensor*kse=qk(linear(C,w.Wks,w.bks,hx),cT,sT); - ggml_tensor*vse=linear(C,w.Wvs,w.bvs,hx); - ggml_tensor*kad=qk(linear(C,w.Wka,w.bka,had),cA,sA); ggml_tensor*vad=linear(C,w.Wva,w.bva,had); - ggml_tensor*kta=qk(linear(C,w.Wkt,w.bkt,ht),cK,sK); ggml_tensor*vta=linear(C,w.Wvt,w.bvt,ht); - ggml_tensor*Vs=tov(vse),*Va=tov(vad),*VT=tov(vta); - ggml_tensor*ss2=ggml_mul_mat(C,kse,q),*sa=ggml_mul_mat(C,kad,q),*sr=ggml_mul_mat(C,kta,q); - ggml_prec_set_acc(ss2,GGML_PREC_F32); ggml_prec_set_acc(sa,GGML_PREC_F32); ggml_prec_set_acc(sr,GGML_PREC_F32); - ggml_tensor*st2=ggml_scale(C,sr,w.rg); - ggml_tensor*scr=ggml_concat(C,ggml_concat(C,ss2,sa,0),st2,0); - ggml_tensor*attn=ggml_soft_max_ext(C,scr,nullptr,hsc,0.0f); - ggml_tensor*Vc=ggml_concat(C,ggml_concat(C,Vs,Va,0),VT,0); - ggml_tensor*kqv=ggml_mul_mat(C,Vc,attn); + ggml_tensor*kqv; + if(joint){ + // rg scales the vision scores; RoPE is linear, so it goes on their keys. + ggml_tensor*kj=ggml_concat(C,ggml_concat(C,linear(C,w.Wks,w.bks,hx),linear(C,w.Wka,w.bka,had),1), + ggml_scale(C,linear(C,w.Wkt,w.bkt,ht),w.rg),1); + ggml_tensor*vj=ggml_concat(C,ggml_concat(C,linear(C,w.Wvs,w.bvs,hx),linear(C,w.Wva,w.bva,had),1), + linear(C,w.Wvt,w.bvt,ht),1); + ggml_tensor*K=qk(ggml_pad(C,kj,0,(int)(LH-LJ),0,0),cJ,sJ); + ggml_tensor*V=tov(ggml_pad(C,vj,0,(int)(LH-LJ),0,0)); + ggml_tensor*sc=ggml_mul_mat(C,K,q); ggml_prec_set_acc(sc,GGML_PREC_F32); + kqv=ggml_mul_mat(C,V,ggml_soft_max_ext(C,sc,hmask,hsc,0.0f)); + }else{ + ggml_tensor*kse=qk(linear(C,w.Wks,w.bks,hx),cT,sT); + ggml_tensor*vse=linear(C,w.Wvs,w.bvs,hx); + ggml_tensor*kad=qk(linear(C,w.Wka,w.bka,had),cA,sA); ggml_tensor*vad=linear(C,w.Wva,w.bva,had); + ggml_tensor*kta=qk(linear(C,w.Wkt,w.bkt,ht),cK,sK); ggml_tensor*vta=linear(C,w.Wvt,w.bvt,ht); + ggml_tensor*Vs=tov(vse),*Va=tov(vad),*VT=tov(vta); + ggml_tensor*ss2=ggml_mul_mat(C,kse,q),*sa=ggml_mul_mat(C,kad,q),*sr=ggml_mul_mat(C,kta,q); + ggml_prec_set_acc(ss2,GGML_PREC_F32); ggml_prec_set_acc(sa,GGML_PREC_F32); ggml_prec_set_acc(sr,GGML_PREC_F32); + ggml_tensor*st2=ggml_scale(C,sr,w.rg); + ggml_tensor*scr=ggml_concat(C,ggml_concat(C,ss2,sa,0),st2,0); + ggml_tensor*attn=ggml_soft_max_ext(C,scr,nullptr,hsc,0.0f); + ggml_tensor*Vc=ggml_concat(C,ggml_concat(C,Vs,Va,0),VT,0); + kqv=ggml_mul_mat(C,Vc,attn); + } ggml_tensor*mg=ggml_reshape_2d(C,ggml_cont(C,ggml_permute(C,kqv,0,2,1,3)),HC,chunk); ggml_tensor*res=ggml_add(C,linear(C,w.Wo,w.bo,mg),hx); hx=ggml_relu(C,linear(C,w.flw,w.flb,layer_norm(C,res,w.flnw,w.flnb,head_ln_eps))); @@ -369,6 +406,7 @@ std::vector VlaAdapterModelArch::predict(const Inputs& in) { gio.t_ids=t_ids; gio.t_proj=t_proj; gio.t_pos=t_pos; gio.t_mask=t_mask; gio.t_state=t_state; gio.t_x0=t_x0; gio.norm_actions=norm_actions; gio.cT=cT; gio.sT=sT; gio.cA=cA; gio.sA=sA; gio.cK=cK; gio.sK=sK; + gio.cJ=cJ; gio.sJ=sJ; gio.hmask=hmask; ggml_cgraph*gf=ggml_new_graph_custom(C,65536,false); ggml_build_forward_expand(gf,norm_actions); return gf; @@ -391,17 +429,33 @@ std::vector VlaAdapterModelArch::predict(const Inputs& in) { if(!ggml_are_same_shape(proj,t_proj)){ std::fprintf(stderr,"vla(vla_adapter): projector output does not match the LM width\n"); return {}; } ggml_backend_tensor_copy(proj,t_proj); if(!gio.consts){ - std::vector pp(SEQ); - for(int64_t i=0;i pp(LSEQ); + for(int64_t i=0;i mk; build_causal_mask(SEQ, mk); + std::vector mk; build_causal_mask(LSEQ, mk); ggml_backend_tensor_set(t_mask,mk.data(),0,ggml_nbytes(t_mask)); std::vector zx((size_t)action_dim*HC*chunk,0.0f); ggml_backend_tensor_set(t_x0,zx.data(),0,ggml_nbytes(t_x0)); auto fill_cs=[&](ggml_tensor*cc,ggml_tensor*ss,int64_t Lh){ std::vector cb,sb; rope_pairwise_table(HD,Lh,head_rope_base,cb,sb); ggml_backend_tensor_set(cc,cb.data(),0,ggml_nbytes(cc)); ggml_backend_tensor_set(ss,sb.data(),0,ggml_nbytes(ss)); }; - fill_cs(cT,sT,chunk); fill_cs(cA,sA,num_tokens+1); fill_cs(cK,sK,NPATCH); + fill_cs(cT,sT,chunk); + if(gio.cJ){ + std::vector cb,sb,cj,sj; + for(int64_t Lh:{chunk,num_tokens+1,NPATCH}){ + rope_pairwise_table(HD,Lh,head_rope_base,cb,sb); + cj.insert(cj.end(),cb.begin(),cb.end()); sj.insert(sj.end(),sb.begin(),sb.end()); + } + const int64_t LH=gio.cJ->ne[1], LJ=(int64_t)cj.size()/HD; + cj.resize((size_t)(HD*LH),1.0f); sj.resize((size_t)(HD*LH),0.0f); + ggml_backend_tensor_set(gio.cJ,cj.data(),0,ggml_nbytes(gio.cJ)); ggml_backend_tensor_set(gio.sJ,sj.data(),0,ggml_nbytes(gio.sJ)); + std::vector hm((size_t)(LH*chunk),0.0f); + for(int64_t r=0;r sv(proprio_dim,0.0f); for(int64_t i=0;i&blk, - int64_t hidden, int64_t heads, int64_t hd, int64_t patch, float eps, bool prefix, bool fused){ + int64_t hidden, int64_t heads, int64_t hd, int64_t patch, float eps, bool prefix, bool fused, bool fa){ ggml_tensor*x=tower_embed(C,pix,pw,pb,pos,cls,reg,hidden,patch,prefix); const int64_t nprefix=prefix?5:0, N=x->ne[1], NP=N-nprefix; for(size_t i=0;inb[1],nprefix*x->nb[1])); return x; @@ -221,6 +225,10 @@ inline ggml_tensor* DualTower::encode(ggml_backend_t backend, graph_cacheggml_cgraph*{ @@ -258,8 +266,8 @@ inline ggml_tensor* DualTower::encode(ggml_backend_t backend, graph_cache +#include #include namespace vla { @@ -64,6 +65,8 @@ void EncStack::declare(WeightLoader & L, const char * prefix, int64_t layers, co w.q_fc1 = L.quant_i8(b + n.fc1 + ".q8", {b + n.fc1 + ".weight"}); w.q_fc2 = L.quant_i8(b + n.fc2 + ".q8", {b + n.fc2 + ".weight"}); } else { + const std::string b = std::string(prefix) + "." + std::to_string(i) + "."; + L.pad_mlp(b + n.fc1, b + n.fc2); w.Wo = L.gemm("%s.%lld.attn_o.weight", prefix, (long long)i); w.Wfc1 = L.gemm("%s.%lld.%s.weight", prefix, (long long)i, n.fc1); w.Wfc2 = L.gemm("%s.%lld.%s.weight", prefix, (long long)i, n.fc2); diff --git a/src/modules/gemma_expert.h b/src/modules/gemma_expert.h index 1c66fe3f..e7641166 100644 --- a/src/modules/gemma_expert.h +++ b/src/modules/gemma_expert.h @@ -525,8 +525,8 @@ inline ggml_tensor * gemma_attn( ggml_tensor * K_full = k_rope; ggml_tensor * V_full = v_h; if (cached_K && cached_V) { - K_full = ggml_concat(ctx, cached_K, k_rope, 2); - V_full = ggml_concat(ctx, cached_V, v_h, 2); + K_full = concat_dim2(ctx, cached_K, k_rope); + V_full = concat_dim2(ctx, cached_V, v_h); } const float scale = 1.f/std::sqrt((float) hd);