Repository navigation
FoldQuant W8A8/W4A4 inference and FoldQuantVLA quantized-model converter - #33
Merged
Merged
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.
hungho77
force-pushed
the
feat/quantized-inference
branch
from
September 30, 2026 07:34
999cb88 to
16c0ce1
Compare
main's performance review (#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.
- 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.
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
1 participant
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.
What
FoldQuant integer inference for GR00T N1.5 / N1.6 / N1.7 and π0.5, and a converter that turns a FoldQuantVLA quantized model into a FoldQuant GGUF.
Runtime
GGML_OP_CUSTOMnodes (src/layers/fq_linear.h):src/foldquant_ref.cpp.src/kernels/foldquant/, with a fused RMSNorm + butterfly + quant prologue, anmma.syncINT8/INT4 GEMM with per-shape tiling, and residual add and head layout fused into the epilogue. The CUDA ops are reached through one extension dispatcher (src/cuda/vla_cuda_ext.cu).WeightLoadergains typed optional and fused declares. A floatgemm()declare of an INT8 tensor fails with the FoldQuant site's name.docs/QUANTIZATION.md.scripts/inspect_gguf_quant.pychecks a file against it.Converter
scripts/convert_quantized_model_to_gguf.pyreads a FoldQuantVLA quantized model and writes the GGUF with no calibration and nothing re-rounded. It accepts both FoldQuantVLA formats:foldquant-quantized-checkpoint): the base checkpoint with.qweight/.weight_scalein place of each quantized.weight;foldquant-quant-state).--check-onnxbyte-compares every site and folded norm gain against the plugin ONNX graphs the TensorRT engines were built from.convert(ckpt, out, writer_factory=...), andscripts/gguf_quant_writer.pyturns their output into a FoldQuant file.convert_pi05_to_gguf.pyalso accepts an OpenPI checkpoint withoutconfig.json.scripts/foldquant_fake_export.py(uncalibrated file for bring-up),scripts/foldquant_ref.py(numpy reference),scripts/foldquant_dequant.py(dequantizing exporter for diagnostics).Why
The stock
Q8_0/Q4_0repack keeps activations in float and dequantizes insideggml_mul_mat, so it saves memory but not compute. FoldQuant runs the LLM and the action head on the integer tensor cores. The converter makes a model calibrated in FoldQuantVLA usable in vla.cpp without recalibrating, and keeps the GGUF bit-exact to the TensorRT deployment of the same arm.Verified
-Wall -Wextra(first-party code)ctestpassesvla_predict_checkdiff), or the change ismeant to move it and a LIBERO sweep is below
Archs and backends tested:
ctestpasses 11/11, including the FoldQuant CPU-op, CUDA-op and GEMM checks. Python unit tests pass (tests/py/test_quantized_model_converter.py,test_foldquant_ref.py,test_converters.py).--check-onnxfound all 304 sites and every folded gain byte-identical to the TensorRT graphs.Open loop against the dataset's actions (6 episodes, 200 steps, horizon 16, raw units):
The W8A8 GGUF was run on a real SO101 arm and behaved correctly.
--check-onnxbyte-identical to their TensorRT graphs.vla_predict_checkfloat diff. The branch touches shared float code (DiT head, Qwen3 LM, N1.7 adaLN precompute), so that box is left open until it is run.