bench: CUDA benchmark on Quadro T2000 (4 GB)

CMakeLists.txt

@@ -286,6 +286,24 @@ if(GGML_CUDA)
     target_compile_definitions(vla_core PUBLIC GGML_USE_CUDA)
 
     enable_language(CUDA)
+
+    # Some kernels need a minimum architecture. Build them for the requested
+    # architectures that reach it, or for the minimum when none do, so an sm_75
+    # or Pascal build still links.
+    function(vla_cuda_min_arch target min)
+        set(archs "")
+        foreach(arch ${CMAKE_CUDA_ARCHITECTURES})
+            string(REGEX MATCH "^[0-9]+" number "${arch}")
+            if(NOT number OR number GREATER_EQUAL min)
+                list(APPEND archs ${arch})
+            endif()
+        endforeach()
+        if(NOT archs)
+            set(archs ${min})
+        endif()
+        set_target_properties(${target} PROPERTIES CUDA_ARCHITECTURES "${archs}")
+    endfunction()
+
     add_library(bitvla_cuda_kernels STATIC
         src/kernels/bitvla/bitnet_kernels.cu
         src/kernels/bitvla/bitvla_lm_cuda.cu
@@ -296,6 +314,7 @@ if(GGML_CUDA)
     # relocatable device code of both CUDA libraries together, and CUDA 12.x
     # rejects the mix because only bitvla compiles with --use_fast_math:
     # "nvlink fatal: Cicc option values for '-ftz' do not match".
+    vla_cuda_min_arch(bitvla_cuda_kernels 75)  # ldmatrix
     set_target_properties(bitvla_cuda_kernels PROPERTIES
         CUDA_SEPARABLE_COMPILATION ON
         CUDA_RESOLVE_DEVICE_SYMBOLS ON
@@ -321,6 +340,7 @@ if(GGML_CUDA)
         src/kernels/foldquant/fq_gemm_i8.cu
         src/kernels/foldquant/fq_gemm_mma.cu
     )
+    vla_cuda_min_arch(vla_fq_kernels 80)  # cp.async, mma
     set_target_properties(vla_fq_kernels PROPERTIES
         CUDA_SEPARABLE_COMPILATION ON
         CUDA_RESOLVE_DEVICE_SYMBOLS ON
@@ -364,6 +384,7 @@ if(GGML_CUDA)
     # which ggml builds with -use_fast_math (ftz, approximate div/tanh), and no
     # other archive here does.
     add_library(vla_cuda_fused STATIC src/cuda/vla_cuda_fused.cu)
+    vla_cuda_min_arch(vla_cuda_fused 80)  # cp.async, mma
     set_target_properties(vla_cuda_fused PROPERTIES
         CUDA_SEPARABLE_COMPILATION ON
         CUDA_RESOLVE_DEVICE_SYMBOLS ON

docs/backend/cuda.md

@@ -22,6 +22,7 @@ WSL2 as in [wsl.md](wsl.md).
 
 | GPU family | Example cards | `CUDA_ARCHITECTURE` |
 |---|---|---|
+| Turing | Quadro T2000, RTX 20-series | `75` |
 | Ampere (Jetson) | Orin Nano, Orin NX | `87` |
 | Ampere (consumer) | RTX 30-series, A40 | `86` |
 | Ada Lovelace | RTX 40-series, L40 | `89` |
@@ -73,8 +74,13 @@ Per-device latency is in [docs/benchmark/](../benchmark): RTX 3060, 3090,
 
 ## The bit-exact fast path
 
+Fused CUDA ops require sm_80+ and a compatible kernel image in the build.
+The runtime checks both before creating fused nodes. On Turing, defaults
+automatically use the plain ggml CUDA graph; `VLA_FUSED_OPS=0` is unnecessary.
+BitVLA uses its CUDA ladder kernels there instead of the Ampere MMA path.
+
 At the default precision (F32 activations, the checkpoint's BF16 or
-F32 weights, no flash attention), every architecture runs a faster path that
+F32 weights, no flash attention), supported devices run a faster path that
 returns **the same bits** as the plain ggml graph. Nothing about the model's
 arithmetic changes: no lower precision, no reordered sums. Success rates
 measured at the defaults therefore carry over unchanged.

docs/benchmark/README.md

@@ -12,7 +12,8 @@ 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
-`bd166b8` (same llama.cpp pin). Each report
+`bd166b8` (same llama.cpp pin). The Quadro T2000 row ran at `4a542ed`
+with the sm_75 changes of PR #82 (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.
 
@@ -35,6 +36,7 @@ means the model does not fit or fails there; each report says which and why.
 | [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 |
+| [Quadro T2000 (CUDA)](t2000.md) | 23.2 | 48.3 | 104.9 | 2330.4 | 513.2 | - | n/r | n/r | 1286.3 | 531.9 | n/r | n/r | n/r | 1114.8 |
 | [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) | 81.7 | 147.8 | 184.6 | 489.7 | 664.0 | 530.2 | 680.9 | 739.6 | N/A | 583.3 | 1117.6 | 1114.0 | 1117.7 | 1653.5 |
 | [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 |
@@ -50,6 +52,8 @@ 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 Quadro T2000 row is the default sm_75 build on a GPU without tensor cores.
+  GR00T N1.5 does not fit at the defaults there.
 - 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).

docs/benchmark/t2000.md

@@ -0,0 +1,184 @@
+# NVIDIA Quadro T2000 (CUDA)
+
+`vla-bench` times `predict()` in-process on synthetic inputs. It measures the
+engine only: no transport, no simulator, and no claim about task success.
+
+## Test setup
+
+| | |
+|---|---|
+| Device | NVIDIA Quadro T2000, 4096 MiB GDDR6, Turing TU117 (sm_75), no tensor cores; CUDA reports 4095 MiB |
+| Host | Intel Core i7-10850H (6 cores, 12 threads), 15.6 GiB RAM visible to WSL2 |
+| Power | AC power; GPU power limit 60 W (the default and the maximum) |
+| OS | Ubuntu 26.04 LTS on WSL2, kernel 6.18.33.2 |
+| Driver / toolkit | 580.92 (WDDM) / CUDA 12.4.131, GCC 13.4.0 with the GCC 15 C++ library path |
+| Commit | `4a542ed` plus the sm_75 changes of PR #82: device and kernel-image checks before fused nodes, BitVLA on its ladder kernels |
+| llama.cpp | `b11223` |
+| Date | 2026-11-05 15:14 to 16:13 (UTC+07) |
+| Build | `-DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=75 -DVLA_BUILD_TESTS=ON` (`GGML_NATIVE=ON`); ggml and BitVLA at sm_75, fused and FoldQuant kernels at sm_80 |
+| Runtime flags | defaults; the fastest flags per model are in the second table. Fusion and BitVLA MMA disable themselves where the device does not support them |
+| 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. |
+
+**Device VRAM** is whole-device `nvidia-smi memory.used`, sampled every 0.5 s,
+because WSL does not expose per-process device memory. **Host RSS** is peak RSS
+from `/usr/bin/time -v`.
+
+## Model configuration
+
+Each model runs at its native view count and input size, with weights as
+shipped. Settings not listed are the defaults.
+
+| Model | Checkpoint | Views | Input | Lang tokens | Action chunk | Action dim | Action head | Steps |
+|---|---|--:|--:|--:|--:|--:|---|--:|
+| SmolVLA | [`smolvla-libero.gguf`](https://huggingface.co/vrfai/smolvla-libero-gguf/blob/349851cd4e872d199c75e1802f037c8026767190/smolvla-libero.gguf) | 2 | 512 | 16 | 50 | 7 | flow matching | 10 |
+| ACT | [`act-aloha-sim-transfer-cube-f32.gguf`](https://huggingface.co/ravediamond/act-aloha-sim-transfer-cube-gguf/blob/f319a4f3c1ea908502f52692624e6db3e663888c/act-aloha-sim-transfer-cube-f32.gguf) | 1 | 640x480 | 16 placeholder (state only) | 100 | 14 | Transformer decoder | 1 |
+| Octo-Small | [`octo-small-libero-f32.gguf`](https://huggingface.co/vrfai/octo-small-libero-gguf/blob/a8739ff83b1f399d324cd6fe36b4327e58bbbfd4/octo-small-libero-f32.gguf) | 2 | 256 + 128 | 16 | 4 | 7 | DDPM diffusion | 20 |
+| TurboVLA | [`turbovla-libero-f32.gguf`](https://huggingface.co/vrfai/turbovla-libero-gguf/blob/3486fe301e08d41f296a298d1553505bba54df90/turbovla-libero-f32.gguf) | 2 | 256 | 16 | 12 | 7 | single-pass decoder | 1 |
+| Evo-1 | [`evo1-libero.gguf`](https://huggingface.co/vrfai/evo1-libero-gguf/blob/c76eb353aa02c82eb6fc298180aae43b064a5918/evo1-libero.gguf) | 1 | 448 | 16 | 50 | 7 | flow matching | 32 |
+| VLA-Adapter | [`libero_object/vla-adapter-libero-object.gguf`](https://huggingface.co/vrfai/vla-adapter-libero-gguf/blob/55fc9fa8b2eb89b57ee899818a2ae822126f21cc/libero_object/vla-adapter-libero-object.gguf) | 1 | 224 | 16 | 8 | 7 | bridge-attention policy | 1 |
+| BitVLA | [`libero_object/bitvla-libero-object.gguf`](https://huggingface.co/vrfai/bitvla-libero-gguf/blob/f011a6082a2cb763609d240dd64d8a7e6dc4f063/libero_object/bitvla-libero-object.gguf) (int2) | 1 | 224 | 16 | 8 | 7 | parallel decoding | 1 |
+| GR00T N1.5 | [`gr00tn1d5-libero-object.gguf`](https://huggingface.co/vrfai/gr00tn1d5-libero-object-gguf/blob/c335d02f7c3155d20bb4a20164a050ee89678ed0/gr00tn1d5-libero-object.gguf) | 1 | 224 | 16 | 16 | 32 | flow matching (DiT) | 4 |
+| VLA-JEPA | [`vla-jepa.gguf`](https://huggingface.co/vrfai/vla-jepa-libero/blob/e6014a9833e89d22ff78d3dbea6946ab930d5385/vla-jepa.gguf) | 1 | 256 | 16 + 32 | 7 | 7 | flow matching (DiT) | 4 |
+
+- **Lang tokens** is the synthetic prompt length. VLA-JEPA adds its 32
+  `<embodied>` tokens with `--extra-token 151697 --extra-count 32`.
+- **Action dim** is the checkpoint's padded width. LIBERO uses the first 7 dims.
+- **GR00T** needs `VLA_GR00T_EMBODIMENT=new_embodiment`. **Octo** needs
+  `VLA_OCTO_UNNORM_DATASET=libero_object`.
+- **Octo** takes a 256 px primary view and a 128 px wrist view.
+- **ACT** uses synthetic state; its language tokens are placeholders.
+- **Evo-1** uses one view with an adapted prompt.
+
+All nine GGUFs match the SHA-256 and byte size published by Hugging Face at the
+pinned revisions linked above.
+
+<details><summary>Full checkpoint SHA-256 values</summary>
+
+| Model | Source revision | SHA-256 |
+|---|---|---|
+| Octo-Small | [a8739ff83b1f](https://huggingface.co/vrfai/octo-small-libero-gguf/blob/a8739ff83b1f399d324cd6fe36b4327e58bbbfd4/octo-small-libero-f32.gguf) | `9b4e719b801119839beb5f81192c2f39844c78e0337ff8e686e64d9a90b0fab5` |
+| ACT | [f319a4f3c1ea](https://huggingface.co/ravediamond/act-aloha-sim-transfer-cube-gguf/blob/f319a4f3c1ea908502f52692624e6db3e663888c/act-aloha-sim-transfer-cube-f32.gguf) | `96c8f432f3db2d0cdfbcd74e9db09bc7b80632be49cce2503b4e79ba287c56a7` |
+| TurboVLA | [3486fe301e08](https://huggingface.co/vrfai/turbovla-libero-gguf/blob/3486fe301e08d41f296a298d1553505bba54df90/turbovla-libero-f32.gguf) | `6f44aa37f65d9210e46732ebd987ebe19ffb6a87a8b1168f5e603d88dd8c6131` |
+| VLA-Adapter | [55fc9fa8b2eb](https://huggingface.co/vrfai/vla-adapter-libero-gguf/blob/55fc9fa8b2eb89b57ee899818a2ae822126f21cc/libero_object/vla-adapter-libero-object.gguf) | `d26b57644ce2f6480d7ec90009f60a66ce5d1129d5ad0518a0f6ac0633f324bb` |
+| SmolVLA | [349851cd4e87](https://huggingface.co/vrfai/smolvla-libero-gguf/blob/349851cd4e872d199c75e1802f037c8026767190/smolvla-libero.gguf) | `6fb2d475c98b4c2cef3e27c4eff4e67b483740cbf983fff320a3b8a5e5f74fe8` |
+| Evo-1 | [c76eb353aa02](https://huggingface.co/vrfai/evo1-libero-gguf/blob/c76eb353aa02c82eb6fc298180aae43b064a5918/evo1-libero.gguf) | `ef61ec7c4041826b16ffceac222934877ebeaae1002fd5b72c3afa8b301dafd3` |
+| BitVLA | [f011a6082a2c](https://huggingface.co/vrfai/bitvla-libero-gguf/blob/f011a6082a2cb763609d240dd64d8a7e6dc4f063/libero_object/bitvla-libero-object.gguf) | `3a1dac4a1b5614881d38bed20c84867ee6dcabdc7848dfead0b52cf67b149e5a` |
+| VLA-JEPA | [e6014a9833e8](https://huggingface.co/vrfai/vla-jepa-libero/blob/e6014a9833e89d22ff78d3dbea6946ab930d5385/vla-jepa.gguf) | `c65a52613ed6850e5191130d0623ff0762a553b85815c77919ffdc6851332cc4` |
+| GR00T N1.5 | [c335d02f7c31](https://huggingface.co/vrfai/gr00tn1d5-libero-object-gguf/blob/c335d02f7c3155d20bb4a20164a050ee89678ed0/gr00tn1d5-libero-object.gguf) | `d12d3ecffae4e95b0274a4351549c536e4fb23e2d0691cbe1b036c4f856a4cab` |
+
+</details>
+
+## Latency and memory
+
+Rows are sorted by latency. `vision ms` is `-` for archs that do not time their
+vision stage separately. Every row runs on CUDA.
+
+| Model | Views | Input | min ms | mean ms | p50 ms | p90 ms | vision ms | Device VRAM MiB | Host RSS MiB |
+|---|--:|--:|--:|--:|--:|--:|--:|--:|--:|
+| Octo-Small | 2 | 256 + 128 | 23.2 | 24.6 | 24.1 | 26.4 | 2.8 | 627 | 400.4 |
+| ACT | 1 | 640x480 | 48.3 | 48.7 | 48.6 | 49.1 | - | 213 | 300.8 |
+| TurboVLA | 2 | 256 | 104.9 | 105.5 | 105.5 | 105.8 | - | 923 | 396.2 |
+| VLA-Adapter | 1 | 224 | 513.2 | 514.3 | 514.4 | 515.1 | 295.8 | 2865 | 1019.6 |
+| SmolVLA | 2 | 512 | 531.9 | 535.0 | 534.4 | 537.4 | 291.0 | 2023 | 414.9 |
+| Evo-1 | 1 | 448 | 1114.8 | 1117.0 | 1116.2 | 1119.9 | 507.3 | 1547 | 465.9 |
+| BitVLA | 1 | 224 | 1286.3 | 1287.5 | 1287.1 | 1288.5 | 167.9 | 1255 | 1129.7 |
+| VLA-JEPA | 1 | 256 | 2330.4 | 2345.8 | 2346.8 | 2355.1 | 106.2 | 3925 | 365.8 |
+
+### Fastest configuration
+
+We screen each model over these runtime flag sets, then re-run the fastest set
+under the full protocol:
+
+- the defaults;
+- `--flash-attn` and `--mm-prec default`, alone and together;
+- `--weight-dtype bf16` and `--weight-dtype f16`, each alone and with
+  `--flash-attn`;
+- `--act-dtype bf16 --flash-attn`, for Evo-1.
+
+The screen runs 2 warmups and 5 reps in one process per flag set. A flag set
+replaces the defaults only if its mean is more than 3% lower, both in the screen
+and in the full re-run. GR00T N1.5 fails at the defaults, so it has nothing to
+compare against: its row is the fastest screen configuration that works, with no
+speedup figure.
+
+Flash attention, bf16 activations and lower-precision weights can move the
+action chunk. Success rates measured at the defaults therefore do not carry over
+to these rows.
+
+| Model | Fastest flags | min ms | mean ms | p50 ms | p90 ms | vision ms | Device VRAM MiB | Host RSS MiB | vs defaults |
+|---|---|--:|--:|--:|--:|--:|--:|--:|--:|
+| Octo-Small | *(defaults)* | 23.2 | 24.6 | 24.1 | 26.4 | 2.8 | 627 | 400.4 | - |
+| ACT | *(defaults)* | 48.3 | 48.7 | 48.6 | 49.1 | - | 213 | 300.8 | - |
+| TurboVLA | *(defaults)* | 104.9 | 105.5 | 105.5 | 105.8 | - | 923 | 396.2 | - |
+| VLA-Adapter | *(defaults)* | 513.2 | 514.3 | 514.4 | 515.1 | 295.8 | 2865 | 1019.6 | - |
+| SmolVLA | *(defaults)* | 531.9 | 535.0 | 534.4 | 537.4 | 291.0 | 2023 | 414.9 | - |
+| Evo-1 | *(defaults)* | 1114.8 | 1117.0 | 1116.2 | 1119.9 | 507.3 | 1547 | 465.9 | - |
+| BitVLA | *(defaults)* | 1286.3 | 1287.5 | 1287.1 | 1288.5 | 167.9 | 1255 | 1129.7 | - |
+| VLA-JEPA | `--weight-dtype f16` | 1397.7 | 1402.3 | 1402.1 | 1405.7 | 406.0 | 3913 | 387.7 | -40.2% |
+| GR00T N1.5 | `--weight-dtype f16` | 2631.2 | 2640.9 | 2639.6 | 2648.2 | 565.9 | 3573 | 375.9 | - |
+
+<details><summary>Screen means (ms) per flag set</summary>
+
+| Model | defaults | `--flash-attn` | `--mm-prec default` | `--flash-attn --mm-prec default` | `--weight-dtype bf16` | `--weight-dtype bf16 --flash-attn` | `--weight-dtype f16` | `--weight-dtype f16 --flash-attn` | `--act-dtype bf16 --flash-attn` |
+|---|--:|--:|--:|--:|--:|--:|--:|--:|--:|
+| Octo-Small | 26.3 | 28.9 | 28.8 | 26.9 | 26.8 | 25.0 | 25.8 | 27.6 | - |
+| ACT | 52.6 | 50.0 | 55.2 | 55.5 | 55.2 | 52.2 | 80.0 | 80.2 | - |
+| TurboVLA | 107.4 | 105.4 | 105.4 | 105.2 | 108.0 | 108.4 | 349.4 | 349.8 | - |
+| VLA-Adapter | 511.1 | 512.5 | 512.8 | 513.5 | 513.5 | 513.9 | 1936.3 | 1936.1 | - |
+| SmolVLA | 530.9 | 535.9 | 573.9 | 579.0 | 574.9 | 579.6 | 1209.8 | 1235.2 | - |
+| Evo-1 | 1113.4 | 1102.6 | 1116.1 | 1103.5 | 1117.6 | 1105.0 | 4026.0 | 4031.9 | 1183.2 |
+| BitVLA | 1287.5 | 1287.4 | 1287.2 | 1288.4 | 1287.2 | 1287.1 | 1287.1 | 1287.3 | - |
+| VLA-JEPA | 2393.0 | 2032.4 | 2790.7 | 2093.4 | 2704.1 | 2147.8 | 1401.8 | 1408.1 | - |
+| GR00T N1.5 | FAIL | FAIL | FAIL | FAIL | FAIL | FAIL | 2652.8 | 2674.5 | - |
+
+</details>
+
+## Not run
+
+- **GR00T N1.5 at the defaults**: does not fit. It loads 3.31 GiB of BF16
+  weights, then a 64 MiB cuBLAS workspace request fails on a 16 MiB pool
+  extension. Turing has no fast BF16 path, so ggml promotes BF16 to F32 and needs
+  this buffer. Explicit F16 avoids it without shrinking the resident weights.
+
+  ```text
+  vla(gr00tn1d5): weights resident in 3.31 GiB (bf16)
+  ggml-cuda.cu:109: CUDA error
+  CUDA error: out of memory
+    current device: 0, in function alloc at ggml-cuda.cu:642
+    cuMemSetAccess(start_ptr, reserve_size, &access, 1)
+  ```
+
+- **GR00T N1.6**: not attempted. The GGUF is 9.16 GB, about 6 GiB resident
+  according to its model card, against 4 GiB here.
+- **GR00T N1.7**: not attempted. Its BF16 weights need 4.68 GiB, against 4 GiB
+  here.
+- **π0, π0.5, OpenVLA-OFT**: not attempted. Their GGUFs are 6.49, 12.59 and
+  15.07 GB, and on the RTX 3090 they peak at 5670, 5880 and 14768 MiB of device
+  memory ([rtx-3090.md](rtx-3090.md)), against 4096 MiB here.
+
+## Notes
+
+- The old `CUSTOM unsupported` failure came from fused kernels compiled for sm_80
+  only. The runtime now checks device and image compatibility before emitting
+  fused nodes, and sm_75 uses plain CUDA. BitVLA selects the sm_75 CUDA ladder path
+  for its int2 checkpoint, and its LM, ViT and action head activate CUDA before
+  prediction.
+- VLA-JEPA F16 changes its action chunk: maximum absolute difference 0.0762 and
+  maximum RMSE 0.0205 over three fixed inputs. Its speedup says nothing about task
+  accuracy.
+- Every selected configuration returned finite actions on those three inputs.
+
+## Reproducing
+
+```bash
+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DGGML_CUDA=ON \
+  -DCMAKE_CUDA_ARCHITECTURES=75 -DVLA_BUILD_TESTS=ON \
+  -DCMAKE_C_COMPILER=gcc-13 -DCMAKE_CXX_COMPILER=g++-13 \
+  -DCMAKE_EXE_LINKER_FLAGS=-L/usr/lib/gcc/x86_64-linux-gnu/15 \
+  -DCMAKE_SHARED_LINKER_FLAGS=-L/usr/lib/gcc/x86_64-linux-gnu/15
+cmake --build build --target vla-bench vla-server vla_predict_check -j 4
+
+# one row, for example SmolVLA
+./build/vla-bench --ckpt smolvla-libero.gguf \
+  --images 2 --size 512 --tokens 16 --warmup 3 --reps 20 --markdown
+```

src/cuda/vla_cuda_fused.cu

@@ -2298,8 +2298,44 @@ extern "C" bool vla_cuda_fused_forward(ggml_tensor * dst, void * stream_v) {
 
 namespace vla {
 
-void cuda_register_fused_ops() {
+bool cuda_register_fused_ops(const char * arch, int device) {
+    cudaDeviceProp prop{};
+    cudaError_t err = cudaGetDeviceProperties(&prop, device);
+    if (err != cudaSuccess) {
+        std::fprintf(stderr, "vla(%s): fused CUDA ops disabled: device query: %s (%d); using plain CUDA\n",
+                     arch, cudaGetErrorString(err), (int) err);
+        cudaGetLastError();
+        return false;
+    }
+    // This archive includes cp.async and TF32/BF16 MMA kernels requiring
+    // Ampere. A usable ggml CUDA backend alone does not imply their support.
+    if (prop.major < 8) {
+        std::printf("vla(%s): fused CUDA ops disabled: device %d sm_%d%d requires sm_80+; using plain CUDA\n",
+                    arch, device, prop.major, prop.minor);
+        return false;
+    }
+    // Also detect builds lacking a compatible image for this device, even
+    // when the hardware has the necessary instructions.
+    cudaFuncAttributes attr{};
+    int previous = 0;
+    err = cudaGetDevice(&previous);
+    if (err == cudaSuccess) {
+        err = cudaSetDevice(device);
+        if (err == cudaSuccess) {
+            err = cudaFuncGetAttributes(&attr, k_bias_res);
+            const cudaError_t restored = cudaSetDevice(previous);
+            if (err == cudaSuccess) err = restored;
+        }
+    }
+    if (err != cudaSuccess) {
+        std::fprintf(stderr,
+                     "vla(%s): fused CUDA ops disabled: device %d sm_%d%d kernel image: %s (%d); using plain CUDA\n",
+                     arch, device, prop.major, prop.minor, cudaGetErrorString(err), (int) err);
+        cudaGetLastError();
+        return false;
+    }
     cuda_ext_add_handler(vla_cuda_fused_forward);
+    return true;
 }
 
 }  // namespace vla