Repository navigation
Merge main into backend/hexagon - #65
Open
thanhdq1o2 wants to merge 84 commits into
Open
thanhdq1o2 wants to merge 84 commits into
thanhdq1o2 wants to merge 84 commits into
Conversation
Runtime - A FoldQuant GGUF carries the GR00T N1.5 / N1.6 / N1.7 and pi0.5 language backbone and action module as INT8 or INT4 codes with per-row scales in a block-Hadamard, SmoothQuant-folded frame; activations are quantized per token and the projections run as integer GEMMs. The format and arithmetic are specified in docs/QUANTIZATION.md. - Each site is two GGML_OP_CUSTOM nodes (src/layers/fq_linear.h). The CPU backend runs the reference in src/foldquant_ref.cpp; CUDA runs src/kernels/foldquant/: a fused RMSNorm + butterfly + quant prologue and an mma.sync INT8/INT4 GEMM with per-shape tiling, residual add and head layout in the epilogue. The CUDA extension ops go through one dispatcher (src/cuda/vla_cuda_ext.cu). Other backends refuse a FoldQuant file at load. - WeightLoader gains typed optional / fused declares; a float gemm() declare of an INT8 tensor fails with the FoldQuant site's name. - GR00T N1.7 precomputes the DiT adaLN conditions per denoising step at load. Converter - scripts/convert_quantized_model_to_gguf.py converts a FoldQuantVLA quantized model, either its quantized checkpoint (.qweight / .weight_scale in place of each quantized .weight) or its earlier fake-quant state, with no calibration and nothing re-rounded. A W4A4 DiT's INT4 adaLN is written dequantized; every quantized projection must become a GGUF site or that adaLN. --check-onnx byte-compares every site and folded norm gain against the TensorRT plugin graphs. - The family converters expose convert(ckpt, out, writer_factory=...); scripts/gguf_quant_writer.py is the quantizing writer. convert_pi05_to_gguf.py accepts an OpenPI checkpoint without config.json. - Tooling: foldquant_fake_export.py (uncalibrated file for bring-up), inspect_gguf_quant.py (contract check), foldquant_ref.py (numpy reference), foldquant_dequant.py (dequantizing exporter). Tests: CPU-op, CUDA-op and GEMM checks for FoldQuant; Python tests for the reference and the converter.
-------- Co-authored-by: Khanh Nguyen <khanh.ng.18@gmail.com>
main's performance review (VinRobotics#32) moved pi0.5's Gemma layers into the shared modules/gemma_expert.h and precomputes the DiT adaLN modulation per denoising step (FlowTimes, ActionExpert::denoise). - GR00T: main's per-step adaLN precompute replaces this branch's adaLN cache (same idea); the FoldQuant LLM / DiT sites and the GEMM prefetch chain are kept on top. - pi0.5: the FoldQuant paths move into gemma_attn / gemma_mlp / gemma_layer. Prefix layers fuse the RMSNorm and folded gamma into the act node and the residual into the o / down epilogues; expert layers quantize the adaRMS output (gated residuals stay out of the epilogue). pi0 has no FoldQuant sites, so its float path is unchanged. - loader: main's resident-type fuse check, plus the same source type for a tensor copied raw (INT8 codes). - convert_pi05_to_gguf.py: convert() and the OpenPI config merge, plus main's QUANTILES-only check. - Docs: FoldQuant under [Unreleased] in the CHANGELOG, its section in docs/MODELS.md, docs/QUANTIZATION.md linked from the README.
…ernels Metal, SYCL, OpenVINO, Hexagon and OpenCL have no implementation of the two FoldQuant custom nodes and vla.cpp drives one backend with no per-op fallback, so a FoldQuant GGUF was refused there. Now foldquant_check_backend switches such a backend (or CUDA/CPU with VLA_FQ_DEQUANT=1) to dequant mode: every site is registered with the loader as a float GEMM weight, rebuilt at upload in the resident type, and the arch takes its stock float path. The activation path is x' = R(x / a) (fold before) or R(x) / a (after), R the block-normalised Sylvester-Hadamard butterfly (symmetric, orthonormal), so row n of the float weight is R(w_n) / a or R(w_n / a) (fq_dequant_rows). The LLM sites' SmoothQuant vector is already folded into the norm gains the file carries, which the float path reads as its norm weights. - WeightLoader::as_float registers a tensor of another file type as a float [K, N] weight; gemm/opt_gemm, fuse_gemm (DiT fused q/k/v, k/v) and upload honour it, converting to the resident type. - The four arch callers hold the spec mutable so the check can set the mode. - test_foldquant_dequant checks the rebuilt weights against the dense product (W8/W4, both fold orders, with and without ascale, a narrowed block). Weight-only quantization on those backends: weights keep FoldQuant's rounding, activations stay float. pi0.5 LIBERO W4A4 (vrfai/pi05-libero-w4a4): action cosine vs bf16 0.99986 dequant, 0.99703 integer path; dequant CUDA vs CPU 1.00000. GR00T N1.7 LIBERO W4A4: integer vs dequant 0.99930, dequant CUDA vs CPU 1.00000.
ggml's OpenVINO backend now runs a FoldQuant GGUF natively on the CPU and GPU plugins instead of reading the sites back as float weights. - src/openvino/foldquant_ov.cpp, compiled into ggml-openvino: fq_act becomes the CPU reference's arithmetic in OpenVINO ops (RMSNorm with the folded gamma, the ascale divide, the block rotation as a MatMul with the normalised Hadamard matrix, per-token amax scale, round-half-even, clamp) and passes codes and scale as floats to fq_gemm, which multiplies them with the weight kept as an i8 / i4 constant (W4 bytes reinterpreted in place), dequantized by wscale in the decompression pattern the plugins keep compressed. - patch_ggml_openvino.py hunk 14: register the translator for GGML_OP_CUSTOM, let supports_op accept FoldQuant's nodes (only those) ahead of the type checks, keep INT8 weights as i8 constants. - foldquant_check_backend: native on OpenVINO CPU/GPU, dequant on the NPU or with VLA_FQ_DEQUANT=1; the head-laid-out epilogue is off there (FqModuleSpec::no_heads), since a translated graph has no raw layout. - test_foldquant_ov_op: one site through the backend against the CPU reference (W8A8, W4A8, W4A4; gamma, ascale before/after, bias, narrowed and no rotation): every code agrees, outputs within 1e-6 relative, on an Intel Core Ultra X7 358H CPU and its Arc B390 iGPU. Whole model, against the CUDA integer path: W8A8 pi0.5 1.00000 (CPU), 0.99999 (GPU). W4A4 pi0.5 / GR00T N1.7 0.9993 / 0.9987 (CPU): with 15 activation levels a code on a rounding tie flips by one when the reductions run in a different float order, and the flips compound. pi0.5 W4A4 on the CPU plugin: 3926 ms vs 6595 ms bf16, 6.7 GB vs 16.5 GB peak RSS; on the iGPU 651 ms vs 560 ms, 3.8 GB vs 11.6 GB.
fq_act now emits the integer-valued codes already multiplied by their per-token scale, and fq_gemm feeds that straight into the MatMul with the decompressed weight. The scale factors out of the GEMM, so this drops the Concat, the two Slices and the scale Multiply at every site with the same arithmetic. Measured on an Intel Core Ultra X7 358H, pi0.5 LIBERO W4A4, against the codes+scale form: iGPU 666 -> 645 ms, CPU 4.69 -> 4.15 s (CPU timings on that machine drift by ~20% between runs). Accuracy against the CUDA integer path is unchanged: W8A8 1.00000 (CPU) / 0.99999 (GPU), W4A4 + o/down INT8 0.99962 / 0.99971, W4A4 0.99936 / 0.99928. test_foldquant_ov_op passes on the Intel CPU, the Arc B390 iGPU and OpenVINO's ARM CPU plugin. Tried and dropped: the plugins' dynamic activation quantization (on by default on the CPU, group 32, and the fastest setting there; no effect on the iGPU), and a group-quantized [N, K/g, g] weight layout (no faster on the iGPU). The iGPU stays behind bf16 because its compressed-weight matmul decompresses to F16 rather than running int8.
…sion hook
pi0.5 W4A4 on an Arc B390 iGPU runs in 291 ms (bf16: 326 ms), W4A4 with
INT8 o/down in 290 ms, W8A8 in 330 ms; every node bit-identical to the CPU
reference.
- scripts/patch_ggml_sycl_ext_hook.py: the SYCL counterpart of the CUDA hook,
two exported pointers (compute, supports_op for GGML_OP_CUSTOM) consulted
first by ggml-sycl; null by default, so a hooked ggml behaves as stock.
- src/sycl/vla_sycl_foldquant.cpp:
- fq_act follows the reference's reduction tree (lane l of a 32-wide
sub-group owns chunks l, l+32, ...). Rows up to K = 8192 spread over a
work-group (the CUDA act_row_kernel, ported); longer ones run one
sub-group per row.
- fq_gemm runs oneDNN's int8 matmul (INT4 weights as s4, INT4 activations
unpacked to s8) into an int32 buffer, then the reference epilogue.
Integer sums are exact in any order. Without oneDNN, or with
VLA_FQ_SYCL_GEMM=native, a GEMV (M <= 32) and an int8 joint_matrix (XMX)
kernel take over; =simple keeps a plain tiled one.
- Exactness: -ffp-contract=off for the source, and the device JIT option
-cl-fp32-correctly-rounded-divide-sqrt (a link option); without it the
driver's division and sqrt move codes across rounding ties. The CPU
reference gets -fp-model=precise under icpx, whose host default
(-fp-model=fast) reassociates its sums.
- foldquant_check_backend registers the kernels on a SYCL backend (plain
[N][T] outputs, so no head layout); VLA_FQ_DEQUANT=1 still selects dequant.
- test_foldquant_sycl_op: W8A8/W4A8/W4A4/W8A4 with gamma, ascale
before/after, bias, residual, narrowed and no rotation, over every GEMM
path and production shapes (K = 16384 / 6144 / 2048, M up to 300): 168
cases, activation bytes and outputs bit-identical.
- Docs: QUANTIZATION.md SYCL section, backend/sycl.md numbers and the Level
Zero loader workaround, MODELS.md, CHANGELOG, the foldquant.h overview.
Whole model vs the CUDA integer path: W8A8 1.00000 action cosine, W4A4 +
o/down INT8 0.99966, W4A4 0.99912 (pi0.5) / 0.9993 (GR00T N1.7); the gap is
the float layers between the sites landing codes on other rounding ties.
- build-sycl: oneAPI 2026.1 (DPC++, oneMKL, oneDNN) on ubuntu-24.04, full build with -Werror, then ctest on the OpenCL CPU device the compiler package ships. ggml-sycl will not start without a GPU, so test_foldquant_sycl_op now hands its two nodes straight to the extension hook on any SYCL device, tensors in USM shared memory, and checks them against the CPU reference as before (GPU: still through ggml-sycl). VLA_FQ_TEST_REQUIRE_DEVICE turns a missing device into a failure. - build-openvino: OpenVINO 2026.4 archive (cached), full build with -Werror, ctest; the CPU plugin runs test_foldquant_ov_op on the runner. - build-gate checks the SYCL hook patch still applies, and is idempotent. - icpx compiles first-party code with -fp-model=precise (its default is fast, which broke test_vision_common's exact float compare); ggml keeps its own flags. The intended -ffp-contract=off override is no longer warned about. - The oneDNN GEMM path falls back to the native kernels, once per queue, on a device oneDNN cannot drive.
The pi0.5 W8A8 rows came from scripts/foldquant_fake_export.py (rotation and per-row RTN, no calibration). Re-measured with FoldQuantVLA's W8A8 arm (w8a8_sr + w8a8_sh) calibrated on the same 128 LIBERO frames as vrfai/pi05-libero-w4a4, converted byte-identical to its TensorRT graphs: SYCL 305 ms (bf16 326), OpenVINO GPU 655 ms / 4.7 GB, CPU 3653 ms / 8.6 GB; actions vs the CUDA integer path 0.99999 (SYCL), 1.00000 (OpenVINO). The tables now name the checkpoints they were measured on.
On Linux gfx1151 with ROCm 7.14.60850 and llama.cpp b11223, CPU and HIP Release builds pass 8/8 CTest cases. SmolVLA and pi0.5 CPU BF16 actions remain byte-identical to upstream main@8eec763; HIP actions are finite and byte-identical across five fresh runs each. HIP-vs-CPU BF16 max abs is 1.992e-3 for SmolVLA and 8.792e-4 for pi0.5. An invalid HIP device ordinal falls back to CPU with identical action bytes. CMake rejects HIP with CUDA/Vulkan or dynamic backend loading; HIP builds exclude vla.cpp CUDA objects and NVIDIA runtime links. Synthetic-input HIP P50 is 395.0 ms for SmolVLA and 387.8 ms for pi0.5 under the recorded high clock setting. Other models and task success need validation on this upstream revision.
Limit the ROCm support matrix to SmolVLA and pi0.5 on Linux gfx1151, ROCm 7.14.60850 and llama.cpp b11223. CPU/HIP Release CTest passes 8/8; CPU actions match upstream main@8eec763 byte for byte, and five HIP runs per model are finite and identical. HIP-vs-CPU BF16 max abs is 1.992e-3 and 8.792e-4 respectively. Document GGUF SHA-256 values, build/runtime paths, fallback checks and three-round synthetic latency using stock vla-bench P50, P90 and Vision mean. Other architectures, server/client execution and robot-task success remain unverified on this revision.
ggml-hip also exports GGML_USE_CUDA. Define GGML_USE_HIP for vla_core and use it to select the HIP rung and exclude VLA CUDA-only BF16 code. Verified on gfx1151, ROCm 7.14.60850 and llama.cpp b11223: CPU and HIP Release rebuilds, 8/8 CTest each. SmolVLA and pi0.5 BF16 fixed-input action bytes match their pre-change CPU/HIP baselines; the HIP banner remains ROCm/HIP. Numeric output did not move. The prior CPU/HIP max-abs differences were 1.992e-3 and 8.792e-4. Performance was not remeasured for this macro-only change; the A9 3x20 engine baseline remains the reference.
- FoldQuant SYCL: oneDNN only on GPU devices. On the CI runner's OpenCL CPU device (an AMD EPYC, no VNNI) oneDNN's int8 matmul sums pairs in saturating 16-bit lanes and the outputs were off; the native kernels (and oneDNN with s4 weights) were exact there. vla.cpp runs ggml-sycl on GPUs only. - test_graph_names: an OpenVINO build already defines GGML_USE_OPENVINO; redefining it broke the -Werror build.
The CI runner's OpenCL CPU compiler crashed (free(): invalid pointer in sycl-kernel-reduce-cross-barrier-values) JIT-compiling the module, which holds every kernel, the XMX one included. With per-kernel images a device compiles only the kernels it launches. No change on the Arc GPU (W4A4 p50 291 ms); 3/3 CPU-device runs pass on the Intel PC. QUANTIZATION.md: W4A4 SYCL vs CUDA is 0.9995 now that host code is built with precise floating point under icpx (bf16 moved by 6e-4 with it).
The CUDA path's diagnostic, ported: after each node, recompute it with the CPU reference on host copies of the same inputs and report byte differences. On an Arc B390, GR00T N1.5 W4A4 (732 nodes) and N1.6 W4A4 (1376 nodes) are byte-identical node for node.
…erence FoldQuant W8A8/W4A4 inference and FoldQuantVLA quantized-model converter
Resolve the Unreleased changelog entries for ROCm/HIP and FoldQuant while retaining the combined CUDA/HIP guard in the registration header. Verified on gfx1151 with ROCm 7.14.60850 and llama.cpp b11223: CPU and HIP Release builds pass 10/10 CTest each. SmolVLA and pi05 BF16 fixed-input CPU/HIP actions remain byte-identical to the pre-merge branch; HIP-vs-CPU max abs errors are 1.992e-3 and 8.792e-4, respectively. Invalid HIP device falls back to CPU with identical actions. Three 20-rep HIP benchmark rounds per model show median P50 395.0/389.1 ms for SmolVLA/pi05. No CUDA-only VLA kernels or NVIDIA libraries are linked in the HIP build.
Document CPU/HIP tests and two-model fixed-input numerics on main 1adf078. Both Release builds pass 10/10 CTest on gfx1151 with ROCm 7.14.60850 and llama.cpp b11223. CPU actions match upstream main byte for byte; HIP actions are finite and repeat across five processes. HIP-vs-CPU max abs errors remain 1.992e-3 for SmolVLA and 8.792e-4 for pi05. Fresh 3x20 synthetic benchmarks report median P50 395.0/389.1 ms and worst P90 396.9/403.1 ms. Mark FoldQuant quantized GGUFs as unmeasured on HIP.
--------- Co-authored-by: Khanh Nguyen <khanh.ng.18@gmail.com>
New ggml-sycl ops behind an in-tree extension hook (src/sycl/vla_sycl_ext.cpp), built as vla_sycl_ops on oneDNN: - BF16 activations: BF16 matmul (BF16 or F32 result), norm, rms_norm, scale. - FLASH_ATTN_EXT on oneDNN's fused SDPA, operands read in place (F16/BF16, arbitrary strides, batched); default on SYCL for pi0 and pi0.5. - Fused elementwise nodes (layers/fused_ew.h): adaRMS, gated residual, geglu, LayerNorm, RoPE + K/V cache write, GEMM with bias/GELU/residual epilogue. - W8A8 linear with per-token INT8 activations (layers/qlinear.h), weights quantized at load; opt-in for pi0.5 (VLA_PI05_W8A8, VLA_PI05_W4). pi0.5 uses them for an in-place F16 K/V cache, fused projections and fused SigLIP blocks with all camera views in one pass. No CUDA, BitVLA or Evo-1 files change. Test: test_bf16_sycl_ops.
…buckets - patch_ggml_openvino.py: a graph fingerprint in front of the naive_compute cache. On a hit the cached decoder is already bound to these tensors, so naive_key and update_io() are skipped. GGML_OPENVINO_EXEC_MODE=performance lets a model opt out of the ACCURACY hint that pins non-MUL_MAT graphs to f32. - pi0.5 on OpenVINO: the language segment is padded to 16-token buckets with the pad masked as keys, so a prompt whose length changes every step compiles one model per bucket instead of one per length; performance exec mode.
GGML_OP_CUSTOM nodes (layers/fused_ops.h, cuda/vla_cuda_fused.cu) that replace chains of ggml ops with one kernel each and give the same bits: LayerNorm / RMSNorm with affine or modulation, bias + GELU / GELU (erf), bias + residual with or without a gate, SiLU / GELU gates, Q/K/V split with RoPE into the attention layout, patchify. gemm_exact (mm_exact): a BF16-weight GEMM on the fastest cuBLASLt algorithm or gemm_skinny kernel that reproduces ggml's cublasGemmEx call bit for bit, checked per shape at first use on two random operand sets; otherwise that call itself. Only where ggml uses cuBLAS (> 16 rows). attn_exact: attention as ggml's chain (two cuBLAS GEMMs, soft_max_f32, permute) with the softmax emulated in one kernel, or a fused FFMA or TF32 kernel where it reproduces that chain (checked per shape). Nothing uses these yet except mm_act once a model turns the nodes on. tests/test_fused_cuda_ops compares every node with the ggml chain it replaces, bitwise. Verified: test_fused_cuda_ops on Jetson AGX Orin (CUDA 12.6, sm_87).
…odes EncStack (SigLIP and GR00T's VL self-attention blocks), the Qwen3 LM, the Qwen3-VL tower and the DiT head take the fused nodes when the model turns them on. EncStack keeps ggml's graph for several views: ggml multiplies 3D activations with one batched cuBLAS call, which one GEMM over all views does not reproduce. GR00T N1.7 also runs its vision tower inside the main graph (all views, patchify on the device), so the tower outputs never leave the device. VLA_GR00T_FUSED_VIT=0 restores the separate vision graph. rope_2d takes a view dimension; with one view it emits the same nodes. Verified on Jetson AGX Orin: bit-identical actions for every model on synthetic requests, and on requests recorded from LIBERO-Object runs (GR00T N1.5, N1.6 with 2 views, N1.7 with 1 and 2 views).
The Gemma prefix and action expert run RMSNorm (pi0.5: adaRMS) into the BF16 GEMM operand, bias-free NEOX RoPE + head layout with the K/V cache, exact attention with MQA and the suffix mask, GEGLU into the BF16 operand, and pi0.5's gated residuals in one kernel. In F32 the SigLIP layers are EncStack's block op for op, so the fused blocks stand in. Verified on Jetson AGX Orin: bit-identical actions on synthetic requests.
SmolVLA's SigLIP layers are EncStack's block op for op and run on the fused blocks (1024-token attention on the TF32 kernel). Evo-1: InternViT (fused QKV split, bias + layer scale + residual, bias + GELU erf), the Qwen2 LM (GQA, padding mask, SiLU gate) and the DiT. VLA_FUSED_OPS=vit, lm or dit limits Evo-1's fused paths to one part. Verified on Jetson AGX Orin: bit-identical actions on synthetic requests and on requests recorded from LIBERO-Object runs (Evo-1 with 3 views).
DINOv2 + SigLIP towers and projector (OpenVLA-OFT, VLA-Adapter); the Llama-2 and Qwen2.5 LMs (RMSNorm into the BF16 operand, bias + NEOX RoPE, exact attention, VLA-Adapter with its causal mask, SiLU gate). BF16 weights only. TurboVLA: DINOv3 layers (bias + half-split RoPE + head layout in one pass; the rotation takes sin itself, uploaded next to the signed table), exact attention over the views, and fused LayerNorm and residual epilogues in the fusion layers, text enhancer and ACT decoder. Octo: the block transformer (with its additive mask) and the diffusion head's LayerNorm and residuals. Verified on Jetson AGX Orin: bit-identical actions on synthetic requests, and on requests recorded from LIBERO-Object runs for OpenVLA-OFT, VLA-Adapter and TurboVLA. Octo's recorded requests do not replay deterministically on the original either, so only its synthetic set is checked.
The ternary GEMM runs on mma.sync.m16n8k32 (int8 x the 2-bit codes as
uint8, -2*rowsum(A) applied once per row), with a cp.async pipeline and
the weights repacked into the fragment layout at load. Products are int8
x {-1, 0, 1} summed in int32, exact in any order; the epilogue (scales,
bias, gelu_tanh, residual, head layout, repeat_kv) keeps the original
kernels' bf16 roundings. q/k/v run in one launch. Softmax, the norms with
the act_quant that follows them, and RoPE with repeat_kv are rewritten
with the original kernels' reduction order and FMA contractions. Image
embeddings stay on the device for the LM input.
VLA_BITVLA_MMA=0 runs the original kernels; VLA_BITVLA_DEVICE_EMBEDS=0
copies the embeddings through the host.
Verified on Jetson AGX Orin: bit-identical actions on synthetic requests
and on requests recorded from LIBERO-Object runs; per-layer LM dumps
identical to the original kernels.
docs/backend/cuda-exact.md describes the fused nodes, gemm_exact and attn_exact, how they keep ggml's output, what each arch runs fused, the switches and how to verify. ARCHITECTURE, USAGE and the README link it. The AGX Orin report is re-run at e6653e8 with its protocol and compared with 1012c05 built the same way. getrusage does not see the CUDA buffers on this board, so the memory column is now the peak cudaMalloc + cuMemCreate total, counted with an LD_PRELOAD shim. Verified: vla_predict_check gives identical actions at 1012c05 and e6653e8 for every arch, 20 view configurations; ctest passes.
…sed ops VLA_GR00T_FUSED_VIT defaulted on even with VLA_FUSED_OPS=0 or off CUDA. The batched plain ggml graph is not bit-identical to the per-view one, so with two or more views the fallback differed from main. It now defaults to the vision fused ops being active. Checked on Jetson AGX Orin: vla_predict_check matches 1012c05 for 1, 2 and 3 views at the defaults, VLA_FUSED_OPS=0/vit/dit and VLA_GR00T_FUSED_VIT=0; ctest 15/15; latency unchanged.
…on the GPU pi0.5 writes the state into the prompt, so most calls bring a new prompt length and every GEMM over it is chosen again. That took about 2.3 s per new length on Orin. - A cuBLASLt pick for the same N, K and views within 64 rows is tried first and kept untimed when it is bit-exact on both operand sets. VLA_GEMM_REUSE=0 turns this off. - The random test operands come from a jump-ahead LCG kernel. The values match the host loop byte for byte, so the selection itself is unchanged. New-length call on pi0.5 SO101 bf16: 2596 ms -> 555 ms p50. Actions stay bit-identical to 1012c05 (80/80 over a prompt-length sweep), verify_all passes, steady-state latencies unchanged.
…UDA path A Gemma layer whose sites are all FoldQuant now takes gemma_attn_fused and gemma_mlp_fused: q/k/v come from the FoldQuant GEMMs, the attention core, RoPE, GEGLU, adaRMS and the gated residuals use the exact fused nodes, and the o / down sites keep their act nodes and residual epilogues. Actions unchanged on the pi0.5 SO101 w4a4 / w4a4od8 / w8a8 files. Orin, 2 views 224, 48 tokens: w4a4 214.9 -> 177.3 ms, w4a4od8 235.0 -> 197.8 ms, w8a8 265.6 -> 228.0 ms.
…=16384 Rows past 8192 (the pi0.5 prefix down input) used the warp-per-row kernel. They now take act_row_kernel with 16 warps a row; the values are formed in the same order, so the blobs are unchanged (test_foldquant_cuda_op passes). M=528, K=16384 alone on Orin: A8 660 -> 319 us, A4 504 -> 308 us.
gelu_mul_act (layers/fused_ops.h, built by fq_act_geglu in layers/fq_linear.h) computes gelu(g) * u as k_gelu_mul does and quantizes the product as the site's act node would, so the [T, ffn] product is never written out and read back. The act chain is spelled with IEEE ops without flush-to-zero, since vla_cuda_fused.cu builds with -use_fast_math while the FoldQuant kernels build with -fmad=false. Used by the pi0.5 Gemma MLPs; VLA_FQ_NO_ACT_FUSE=1 keeps the two nodes. test_fused_cuda_ops compares the blob bytes against geglu + fq_act for A8 and A4, rotation blocks 1/32/64 and ascale before or after the butterfly. pi0.5 SO101 actions unchanged; Orin, 2 views 224, 48 tokens: w4a4 177.7 -> 164.4 ms, w8a8 226.3 -> 212.4 ms.
Two mma tilings join the candidates: 64x128x128 over 8 warps (variant 17) and 128x128x128 with 64x64 warp tiles, 2 stages (variant 18). Prefill sites (M > 64) with N a multiple of 128 take 17, and W8 sites past 512 rows take 18; the rest keep 128x64x128. The accumulation is integer, so the outputs are unchanged (test_foldquant_cuda_op passes). Measured in the models on Orin (prefill GEMM time): pi0.5 2 views W8 73 -> 54 ms, W4 37 -> 32 ms; pi0.5 1 view W8 67 -> 61 ms, W4 39 -> 36 ms; GR00T N1.7 W8A8 / W4A4 at 1-2 views unchanged or up to 1 ms faster.
…pe on FQ Qwen3 layers glu_act generalizes gelu_mul_act to gelu(g) * u, silu(g) * u and gelu(g), each quantized as the act node of the site it feeds. Qwen3 FQ FFNs use silu_mul, the DiT FQ FFN uses gelu; the DiT plain LayerNorm and the Qwen3 q/k norm + IMROPE use the existing fused nodes on FQ layers. GR00T N1.7 FQ (W8A8, W4A4) actions are bit-identical to upstream; 2-view latency drops 95.8 -> 91.4 ms (W8) and 79.9 -> 76.2 ms (W4).
The cuBLAS row threshold only matters for float layers, where ggml picks a different BF16 GEMM below 17 rows. FoldQuant layers run the in-tree INT GEMMs, so the pi0.5 LIBERO expert (10 rows) now takes the fused path. Bit-identical: pi0.5 LIBERO w4a4 152.8 -> 144.9 ms on AGX Orin.
…xpert adaRMS norm_act runs the norm feeding a FoldQuant site inside its act node, with ggml's 1024-thread reduction order, so the normed rows are never written out. The N1.7 cross-attention q site writes its output in head layout. Both are bit-identical (test_fused_cuda_ops, verify_all, action md5 of every calibrated and SO101 checkpoint). AGX Orin, w4a4: N1.7 58.0 -> 55.6 ms, N1.6 73.1 -> 71.9, N1.5 68.3 -> 67.5, pi0.5 145.0 -> 142.7.
…dquant pi0.5 SYCL: FoldQuant fast path on main's SYCL stack; cached vision graph; vla_encode_images
vla_cuda_fused.cu compiles with -use_fast_math; device-linked together with
the other vla_cuda_ops sources, CUDA 12.9's nvlink rejects the mix
("Cicc option values for '-ftz' do not match"), as it did for bitvla. Same
compile flags, so the same code.
- GEMM/attention scratch (path B, BF16 staging, cuBLASLt workspace) is one buffer per (device, stream) and slot instead of one per shape, found with the stream's handles through a thread-local lookup; plans are per device. pi0.5 with a prompt length that changes every call: +447 MiB -> +29 MiB over 21 lengths. - Attention tries a nearby shape's path A/C choice, checked bit-exact against path B, without timing, as GEMMs do; the bit-exact checks compare on the device. New-length request 445 -> 357 ms (repeat 298). - Without a plan under graph capture, attn_exact runs path B. - BitVLA: a freed model's weights leave the repacked set, so a reload in the same process repacks again (it returned wrong actions before). - Fused nodes only for weights in a CUDA buffer (fused_ops_for); the FoldQuant glu_act in the DiT and Qwen3 follows cfg.fused_ops. - bf16_operand with the operand rows in the DiT, Qwen3 and vision towers. - mm_exact leaves batched operands to ggml_mul_mat. - gemma_attn_fused: k_out holds the new tokens only when a cache is given. - Evo-1: VLA_FUSED_OPS parts matched exactly. Actions unchanged: bit-identical to 1012c05 on all archs (1/2/3 views), verify_all, test_fused_cuda_ops (36 cases), FoldQuant md5. Fixed-shape latency unchanged (A/B against the previous head).
…nels CUDA: bit-exact fused kernels for BF16 and FoldQuant, 23-40% faster on Jetson AGX Orin
Add AMD ROCm/HIP backend
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Sync only, no new code: brings
backend/hexagonup tomain(aeef71b) so Hexagon work builds on #30 (Snapdragon X support, CPU fallback wrapper) and #32.docs/backend/hexagon.md(add/add), resolved with main's rewritten version.main(git diff origin/mainis empty).-DVLA_BUILD_TESTS=ON -DVLA_BUILD_SERVER=OFF -DVLA_SPM=OFF) passes,ctest10/10.