Add NVFP4 CPU kernels results, integrate model card, add handoff

README.md

@@ -12,24 +12,28 @@ tags:
 - amx
 - moe
 - flint
+- nvfp4
 ---
 
-# flint CPU optimization (llama.cpp patch + two GGUF files)
+# flint CPU optimization (llama.cpp patch + GGUF files)
 
 [Recursive Flint](https://huggingface.co/RecursiveSI/flint) is a 39.8B-parameter mixture-of-experts coding model with about 2.5B
 active parameters per token (256 experts, top-6, 49 layers). This repo holds a **llama.cpp CPU patch** that speeds it up on
-Intel Xeon (AVX512-VNNI / AMX) and the two GGUF files it was tuned and measured on. **The GGUF files are unmodified**; every gain
+Intel Xeon (AVX512-VNNI / AMX) and the GGUF files it was tuned and measured on. **The GGUF files are unmodified**; every gain
 comes from code (`patches/flint-cpu-opt.patch`, base commit `de5122d87`). It is in the style of
 [HelloSun/Qwen3.6-35B-A3B-GGUF_CPU_OPT](https://huggingface.co/HelloSun/Qwen3.6-35B-A3B-GGUF_CPU_OPT).
 
-| file | source | quant | size |
-|---|---|---|---|
-| `flint-UD-Q4_K_XL.gguf` | [unsloth/flint-GGUF](https://huggingface.co/unsloth/flint-GGUF) | Q5_K 11.6 GiB + IQ4_XS 9.0 GiB + Q8_0 2.0 GiB (+ small Q4_K/Q6_K/IQ3_S) | 23.3 GiB |
-| `flint-int4-q4_1.gguf` | [RecursiveSI/flint-int4-gguf](https://huggingface.co/RecursiveSI/flint-int4-gguf) | Q4_1 for every weight (token_embd and router F32) | 24.2 GiB |
+| file | source | quant | size | in this repo |
+|---|---|---|---|---|
+| `flint-UD-Q4_K_XL.gguf` | [unsloth/flint-GGUF](https://huggingface.co/unsloth/flint-GGUF) | Q5_K 11.6 GiB + IQ4_XS 9.0 GiB + Q8_0 2.0 GiB (+ small Q4_K/Q6_K/IQ3_S) | 23.3 GiB | yes |
+| `flint-int4-q4_1.gguf` | [RecursiveSI/flint-int4-gguf](https://huggingface.co/RecursiveSI/flint-int4-gguf) | Q4_1 for every weight (token_embd and router F32) | 24.2 GiB | yes |
+| `flint-nvfp4.gguf` | [RecursiveSI/flint-nvfp4-gguf](https://huggingface.co/RecursiveSI/flint-nvfp4-gguf) | NVFP4 weights (485 tensors, 20.7 GiB), F32 per-tensor `.scale`, BF16 token_embd | 22.9 GB | **no**, download from the link |
 
 ## Which file to use
 
-Measured on Intel Xeon Platinum 8559C, 16 threads, with the patch (median of 3 interleaved rounds, tokens/s):
+Measured on Intel Xeon Platinum 8559C, 16 threads, with the patch (median of 3 interleaved rounds, tokens/s).
+Q4_K_XL and Q4_1 rows are from the earlier measurement on a quieter host; the NVFP4 rows are from a later run on a busier
+host (see "Measurement notes"), so compare NVFP4 only with the NVFP4 rows.
 
 | file | prefill pp256 | prefill pp2048 | decode tg128 | PPL (wikitext-2, 32x512) |
 |---|---|---|---|---|
@@ -37,11 +41,15 @@ Measured on Intel Xeon Platinum 8559C, 16 threads, with the patch (median of 3 i
 | **Q4_K_XL, patched** | **296** | **307** | 42.8 | 18.277 +- 0.787 |
 | Q4_1, stock llama.cpp | 183 | 189 | 48.3 | 18.260 +- 0.780 |
 | **Q4_1, patched** | 266 | 297 | **55.7** | 18.240 +- 0.779 |
+| NVFP4, stock llama.cpp | 83.7 | 83.6 | 30.6 | 24.417 +- 1.566 (wikitext-2, 16x512) |
+| **NVFP4, patched** | **178.5** | **211.1** | 32.1 | 24.414 +- 1.568 (wikitext-2, 16x512) |
 
-- **Long prompts / large contexts (prefill-heavy): Q4_K_XL** still prefills a little faster, but patched Q4_1 is now within ~10 % (it uses the tiled VNNI matmul too).
-- **Chat and agent loops with a cached prefix (decode-heavy): Q4_1**. It reads fewer bytes per weight (0.625) and decodes ~30 % faster.
-- Quality is equivalent within measurement error: the PPL differences above are smaller than the +-0.8 standard error of a 32-chunk test.
-- Single runs on this shared host vary by about +-8 %; the numbers are medians, raw data in `benchmarks/`.
+- **Long prompts / large contexts (prefill-heavy): Q4_K_XL** still prefills fastest of the three in the earlier measurement,
+  and patched Q4_1 is within about 10 %.
+- **Chat and agent loops with a cached prefix (decode-heavy): Q4_1**. It decodes about 30 % faster than Q4_K_XL.
+- **NVFP4 on CPU**: prefill is about 2.1x (pp256) to 2.5x (pp2048) faster than stock after the patch. Decode is unchanged
+  within noise. NVFP4 is primarily intended for CUDA; on this CPU it is still slower than Q4_1 for both phases.
+- Quality is equivalent within measurement error: PPL differences are much smaller than the standard errors.
 
 ## Gain of the patch (vs stock llama.cpp, same file)
 
@@ -49,19 +57,24 @@ Measured on Intel Xeon Platinum 8559C, 16 threads, with the patch (median of 3 i
 |---|---|---|---|
 | Q4_K_XL | 1.28x | 1.12x | 1.12x |
 | Q4_1 | 1.46x | 1.57x | 1.15x |
+| NVFP4 | 2.13x | 2.52x | 1.05x (noise) |
 
-For Q4_K_XL most of this comes from the Qwen3.6 patch kernels (Q4_K/Q5_K/Q6_K VNNI, AMX tweaks, MoE tile grouping) also
-applying to flint; against the Qwen3.6 patch alone the flint-specific part is small (pp256 +12 %, decode +2 %, pp2048 equal).
+## What the patch changes
+
+For Q4_K_XL most of the gain comes from the Qwen3.6 patch kernels (Q4_K/Q5_K/Q6_K VNNI, AMX tweaks, MoE tile grouping) also
+applying to flint. Against the Qwen3.6 patch alone the flint-specific part is small (pp256 +12 %, decode +2 %, pp2048 equal).
 The flint-specific changes:
 
 1. **IQ4_XS x Q8_K `vec_dot` in AVX512-VNNI** (x86 had AVX2 only). Unsigned LUT (+128) with `vpdpbusd`, bias removed through `bsums`, L2 prefetch.
 2. **Q4_1 x Q8_1 `vec_dot` in AVX-512VL VNNI**. Two blocks per iteration, packed into one ymm with memory-source `vinserti128`, low/high nibbles each through one `vpdpbusd`, fp16 headers converted together. A 512-bit version using cross-lane shuffles was slower than stock and dropped.
 3. **Q4_1 in the tiled VNNI GEMM** for prefill. Q4_1 has no super-block scale, so the tiles gain per-32-block f32 scale/min arrays (`sd`/`sm`) and a `Q41` kernel flag; activations reuse the existing Q8_1 `vec_dot_type`. Per 32-block: `d_a * (d_w * dot(q, code) + m_w * sum(code))`, integer dot/sum exact. This lifts pp256/pp2048 by ~34 %/42 % over the `vec_dot` path.
 4. Q5_K 5th-bit merge with `vptestmb` + masked add (bit-identical to the previous version, no speed change).
+5. **NVFP4 x Q8_0 `vec_dot` in AVX512-VNNI** (decode). Codes `|w|` (0..12) form the unsigned operand, the weight sign is moved onto the int8 activations. Each 64-element block has four 16-element subblocks with UE4M3 scales.
+6. **NVFP4 in the tiled VNNI GEMM** (prefill). Weights are stored as `LUT + 12` (unsigned 0..24) with one f32 scale per 16-wide subblock (`sd16`). The bias is removed exactly through the per-16 activation code sums. Activations are Q8_0 with one scale per 32-block. Enabled only in VNNI builds.
 
 ## Accuracy
 
-Wikitext-2 test, 32 chunks x 512, KL divergence against the stated baseline logits of the same file:
+Wikitext-2 test, KL divergence against the stated baseline logits of the same file (Q4_K_XL / Q4_1):
 
 | comparison | KLD | same top-1 | noise floor: stock vs stock with `-fa off` |
 |---|---|---|---|
@@ -70,22 +83,19 @@ Wikitext-2 test, 32 chunks x 512, KL divergence against the stated baseline logi
 | Q4_1, tiled GEMM vs `vec_dot` patch | 0.0372 +- 0.0026 | 94.2 % | 0.0385 |
 
 The KLD is **inside the noise floor of stock itself**: flint is very sensitive to tiny numeric differences (changing only the
-attention path of stock gives the same KLD), so it does not indicate a defect. Supporting checks: a graph-level `mul_mat` /
-`mul_mat_id` test gives the same relative L2 error versus a float reference as stock (Q4_K/Q5_K/Q6_K/IQ4_XS ~7e-3, Q4_1 ~5.4e-3, Q8_0
-~5.4e-3; the tiled Q4_1 path gives the same 5.29e-3/5.35e-3 as the Q4_1 `vec_dot` path). Integer dot products are exact; only
-float accumulation order changes. Why small differences get amplified was not verified (top-6-of-256 expert routing ties is a
-guess). The earlier version of this repo reported KLD 0.2058 for a different patch.
-
-## Where the time goes (why further kernel work is limited)
-
-Measured DRAM read bandwidth on this host is 125-175 GB/s with 16 threads (single thread 15 GB/s). Decode reads about 2 GB of
-weights per token (~1.6 GB dense Q8_0 attention/shared-expert/lm_head, ~0.5 GB for the 6 selected experts), i.e. ~90 GB/s effective:
-mostly memory-bound. Per-type effective bandwidth during decode (Q4_K_XL): Q8_0 dense ~100-170 GB/s (lm_head ~170), IQ4_XS experts ~100,
-Q5_K experts ~50. The Q5_K gap is not kernel throughput (cache-resident Q5_K and IQ4_XS run the same, 232 vs 218 GB/s); a
-prefetch-all-experts experiment did not help and was dropped; the cause was not found. Prefill is compute-bound; once Q4_1 is
-in the tiled GEMM, FlashAttention and the tiled MoE matmul are the largest items (the per-row Q4_1 `vec_dot` remains only on the
-decode path). The Q4_1 tiled kernel carries per-32-block f32 scales/mins in the tiles (`sd`/`sm`) on top of the integer
-sub-block layout used by the K-quants.
+attention path of stock gives the same KLD), so it does not indicate a defect.
+
+NVFP4 (not KLD-tested): perplexity on wikitext-2 test, 16 chunks x 512 tokens, `llama-perplexity`: stock 24.4165 +- 1.566,
+patched 24.4135 +- 1.568. Kernel-level checks: AVX512 `vec_dot` vs generic reference, relative error <= 1.3e-6; tiled vs
+`vec_dot`, relative L2 1.4e-7 .. 2.5e-7 over several shapes including ragged row counts and the narrow path.
+Integer dot products are exact; only float accumulation order changes.
+
+## Where the time goes
+
+Measured DRAM read bandwidth on this host is 125-175 GB/s with 16 threads (single thread 15 GB/s). Decode of the Q4_K_XL / Q4_1
+files is mostly memory-bound (about 90 GB/s effective). Prefill is compute-bound. For NVFP4, `GGML_PROF` per-op timing showed
+that prefill `MUL_MAT` / `MUL_MAT_ID` with 256 rows dominated the run time before the tiled kernel was added; decode spends its
+time spread over many small ops.
 
 ## Build and run
 
@@ -94,19 +104,24 @@ sub-block layout used by the K-quants.
 taskset -c 0-15 llama.cpp/build/bin/llama-cli -m flint-UD-Q4_K_XL.gguf -t 16 -fa on -lm mlock
 ```
 
+For NVFP4, download `flint-nvfp4.gguf` from RecursiveSI and run the same binary with `-m flint-nvfp4.gguf`.
+
 `GGML_IQ4XS_PF`, `GGML_Q41_PF`, `GGML_K4_PF`, `GGML_K5_PF`, `GGML_Q6_PF` tune the prefetch distances (defaults are the tuned
-values). `patches/optional-profiler.patch` adds `GGML_PROF=1` per-op timing.
+values). `patches/optional-profiler.patch` adds `GGML_PROF=1` per-op timing. `GGML_CPU_TILED_MM=0` disables the tiled GEMM.
 
-## Limitations
+## Measurement notes
 
 - Validated only on Intel Xeon Platinum 8559C (AVX512-VNNI + AMX), 16 threads, on a shared host. Other CPUs are untested; the new kernels need AVX512-VNNI (the Q4_1 one also AVX512VL/F16C).
-- The Q4_K_XL numbers were measured before the Q4_1 kernels were added; the K-quant code path is unchanged except for removing unused constants, and was not re-benchmarked (the Q4_1 tiled kernel was measured after).
-- PPL and KLD use a short test (32 x 512 tokens); no downstream task evaluation was run.
+- Single runs on this shared host vary widely. The NVFP4 round-to-round spread was up to 2x for decode while the host load average was 55-63. Use medians of interleaved rounds; the raw data is in `benchmarks/nvfp4/`.
+- The Q4_K_XL numbers were measured before the Q4_1 kernels were added; the K-quant code path is unchanged except for removing unused constants, and was not re-benchmarked.
+- PPL and KLD use short tests (16 or 32 chunks of 512 tokens); no downstream task evaluation was run.
 
 ## Files
 
 - `patches/flint-cpu-opt.patch` (base `patches/BASE_COMMIT`), `patches/optional-profiler.patch`, `build.sh`, `machine.txt`
 - `benchmarks/summary.md` (Q4_K_XL), `benchmarks/q41-summary.md` (Q4_1), raw A/B CSVs (`ab-interleaved-flint.csv`, `q41-ab-interleaved.csv`, `q41-tiled-ab-interleaved.csv`), KLD logs (`kld_vs_stock.log`, `q41-kld.log`, `q41-tiled-kld.log`), op profiles, KLD bisect, `graph_accuracy_test.cpp`
+- `benchmarks/nvfp4/` (NVFP4 summary `README.md`, raw llama-bench CSV per round `bin_*_r*.csv`, flat `summary.csv`, PPL logs `ppl_*.log`, run script, kernel test sources)
+- `handoff/HANDOFF.md`: state of the work and next steps for the next agent
 
 ## Credits and license
 

benchmarks/nvfp4/README.md

@@ -0,0 +1,62 @@
+# RecursiveSI/flint-nvfp4-gguf (flint-nvfp4.gguf, 22.9 GB) with the CPU patch
+
+NVFP4 weights (485 tensors, 20.7 GiB) plus F32 per-tensor `.scale` tensors and BF16 token_embd. Only the NVFP4 file was
+measured here; the Q4_1 and Q4_K_XL numbers in `README.md` / `q41-summary.md` were not re-run.
+
+## What was added to the patch (`patches/flint-cpu-opt.patch`)
+
+1. **AVX512-VNNI `vec_dot` for NVFP4 x Q8_0** (decode, `ggml/src/ggml-cpu/arch/x86/quants.c`). Weight codes are
+   `|w|` (0..12, unsigned operand of `vpdpbusd`), the sign of `w` is moved onto the int8 activations
+   (`_mm256_sign_epi8`). The 64-element block is split into 4 subblocks of 16 with per-subblock UE4M3 scales.
+2. **Tiled VNNI GEMM for NVFP4** (prefill, `ggml/src/ggml-cpu/tiled/`). Same framework as the Q4_1 path. Weights are
+   stored as `LUT + 12` (unsigned, 0..24) with one f32 scale per 16-wide subblock (`sd16`). The bias is removed exactly
+   with the per-16 activation code sums (`bsums`): `sum(w*a) = sum((w+12)*a) - 12*sum(a)`. Activations use Q8_0 with
+   one scale per 32-block. `tiled_unpk_fp4_64` is an AVX2 unpacker (scalar fallback).
+   Enabled only in VNNI builds (`__AVX512VNNI__ && __AVX512VL__`).
+
+Integer dot products are exact; only the float accumulation order changes.
+
+## Results
+
+Xeon Platinum 8559C, `taskset -c 0-15`, `-t 16 -fa 1 -b 1024 -ub 1024 -lm mmap+mlock -p 256,2048 -n 128 -r 3`.
+Three interleaved rounds per build, median of the per-round means (tokens/s). Raw CSV: `bin_*_r{1,2,3}.csv`,
+flat list: `summary.csv`.
+
+| build | pp256 | pp2048 | tg128 |
+|---|---|---|---|
+| **before**: stock llama.cpp (de5122d87), no patch | 83.7 | 83.6 | 30.6 |
+| patch (Q4_1 kernels), NVFP4 path unchanged | 82.3 | 84.4 | 31.8 |
+| **after**: patch + NVFP4 VNNI vec_dot + tiled VNNI GEMM | **178.5** | **211.1** | 32.1 |
+
+| gain | pp256 | pp2048 | tg128 |
+|---|---|---|---|
+| after vs stock | 2.13x | 2.52x | 1.05x (noise) |
+| after vs patch-before | 2.17x | 2.50x | 1.01x (noise) |
+
+Per-round values: pp256 after = 153.1 / 178.5 / 181.6, pp2048 after = 206.7 / 218.6 / 211.1, tg128 after = 32.1 / 29.8 / 32.4.
+
+Decode (tg128) did not change measurably. The VNNI vec_dot was about 1 % faster in an earlier single-run check, which is within
+noise. Per-op profiling (`GGML_PROF`) showed that prefill `MUL_MAT`/`MUL_MAT_ID` with 256 rows dominated the run time.
+
+## Correctness
+
+- `vec_dot` AVX512 vs generic reference (random blocks, n = 1920..12352): relative error <= 1.3e-6 (float ordering only).
+- Dense `mul_mat`, tiled (M=64) vs vec_dot (M=1): relative L2 1.4e-7 .. 2.5e-7 for K in {1024, 2048, 4096, 8192} and N in {8, 16, 24, 100, 256, 512}
+  (covers ragged row counts and the narrow path). Versus f32 reference: about 3.7e-3 (same order as Q4_1 in the HANDOFF).
+- wikitext-2 test, 16 chunks x 512, llama-perplexity:
+
+| build | PPL |
+|---|---|
+| stock | 24.4165 +- 1.566 |
+| patch + NVFP4 opt | 24.4135 +- 1.568 |
+
+## Notes
+
+- The host was shared and heavily loaded during these runs (load average 55..63). Single runs ranged widely
+  (stock tg128: 14.1 / 32.3 / 30.6), so use the medians. Absolute numbers are lower than the Q4_1 reference figures in
+  `q41-summary.md`, which were measured on a quieter host. Compare only within this table.
+- After the patch, NVFP4 prefill (pp2048 211) is still below Q4_1 (pp2048 297 in `q41-summary.md`). Decode is about 32 t/s versus
+  55.7 t/s for Q4_1 on the reference host, which is not a fair comparison because the hosts differ.
+- Reproduce: `bench_run.sh` runs three copies of llama-bench (stock, patch before, patch after) from the directories
+  set in `LD_LIBRARY_PATH`/paths in the script. Those copies live under `/tmp/opencode` and are not kept. Rebuild them from
+  `de5122d87` (stock) and from `build.sh` (patched). `test_nvfp4.c` and `tiled_test.c` are the kernel checks above.

benchmarks/nvfp4/bench_run.sh

@@ -0,0 +1,11 @@
+#!/bin/bash
+# interleaved A/B: each round runs stock, patched-before, patched-after for NVFP4 only
+M=/root/flint_cpu_opt/flint-nvfp4.gguf
+OUT=/tmp/opencode/bench
+for R in 1 2 3; do
+  for B in bin_stock bin_patched_before bin_ppl_new; do
+    LD_LIBRARY_PATH=/tmp/opencode/$B taskset -c 0-15 /tmp/opencode/$B/llama-bench -m $M -t 16 -fa 1 -b 1024 -ub 1024 -lm mmap+mlock -p 256,2048 -n 128 -r 3 -o csv > $OUT/${B}_r$R.csv 2> $OUT/${B}_r$R.err
+    echo "$B round$R exit $? $(uptime | sed 's/.*average://')" >> $OUT/done.log
+  done
+done
+echo ALLDONE >> $OUT/done.log

benchmarks/nvfp4/bin_patched_before_r1.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T20:02:06Z","3312164723","272821630","77.654558","6.627089"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T20:02:21Z","24422205324","3043096286","83.861574","0.659045"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:04:01Z","3997120387","118519237","32.041668","0.941769"

benchmarks/nvfp4/bin_patched_before_r2.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T20:07:24Z","3055729123","36750733","83.785192","1.014280"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T20:07:37Z","24271223135","273666724","84.386872","0.946025"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:09:14Z","4029829668","17507005","31.763528","0.137680"

benchmarks/nvfp4/bin_patched_before_r3.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T20:12:41Z","3115243871","174625727","82.344664","4.499385"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T20:12:54Z","24047580535","123875970","85.166000","0.438822"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:14:30Z","4156609034","39183991","30.796166","0.291850"

benchmarks/nvfp4/bin_ppl_new_r1.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T20:04:17Z","1701202055","265548301","153.080592","25.051629"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T20:04:24Z","9925366157","487402521","206.674830","10.231588"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:05:03Z","3984507398","114876303","32.142090","0.919293"

benchmarks/nvfp4/bin_ppl_new_r2.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T20:09:31Z","1440909288","122602118","178.486780","14.476883"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T20:09:37Z","9376213979","351329487","218.631981","8.285264"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:10:14Z","4316434407","422247907","29.834551","2.767181"

benchmarks/nvfp4/bin_ppl_new_r3.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T20:14:47Z","1409540222","17310901","181.637662","2.217434"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T20:14:54Z","9740323935","742766637","211.051780","15.582107"
+"de5122d87","1","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:15:32Z","3955095681","36151817","32.365110","0.294863"

benchmarks/nvfp4/bin_stock_r1.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T19:59:33Z","3051342444","30706918","83.903179","0.846589"
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T19:59:46Z","26896699672","1704117578","76.355246","5.020894"
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:01:31Z","10405205386","1600748572","14.122672","6.412075"

benchmarks/nvfp4/bin_stock_r2.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T20:05:19Z","3057753126","7591527","83.721949","0.207641"
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T20:05:31Z","24160641638","134679578","84.767711","0.471008"
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:07:09Z","3961015788","80138381","32.323836","0.659357"

benchmarks/nvfp4/bin_stock_r3.csv

@@ -0,0 +1,4 @@
+build_commit,build_number,cpu_info,gpu_info,backends,model_filename,model_type,model_size,model_n_params,n_batch,n_ubatch,n_threads,cpu_mask,cpu_strict,poll,type_k,type_v,n_gpu_layers,n_cpu_moe,split_mode,main_gpu,no_kv_offload,flash_attn,devices,tensor_split,tensor_buft_overrides,load_mode,lazy_mode,embeddings,no_op_offload,no_host,repack,fit_target,fit_min_ctx,n_prompt,n_gen,n_depth,test_time,avg_ns,stddev_ns,avg_ts,stddev_ts
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","256","0","0","2026-11-02T20:10:31Z","3526759564","83460948","72.614765","1.704199"
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","2048","0","0","2026-11-02T20:10:45Z","24491094581","185574398","83.625447","0.636362"
+"unknown","0","Intel(R) Xeon(R) Platinum 8559C","","CPU","/root/flint_cpu_opt/flint-nvfp4.gguf","flint ?B NVFP4","22857561568","39785323896","1024","1024","16","0x0","0","50","f16","f16","-1","0","layer","0","0","1","auto","0.00","none","mmap+mlock","auto","0","0","0","1","0","0","0","128","0","2026-11-02T20:12:24Z","4183101976","62041693","30.603761","0.450938"

benchmarks/nvfp4/ppl_nvfp4_opt.log

@@ -0,0 +1,10 @@
+0.01.321.361 I cmn          init: llama threadpool init, n_threads = 16
+0.01.446.694 I 
+0.01.446.805 I system_info: n_threads = 16 (n_threads_batch = 16) / 192 | CPU : SSE3 = 1 | SSSE3 = 1 | AVX = 1 | AVX_VNNI = 1 | AVX2 = 1 | F16C = 1 | FMA = 1 | BMI2 = 1 | AVX512 = 1 | AVX512_VBMI = 1 | AVX512_VNNI = 1 | AVX512_BF16 = 1 | AMX_INT8 = 1 | LLAMAFILE = 1 | OPENMP = 1 | REPACK = 1 | 
+0.01.446.823 I perplexity: tokenizing the input ..
+0.02.276.737 I perplexity: tokenization took 829.903 ms
+0.02.276.945 I perplexity: calculating perplexity over 16 chunks, n_ctx=512, batch_size=512, n_seq=1
+0.05.073.679 I perplexity: 2.80 seconds per pass - ETA 0.73 minutes
+[1]18.2490,[2]22.6034,[3]23.5732,[4]24.1180,[5]22.7978,[6]24.8210,[7]25.0443,[8]26.6022,[9]26.3313,[10]26.4278,[11]26.0816,[12]26.7523,[13]26.3263,[14]25.5603,[15]26.2636,[16]24.4135,
+0.46.897.942 I Final estimate: PPL = 24.4135 +/- 1.56769
+

benchmarks/nvfp4/ppl_stock.log

@@ -0,0 +1,10 @@
+0.05.265.549 I cmn          init: llama threadpool init, n_threads = 16
+0.05.391.697 I 
+0.05.391.798 I system_info: n_threads = 16 (n_threads_batch = 16) / 192 | CPU : SSE3 = 1 | SSSE3 = 1 | AVX = 1 | AVX_VNNI = 1 | AVX2 = 1 | F16C = 1 | FMA = 1 | BMI2 = 1 | AVX512 = 1 | AVX512_VBMI = 1 | AVX512_VNNI = 1 | AVX512_BF16 = 1 | AMX_INT8 = 1 | LLAMAFILE = 1 | OPENMP = 1 | REPACK = 1 | 
+0.05.391.816 I perplexity: tokenizing the input ..
+0.06.315.987 I perplexity: tokenization took 924.156 ms
+0.06.316.239 I perplexity: calculating perplexity over 16 chunks, n_ctx=512, batch_size=512, n_seq=1
+0.13.030.958 I perplexity: 6.71 seconds per pass - ETA 1.78 minutes
+[1]18.5652,[2]22.5217,[3]23.1922,[4]24.0153,[5]22.5697,[6]24.5359,[7]24.7302,[8]26.3839,[9]26.1684,[10]26.2347,[11]25.8563,[12]26.5899,[13]26.1893,[14]25.4773,[15]26.1958,[16]24.4165,
+1.49.754.158 I Final estimate: PPL = 24.4165 +/- 1.56606
+

benchmarks/nvfp4/summary.csv

@@ -0,0 +1,28 @@
+build,round,test,t_s
+stock,1,pp256,83.903
+stock,1,pp2048,76.355
+stock,1,tg128,14.123
+stock,2,pp256,83.722
+stock,2,pp2048,84.768
+stock,2,tg128,32.324
+stock,3,pp256,72.615
+stock,3,pp2048,83.625
+stock,3,tg128,30.604
+patch_before,1,pp256,77.655
+patch_before,1,pp2048,83.862
+patch_before,1,tg128,32.042
+patch_before,2,pp256,83.785
+patch_before,2,pp2048,84.387
+patch_before,2,tg128,31.764
+patch_before,3,pp256,82.345
+patch_before,3,pp2048,85.166
+patch_before,3,tg128,30.796
+patch_nvfp4_opt,1,pp256,153.081
+patch_nvfp4_opt,1,pp2048,206.675
+patch_nvfp4_opt,1,tg128,32.142
+patch_nvfp4_opt,2,pp256,178.487
+patch_nvfp4_opt,2,pp2048,218.632
+patch_nvfp4_opt,2,tg128,29.835
+patch_nvfp4_opt,3,pp256,181.638
+patch_nvfp4_opt,3,pp2048,211.052
+patch_nvfp4_opt,3,tg128,32.365

benchmarks/nvfp4/test_nvfp4.c

@@ -0,0 +1,48 @@
+// Compare ggml_vec_dot_nvfp4_q8_0 (arch path) against ggml_vec_dot_nvfp4_q8_0_generic on random data.
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <math.h>
+#include <stddef.h>
+
+typedef struct { uint8_t d[4]; uint8_t qs[32]; } block_nvfp4;   // 64 elems: 4 UE4M3 scales + 32 nibble bytes
+typedef struct { _Float16 d; int8_t qs[32]; } block_q8_0;       // 32 elems
+
+void ggml_vec_dot_nvfp4_q8_0(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc);
+void ggml_vec_dot_nvfp4_q8_0_generic(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc);
+
+static uint32_t rng_state = 12345;
+static uint32_t rnd(void) { rng_state = rng_state * 1664525u + 1013904223u; return rng_state >> 8; }
+
+void ggml_cpu_init(void);
+int main(void) {
+    ggml_cpu_init();
+    int fails = 0;
+    for (int trial = 0; trial < 20; ++trial) {
+        const int n = 64 * (1 + (trial * 37) % 200);
+        const int nb = n / 64;
+        block_nvfp4 * w = aligned_alloc(64, sizeof(block_nvfp4) * nb);
+        block_q8_0  * a = aligned_alloc(64, sizeof(block_q8_0) * (n / 32));
+        for (int i = 0; i < nb; ++i) {
+            for (int s = 0; s < 4; ++s) w[i].d[s] = (uint8_t)(rnd() % 127);  // avoid NaN code 0x7F
+            for (int j = 0; j < 32; ++j) w[i].qs[j] = (uint8_t)rnd();
+        }
+        for (int i = 0; i < n / 32; ++i) {
+            float amax = 0.f; float xs[32];
+            for (int j = 0; j < 32; ++j) { xs[j] = ((int)(rnd() % 20001) - 10000) / 1000.0f; if (fabsf(xs[j]) > amax) amax = fabsf(xs[j]); }
+            float d = amax / 127.0f; float id = d ? 1.0f / d : 0.0f;
+            a[i].d = (_Float16) d;
+            for (int j = 0; j < 32; ++j) a[i].qs[j] = (int8_t) lrintf(xs[j] * id);
+        }
+        float r_gen = 0.f, r_arch = 0.f;
+        ggml_vec_dot_nvfp4_q8_0_generic(n, &r_gen, 0, w, 0, a, 0, 1);
+        ggml_vec_dot_nvfp4_q8_0(n, &r_arch, 0, w, 0, a, 0, 1);
+        double rel = fabs((double) r_arch - r_gen) / (fabs((double) r_gen) + 1e-9);
+        printf("n=%6d generic=% .6f arch=% .6f rel=%.3e %s\n", n, r_gen, r_arch, rel, rel < 1e-5 ? "OK" : "FAIL");
+        if (rel >= 1e-5) fails++;
+        free(w); free(a);
+    }
+    printf("%s (%d failures)\n", fails ? "FAILED" : "PASSED", fails);
+    return fails != 0;
+}

benchmarks/nvfp4/tiled_test.c

@@ -0,0 +1,80 @@
+// Dense mul_mat check for NVFP4: tiled path (M=64 columns) vs vec_dot path (M=1) vs float reference.
+#include "ggml.h"
+#include "ggml-cpu.h"
+#include <math.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+void ggml_cpu_init(void);
+void dequantize_row_nvfp4(const void * x, float * y, int64_t k);
+
+static uint32_t st = 777;
+static uint32_t rnd(void) { st = st * 1664525u + 1013904223u; return st >> 8; }
+static float frnd(void) { return ((int) (rnd() % 20001) - 10000) / 10000.0f; }
+
+int main(int argc, char ** argv) {
+    const int K = argc > 1 ? atoi(argv[1]) : 2048;
+    const int N = argc > 2 ? atoi(argv[2]) : 512;
+    const int M = 64;
+
+    ggml_cpu_init();
+    struct ggml_init_params ip = { .mem_size = 512u * 1024 * 1024, .mem_buffer = NULL, .no_alloc = false };
+    struct ggml_context * ctx = ggml_init(ip);
+
+    struct ggml_tensor * w  = ggml_new_tensor_2d(ctx, GGML_TYPE_NVFP4, K, N);
+    struct ggml_tensor * xM = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, K, M);
+    struct ggml_tensor * x1 = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, K, 1);
+
+    // random valid nvfp4 blocks (UE4M3 codes below the NaN code)
+    uint8_t * wd = (uint8_t *) w->data;
+    const size_t blocks = (size_t) N * (K / 64);
+    for (size_t b = 0; b < blocks; b++) {
+        uint8_t * blk = wd + b * 36;
+        for (int s = 0; s < 4; s++) blk[s] = (uint8_t) (rnd() % 0x70);
+        for (int j = 0; j < 32; j++) blk[4 + j] = (uint8_t) rnd();
+    }
+    float * xd = (float *) xM->data;
+    for (int m = 0; m < M; m++) for (int k = 0; k < K; k++) xd[(size_t) m * K + k] = frnd();
+    for (int k = 0; k < K; k++) ((float *) x1->data)[k] = xd[k];  // column 0
+
+    struct ggml_tensor * yM = ggml_mul_mat(ctx, w, xM);
+    struct ggml_tensor * y1 = ggml_mul_mat(ctx, w, x1);
+    struct ggml_cgraph * gM = ggml_new_graph(ctx);
+    ggml_build_forward_expand(gM, yM);
+    ggml_graph_compute_with_ctx(ctx, gM, 16);
+    struct ggml_cgraph * g1 = ggml_new_graph(ctx);
+    ggml_build_forward_expand(g1, y1);
+    ggml_graph_compute_with_ctx(ctx, g1, 16);
+
+    // float reference from the dequantized weights
+    float * wf = malloc(sizeof(float) * (size_t) K * N);
+    for (int n = 0; n < N; n++) dequantize_row_nvfp4(wd + (size_t) n * (K / 64) * 36, wf + (size_t) n * K, K);
+
+    double e_tiled = 0, n_tiled = 0, e_ref = 0, n_ref = 0, e_ref1 = 0;
+    const float * yMd = (const float *) yM->data;
+    const float * y1d = (const float *) y1->data;
+    for (int n = 0; n < N; n++) {
+        double ref = 0;
+        for (int k = 0; k < K; k++) ref += (double) wf[(size_t) n * K + k] * xd[k];
+        // column 0 of the tiled result vs the vec_dot result
+        const double a = yMd[n], b = y1d[n];
+        e_tiled += (a - b) * (a - b); n_tiled += b * b;
+        e_ref += (b - ref) * (b - ref); n_ref += ref * ref;
+        for (int m = 1; m < M; m++) {
+            double r = 0;
+            for (int k = 0; k < K; k++) r += (double) wf[(size_t) n * K + k] * xd[(size_t) m * K + k];
+            const double t = yMd[(size_t) m * N + n];
+            e_ref1 += (t - r) * (t - r);
+            n_ref  += r * r;
+        }
+    }
+    printf("K=%d N=%d M=%d\n", K, N, M);
+    printf("tiled(M=64) vs vec_dot(M=1), col0: rel L2 = %.3e\n", sqrt(e_tiled / n_tiled));
+    printf("vec_dot(M=1) vs f32 ref (col0):    rel L2 = %.3e\n", sqrt(e_ref / n_ref));
+    printf("tiled(M=64) vs f32 ref (cols 1..): rel L2 = %.3e\n", sqrt(e_ref1 / n_ref));
+    free(wf);
+    ggml_free(ctx);
+    return 0;
+}

handoff/HANDOFF.md

@@ -0,0 +1,92 @@
+# HANDOFF: flint CPU optimization (Q4_1 + NVFP4)
+
+Last updated: 2026-11-02. Repo: `/root/flint_cpu_opt` (mirrored to `HelloSun/flint_cpu_opt` on HF).
+Read `/root/flint_cpu_opt/README.md` (model card) for the public numbers. This file is for continuing the work.
+
+## 1. Current state
+
+| item | state |
+|---|---|
+| llama.cpp source | `/root/flint_cpu_opt/llama.cpp`, base `de5122d879d00d2895c36be7e0cee94105aeeabd`, working tree = base + `patches/flint-cpu-opt.patch` (verified: `git diff` == patch file) |
+| patch | Q4_1 kernels (vec_dot + tiled), IQ4_XS, Q4_K/Q5_K/Q6_K tweaks, **NVFP4 vec_dot + tiled (new)** |
+| models | `flint-int4-q4_1.gguf` (24.2 GiB, served), `flint-nvfp4.gguf` (22.9 GB, sha256 `97cc8009...b8bde`), `flint-UD-Q4_K_XL.gguf` not in local dir |
+| llama-server | running on `127.0.0.1:20000`, alias `flint-q4_1`, `-t 16 -fa on --jinja -c 131072`, PID in `/root/llama-server.pid`, log `/root/llama-server.log` |
+| built binaries | `llama.cpp/build/bin/` (llama-server, llama-bench, llama-perplexity; rebuilt with final code) |
+| benchmarks | NVFP4 in `benchmarks/nvfp4/` (3 interleaved rounds, medians in README). Q4_1 / Q4_K_XL in `benchmarks/*.md` (from HF, not re-run) |
+| HF | `HelloSun/flint_cpu_opt`: model card + patch + handoff uploaded (see section 6) |
+
+## 2. NVFP4 results (median of 3 interleaved rounds, tokens/s)
+
+| build | pp256 | pp2048 | tg128 |
+|---|---|---|---|
+| stock de5122d87 | 83.7 | 83.6 | 30.6 |
+| patch, NVFP4 untouched | 82.3 | 84.4 | 31.8 |
+| patch + NVFP4 kernels | 178.5 | 211.1 | 32.1 |
+
+Correctness: AVX512 vec_dot vs generic rel err <= 1.3e-6; tiled vs vec_dot rel L2 ~2e-7; PPL wikitext-2 16x512: stock 24.4165, new 24.4135.
+
+Q4_1 reference (from earlier, quieter host): pp256 266, pp2048 297, tg128 55.7. Decode of NVFP4 (32 t/s) is far below Q4_1, and NVFP4 prefill is still below Q4_1.
+
+## 3. Code map (NVFP4 work)
+
+- `ggml/src/ggml-cpu/arch/x86/quants.c`: `ggml_vec_dot_nvfp4_q8_0`, first branch `#if AVX512F && AVX512VL && AVX512VNNI && AVX2` (one 64-block per iteration; `alo/ahi` built with `_mm512_permutexvar_epi64`; lane k of the int32 sum = subblock k/2).
+- `ggml/src/ggml-cpu/tiled/tiled-kernel.h`: `tiled_tile_src0::sd16` (per-16 f32 scale), `tiled_unpk_fp4_64` (AVX2 + scalar).
+- `ggml/src/ggml-cpu/tiled/tiled-kernel.cpp`: `tiled_run_micro_nv4_vnni_8x16` / `tiled_run_microtile_nv4_vnni`; routed from `tiled_run_microtile` when `Q41 && SUBBLK==16`; explicit instantiation `<16,false,0,false,true>` at the bottom.
+- `ggml/src/ggml-cpu/tiled/tiled.cpp`: `tiled_unpack_src0(const block_nvfp4*)`, `tiled_unpack_src1(const block_q8_0*)`, NVFP4 in `ggml_tiled_supported` list and in `ggml_tiled_matmul_type_dispatch` (VNNI-only guard).
+- Kernel tests: `benchmarks/nvfp4/test_nvfp4.c` (vec_dot vs generic) and `benchmarks/nvfp4/tiled_test.c` (tiled vs vec_dot vs f32 ref). Build with `gcc ... -L<build>/bin -lggml-cpu -lggml-base -Wl,-rpath,<build>/bin` and call `ggml_cpu_init()` first (the UE4M3/fp16 tables are otherwise zero, which gives a false pass).
+
+Bias trick (important for any change): codes are `LUT+12` (0..24, unsigned), so `sum(w*a) = sum((w+12)*a) - 12*bsum16`. `bsums` is the per-16 sum of the int8 activation codes.
+
+## 4. How to reproduce
+
+```bash
+cd /root/flint_cpu_opt
+# rebuild (after editing llama.cpp sources)
+cmake --build llama.cpp/build -j16 --target llama-bench llama-server llama-perplexity
+# verify the patch file matches the tree
+(cd llama.cpp && git diff) | cmp - patches/flint-cpu-opt.patch
+# nvfp4 bench (copy binaries first; the script expects bin_* copies, see benchmarks/nvfp4/bench_run.sh)
+taskset -c 0-15 llama.cpp/build/bin/llama-bench -m flint-nvfp4.gguf -t 16 -fa 1 -b 1024 -ub 1024 -lm mmap+mlock -p 256,2048 -n 128 -r 3 -o csv
+# PPL check (needs wikitext-2 test; HF: datasets/Salesforce/wikitext wikitext-2-raw-v1)
+taskset -c 0-15 llama.cpp/build/bin/llama-perplexity -m flint-nvfp4.gguf -f wt2_test.txt -c 512 --chunks 16 -t 16 -fa 1 -b 512 -ub 512
+```
+
+Stock baseline: `git archive de5122d879d00d2895c36be7e0cee94105aeeabd | tar -x -C <dir>` and build with `cmake -B build -DGGML_NATIVE=ON -DGGML_CUDA=OFF -DLLAMA_CURL=OFF`.
+
+Restart server (same args as running server):
+```bash
+setsid nohup llama.cpp/build/bin/llama-server -m /root/flint_cpu_opt/flint-int4-q4_1.gguf --alias flint-q4_1 -t 16 -fa on --jinja -c 131072 --host 127.0.0.1 --port 20000 > /root/llama-server.log 2>&1 < /dev/null &
+```
+Note: `build.sh` defaults to `CONTEXT_SIZE=32768`; the running server uses 131072.
+
+## 5. Protocol and gotchas
+
+- The host is shared. Load average was 55-63 during the NVFP4 runs; decode single runs swung from 14 to 32 t/s. Use interleaved rounds (A,B,C,A,B,C,...) and medians. Differences under ~8 % are noise.
+- Stop `llama-server` (port 20000) before benchmarks (it uses ~25 GB RSS and CPUs); restart afterwards. The server serves the user's pi agent.
+- Memory guard: processes you start must stay below 80 GiB total; use `mmap+mlock` on the 22 GB file, do not load it twice in parallel.
+- Downloads: HF single-stream curl was about 7-19 MB/s. 8 parallel `-r` range requests were much faster (about 15 GB in 2 min). Use `setsid nohup` for background jobs; a plain `&` in the tool shell hung the shell.
+- Do not `pkill -f <pattern>` with a pattern that appears in your own command line (it killed the shell once).
+- The GGML_PROF profiler patch (`optional-profiler.patch`) changes `ggml-cpu.c`; apply it only temporarily, build to a copy directory, then `git apply -R`.
+- The HF token was pasted in chat. Rotate it when convenient. Never write it into files; use `HF_TOKEN` env.
+- `git diff` of the llama.cpp tree is the source of truth for the patch. Regenerate with `git diff > patches/flint-cpu-opt.patch` after every kernel change, and check it applies to `BASE_COMMIT` (`git archive` into a temp dir + `git apply --check`).
+
+## 6. HF upload
+
+Uploaded to `HelloSun/flint_cpu_opt`: `README.md` (model card, integrated with NVFP4), `patches/flint-cpu-opt.patch` (updated), `benchmarks/nvfp4/*`, `handoff/*`. The 22.9 GB NVFP4 GGUF is NOT in this repo (link to RecursiveSI). Old Q4_K_XL / Q4_1 files and benchmarks are unchanged.
+
+## 7. Open items / next steps (priority order)
+
+1. **NVFP4 decode is not improved** (32 t/s, ~1.0x). The vec_dot per 64 block has about 25 ops and the tiled path is not used at batch 1. Next: microbenchmark `ggml_vec_dot_nvfp4_q8_0` on a cold MoE shape (`mul_mat_id` K=2048, 256 experts, 6 active) and check memory-bandwidth share (the Q4_1 patch got +15 % this way with prefetch, `GGML_Q41_PF`). Try the same prefetch knob for NVFP4.
+2. **NVFP4 prefill is still ~30 % below Q4_1** (pp2048 211 vs 297, different host state). Profile `MUL_MAT_ID` with `GGML_PROF=1` (apply the profiler patch temporarily). Check unpack cost (`tiled_unpk_fp4_64` is per 64 elements) and the per-16 accumulation (2 int-sums per 32 block).
+3. Accuracy: run KLD for NVFP4 patched vs stock (save baseline logits with llama-perplexity, then compare; check `llama-perplexity --help` for the KL flags, and follow the method in `q41-kld.log`). Only PPL was measured.
+4. Re-run Q4_1 / Q4_K_XL only if asked. They were not re-run in this round.
+5. Non-VNNI x86 (AVX2 only): NVFP4 tiled is intentionally disabled there; the AVX2 vec_dot path is unchanged. Not tested on AVX2-only hardware.
+6. Optional: update `build.sh` default `CONTEXT_SIZE` note or leave as is.
+
+## 8. Where things are
+
+- Work tree: `/root/flint_cpu_opt/llama.cpp`
+- Model card: `/root/flint_cpu_opt/README.md`; handoff: `/root/flint_cpu_opt/handoff/HANDOFF.md`
+- NVFP4 benchmarks + tests: `/root/flint_cpu_opt/benchmarks/nvfp4/`
+- Temporary binaries and downloads: `/tmp/opencode/` (bin_stock, bin_patched_before, bin_ppl_new, nvfp4_test, bench). These are not persistent; rebuild if needed.
+- Earlier handoff (Q4_1 only): on HF `HANDOFF.md`, kept for history.

patches/flint-cpu-opt.patch

@@ -172,7 +172,7 @@ index 9f3a744b5..d23e4d971 100644
  
          parallel_for_ggml(params, n_batch * NB, [&](int begin, int end) {
 diff --git a/ggml/src/ggml-cpu/arch/x86/quants.c b/ggml/src/ggml-cpu/arch/x86/quants.c
-index 21cd6312a..07d4feb37 100644
+index 21cd6312a..37efb7fae 100644
 --- a/ggml/src/ggml-cpu/arch/x86/quants.c
 +++ b/ggml/src/ggml-cpu/arch/x86/quants.c
 @@ -941,6 +941,49 @@ void ggml_vec_dot_q4_1_q8_1(int n, float * GGML_RESTRICT s, size_t bs, const voi
@@ -234,7 +234,52 @@ index 21cd6312a..07d4feb37 100644
  #else
      UNUSED(nb);
      UNUSED(x);
-@@ -2118,7 +2161,89 @@ void ggml_vec_dot_q4_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const voi
+@@ -1080,7 +1123,43 @@ void ggml_vec_dot_nvfp4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const vo
+     int ib = 0;
+     float sumf = 0;
+ 
+-#if defined(__AVX2__)
++#if defined(__AVX512F__) && defined(__AVX512VL__) && defined(__AVX512VNNI__) && defined(__AVX2__)
++    // AVX512-VNNI path: one 64-element block per iteration. |w| (0..12) is the unsigned vpdpbusd operand and the
++    // sign of w moves onto the int8 activations. Int32 lane k sums sub-block k/2 (lo + hi nibbles of 4 bytes).
++    {
++        const __m256i lut  = _mm256_broadcastsi128_si256(_mm_loadu_si128((const __m128i *) kvalues_fp4));
++        const __m256i m4b  = _mm256_set1_epi8(0x0f);
++        const __m512i perm_lo = _mm512_setr_epi64(0, 2, 4, 6, 0, 0, 0, 0);
++        const __m512i perm_hi = _mm512_setr_epi64(1, 3, 5, 7, 0, 0, 0, 0);
++        __m256 acc = _mm256_setzero_ps();
++        for (; ib < nb; ++ib) {
++            const __m256i qbits = _mm256_loadu_si256((const __m256i *) x[ib].qs);
++            const __m256i wlo = _mm256_shuffle_epi8(lut, _mm256_and_si256(qbits, m4b));
++            const __m256i whi = _mm256_shuffle_epi8(lut, _mm256_and_si256(_mm256_srli_epi16(qbits, 4), m4b));
++
++            // activation bytes in weight byte order: lo -> elems 16s+j, hi -> elems 16s+8+j
++            const __m512i z = _mm512_inserti64x4(_mm512_castsi256_si512(_mm256_loadu_si256((const __m256i *) y[2*ib + 0].qs)),
++                                                 _mm256_loadu_si256((const __m256i *) y[2*ib + 1].qs), 1);
++            const __m256i alo = _mm512_castsi512_si256(_mm512_permutexvar_epi64(perm_lo, z));
++            const __m256i ahi = _mm512_castsi512_si256(_mm512_permutexvar_epi64(perm_hi, z));
++
++            const __m256i dot = _mm256_dpbusd_epi32(_mm256_dpbusd_epi32(_mm256_setzero_si256(),
++                                    _mm256_abs_epi8(wlo), _mm256_sign_epi8(alo, wlo)),
++                                    _mm256_abs_epi8(whi), _mm256_sign_epi8(ahi, whi));
++
++            const float dy0 = GGML_CPU_FP16_TO_FP32(y[2*ib + 0].d);
++            const float dy1 = GGML_CPU_FP16_TO_FP32(y[2*ib + 1].d);
++            const float s0 = GGML_CPU_UE4M3_TO_FP32(x[ib].d[0]) * dy0;
++            const float s1 = GGML_CPU_UE4M3_TO_FP32(x[ib].d[1]) * dy0;
++            const float s2 = GGML_CPU_UE4M3_TO_FP32(x[ib].d[2]) * dy1;
++            const float s3 = GGML_CPU_UE4M3_TO_FP32(x[ib].d[3]) * dy1;
++            const __m256 scales = _mm256_set_ps(s3, s3, s2, s2, s1, s1, s0, s0);
++
++            acc = _mm256_fmadd_ps(_mm256_cvtepi32_ps(dot), scales, acc);
++        }
++        sumf = hsum_float_8(acc);
++    }
++#elif defined(__AVX2__)
+ 
+     const __m128i values128 = _mm_loadu_si128((const __m128i*)kvalues_fp4);
+     const __m128i m4b  = _mm_set1_epi8(0x0f);
+@@ -2118,7 +2197,89 @@ void ggml_vec_dot_q4_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const voi
  
      uint32_t utmp[4];
  
@@ -325,7 +370,7 @@ index 21cd6312a..07d4feb37 100644
  
      const __m256i m4 = _mm256_set1_epi8(0xF);
  
-@@ -2296,7 +2421,104 @@ void ggml_vec_dot_q5_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const voi
+@@ -2296,7 +2457,104 @@ void ggml_vec_dot_q5_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const voi
  
      uint32_t utmp[4];
  
@@ -431,7 +476,7 @@ index 21cd6312a..07d4feb37 100644
  
      const __m256i m4 = _mm256_set1_epi8(0xF);
      const __m128i mzero = _mm_setzero_si128();
-@@ -2500,7 +2722,81 @@ void ggml_vec_dot_q6_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const voi
+@@ -2500,7 +2758,81 @@ void ggml_vec_dot_q6_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const voi
  
      const int nb = n / QK_K;
  
@@ -514,7 +559,7 @@ index 21cd6312a..07d4feb37 100644
  
      const __m256i m3 = _mm256_set1_epi8(3);
      const __m256i m15 = _mm256_set1_epi8(15);
-@@ -4078,7 +4374,75 @@ void ggml_vec_dot_iq4_xs_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const v
+@@ -4078,7 +4410,75 @@ void ggml_vec_dot_iq4_xs_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const v
  
      const int nb = n / QK_K;
  
@@ -758,10 +803,10 @@ index 9282ca299..0431b7ada 100644
                  return &q4_K_8x8_q8_K;
              }
 diff --git a/ggml/src/ggml-cpu/tiled/tiled-kernel.cpp b/ggml/src/ggml-cpu/tiled/tiled-kernel.cpp
-index 5819ac07c..4d4eaae6f 100644
+index 5819ac07c..b28255b79 100644
 --- a/ggml/src/ggml-cpu/tiled/tiled-kernel.cpp
 +++ b/ggml/src/ggml-cpu/tiled/tiled-kernel.cpp
-@@ -204,6 +204,80 @@ static void tiled_run_microtile_vnni(const tiled_tile_src0 & src0, const tiled_t
+@@ -204,6 +204,147 @@ static void tiled_run_microtile_vnni(const tiled_tile_src0 & src0, const tiled_t
      tiled_run_micro_vnni_8x16<SUBBLK, HAS_MIN, BIAS, NK>(src0, src1, i0 + 8,  j0, num_k, slab, buf, buf_stride);
  }
  
@@ -838,11 +883,78 @@ index 5819ac07c..4d4eaae6f 100644
 +    tiled_run_micro_q41_vnni_8x16<NK>(src0, src1, i0,      j0, num_k, slab, buf, buf_stride);
 +    tiled_run_micro_q41_vnni_8x16<NK>(src0, src1, i0 + 8,  j0, num_k, slab, buf, buf_stride);
 +}
++
++// nvfp4: 16 subblocks of 16 per 256-slab, each with its own f32 weight scale (sd16); q8_0 activations carry one
++// f32 scale per 32-block (sd). Codes are LUT + 12 (unsigned), so the bias is removed per subblock with the
++// int32 code sums (bsums): sum(w*a) = sum((w+12)*a) - 12*sum(a). Integer sums are exact.
++template <int NK>
++static void tiled_run_micro_nv4_vnni_8x16(const tiled_tile_src0 & src0, const tiled_tile_src1 & src1,
++                                          int i0, int j0, int num_k, int slab, float * buf, int buf_stride) {
++    constexpr int NB32 = TILED_TILE_K / 32;  // 8 activation blocks
++    constexpr int NB16 = TILED_TILE_K / 16;  // 16 weight subblocks
++    constexpr int NG   = 4;                  // 4-element groups per 16-subblock
++    constexpr int NUM_ROWS = 8;
++
++    const int nkr  = (NK > 0) ? NK : num_k;
++    const int seff = (NK > 0) ? 0 : slab;
++    const int qk_stride = nkr * TILED_TILE_K;
++    const int qk_off = seff * TILED_TILE_K;
++    const int nb_stride = NB16 * nkr;
++    const int nb_off = seff * NB16;
++    const int bs_stride = (nkr == 1) ? TILED_TILE_ROWS : TILED_MICRO;
++
++    __m512 facc[NUM_ROWS];
++    for (int t = 0; t < NUM_ROWS; t++) { facc[t] = _mm512_setzero_ps(); }
++
++    const int32_t * bsums = src1.bsums;
++    const __m512i twelve = _mm512_set1_epi32(12);
++    for (int sb = 0; sb < NB16; sb++) {
++        const __m512i bsum12 = _mm512_mullo_epi32(twelve,
++            _mm512_load_si512((const __m512i *) &bsums[(seff * TILED_MICRO + sb) * bs_stride + j0]));
++
++        __m512i acc16[NUM_ROWS];
++        for (int t = 0; t < NUM_ROWS; t++) { acc16[t] = _mm512_setzero_si512(); }
++
++#ifdef __GNUC__
++        #pragma GCC unroll 4
++#endif
++        for (int g = 0; g < NG; g++) {
++            const int kg = sb * NG + g;
++            const __m512i codes = _mm512_load_si512((const __m512i *) &src1.q[(j0 / TILED_MICRO) * (TILED_MICRO * qk_stride)
++                + qk_off + (kg % TILED_MICRO) * qk_stride + (kg / TILED_MICRO) * (TILED_MICRO * 4)]);
++            for (int t = 0; t < NUM_ROWS; t++) {
++                const uint32_t u4 = *(const uint32_t *) &src0.q[(i0 + t) * qk_stride + qk_off + kg * 4];
++                const __m512i u4b = _mm512_set1_epi32((int) u4);
++                acc16[t] = _mm512_dpbusd_epi32(acc16[t], u4b, codes);
++            }
++        }
++
++        const __m512 dav = _mm512_loadu_ps(&src1.sd[(seff * NB32 + sb / 2) * bs_stride + j0]);
++        for (int t = 0; t < NUM_ROWS; t++) {
++            const int ar = i0 + t;
++            const float dw = src0.sd16[ar * nb_stride + nb_off + sb];
++            const __m512 raw = _mm512_cvtepi32_ps(_mm512_sub_epi32(acc16[t], bsum12));
++            facc[t] = _mm512_fmadd_ps(raw, _mm512_mul_ps(_mm512_set1_ps(dw), dav), facc[t]);
++        }
++    }
++
++    for (int t = 0; t < NUM_ROWS; t++) {
++        float * p = &buf[(i0 + t) * buf_stride + j0];
++        _mm512_store_ps(p, _mm512_add_ps(_mm512_load_ps(p), facc[t]));
++    }
++}
++
++template <int NK>
++static void tiled_run_microtile_nv4_vnni(const tiled_tile_src0 & src0, const tiled_tile_src1 & src1,
++                                         int i0, int j0, int num_k, int slab, float * buf, int buf_stride) {
++    tiled_run_micro_nv4_vnni_8x16<NK>(src0, src1, i0,      j0, num_k, slab, buf, buf_stride);
++    tiled_run_micro_nv4_vnni_8x16<NK>(src0, src1, i0 + 8,  j0, num_k, slab, buf, buf_stride);
++}
 +
  #endif // __AVX512VNNI__ && __AVX512VL__
  
  #if defined(__AVX2__)
-@@ -496,23 +570,79 @@ static void tiled_run_microtile_avx(const tiled_tile_src0 & src0, const tiled_ti
+@@ -496,23 +637,88 @@ static void tiled_run_microtile_avx(const tiled_tile_src0 & src0, const tiled_ti
  
  #endif // __AVX__ && !__AVX2__
  
@@ -898,6 +1010,15 @@ index 5819ac07c..4d4eaae6f 100644
      const bool standard = (num_k == 1);
  #if defined(__AVX512VNNI__) && defined(__AVX512VL__) && defined(__AVX512DQ__)
 +    if constexpr (Q41) {
++        if constexpr (SUBBLK == 16) {
++            // nvfp4 (per-16 weight scales), see tiled_run_micro_nv4_vnni_8x16
++            if (standard) {
++                tiled_run_microtile_nv4_vnni<1>(src0, src1, i0, j0, num_k, slab, buf, buf_stride);
++            } else {
++                tiled_run_microtile_nv4_vnni<0>(src0, src1, i0, j0, num_k, slab, buf, buf_stride);
++            }
++            return;
++        }
 +        if (standard) {
 +            tiled_run_microtile_q41_vnni<1>(src0, src1, i0, j0, num_k, slab, buf, buf_stride);
 +        } else {
@@ -924,7 +1045,7 @@ index 5819ac07c..4d4eaae6f 100644
      if (standard) {
          tiled_run_microtile_avx2<SUBBLK, HAS_MIN, BIAS, 1, ACTBIAS>(src0, src1, i0, j0, num_k, slab, buf, buf_stride);
      } else {
-@@ -531,6 +661,7 @@ void tiled_run_microtile(const tiled_tile_src0 & src0, const tiled_tile_src1 & s
+@@ -531,6 +737,7 @@ void tiled_run_microtile(const tiled_tile_src0 & src0, const tiled_tile_src1 & s
          tiled_run_microtile_scalar<SUBBLK, HAS_MIN, BIAS, 0>(src0, src1, i0, j0, num_k, slab, buf, buf_stride);
      }
  #endif
@@ -932,21 +1053,24 @@ index 5819ac07c..4d4eaae6f 100644
  }
  
  // explicit instantiations for the in-use formats (q4_K and q5_K share the constants)
-@@ -556,6 +687,9 @@ template void tiled_run_microtile<16, false, 4, true>(const tiled_tile_src0 & sr
+@@ -556,6 +763,12 @@ template void tiled_run_microtile<16, false, 4, true>(const tiled_tile_src0 & sr
  // q2_K: BIAS = 0
  template void tiled_run_microtile<16, true, 0, false>(const tiled_tile_src0 & src0, const tiled_tile_src1 & src1,
                                                 int i0, int j0, int num_k, int slab, float * buf, int buf_stride);
 +// q4_1: per-32-block f32 scales, q8_1 activations
 +template void tiled_run_microtile<32, false, 0, false, true>(const tiled_tile_src0 & src0, const tiled_tile_src1 & src1,
++                                                   int i0, int j0, int num_k, int slab, float * buf, int buf_stride);
++// nvfp4: per-16-subblock f32 scales, q8_0 activations
++template void tiled_run_microtile<16, false, 0, false, true>(const tiled_tile_src0 & src0, const tiled_tile_src1 & src1,
 +                                                   int i0, int j0, int num_k, int slab, float * buf, int buf_stride);
  
  
  #define MIN(a, b) ((a) < (b) ? (a) : (b))
 diff --git a/ggml/src/ggml-cpu/tiled/tiled-kernel.h b/ggml/src/ggml-cpu/tiled/tiled-kernel.h
-index 1a97f406a..4bd9642cc 100644
+index 1a97f406a..c5d081420 100644
 --- a/ggml/src/ggml-cpu/tiled/tiled-kernel.h
 +++ b/ggml/src/ggml-cpu/tiled/tiled-kernel.h
-@@ -29,12 +29,16 @@
+@@ -29,12 +29,18 @@
  // SUBBLK=32 formats index at stride 8 and leave the slack unused.
  struct tiled_tile_src0 {
      static constexpr int NB_MAX = TILED_TILE_K / 16; // max subblocks per 256-elem block
@@ -960,10 +1084,12 @@ index 1a97f406a..4bd9642cc 100644
 +    // q4_1: no super-block scale, each 32-wide block carries its own f32 scale/min
 +    float    sd[TILED_TILE_ROWS * NB32];  // per-32-block scale
 +    float    sm[TILED_TILE_ROWS * NB32];  // per-32-block min
++    // nvfp4: one f32 scale per 16-wide subblock (UE4M3 decoded), codes are LUT + 12 (0..24)
++    float    sd16[TILED_TILE_ROWS * NB_MAX];
  };
  
  // src1 tile: built from q8_K (wdata)
-@@ -45,6 +49,8 @@ struct tiled_tile_src1 {
+@@ -45,6 +51,8 @@ struct tiled_tile_src1 {
      alignas(64) int32_t bsums[(TILED_TILE_K / 16) * TILED_TILE_ROWS];
      // f32 (not f16): q8_k stores fp16, the unpack converts once
      float       d[TILED_TILE_ROWS];
@@ -972,10 +1098,24 @@ index 1a97f406a..4bd9642cc 100644
  };
  
  // per-thread workspace: all tiled state lives here, allocated in wdata (one slot per thread)
-@@ -71,6 +77,12 @@ inline void tiled_unpk_nib4(const uint8_t * src, uint8_t * lo, uint8_t * hi) {
+@@ -71,6 +79,26 @@ inline void tiled_unpk_nib4(const uint8_t * src, uint8_t * lo, uint8_t * hi) {
      _mm256_storeu_si256((__m256i *) lo, _mm256_and_si256(v, _mm256_set1_epi8(0x0F)));
      _mm256_storeu_si256((__m256i *) hi, _mm256_srli_epi32(_mm256_and_si256(v, _mm256_set1_epi8((int8_t) 0xF0)), 4));
  }
++// nvfp4: 32 packed bytes (one 64-element block) -> 64 codes, LUT(kvalues_fp4) + 12 in 0..24.
++// byte B = 8s+j: lo -> elem 16s+j, hi -> elem 16s+8+j
++inline void tiled_unpk_fp4_64(const uint8_t * src, uint8_t * dst) {
++    const __m256i lut = _mm256_setr_epi8(12, 13, 14, 15, 16, 18, 20, 24, 12, 11, 10, 9, 8, 6, 4, 0,
++                                         12, 13, 14, 15, 16, 18, 20, 24, 12, 11, 10, 9, 8, 6, 4, 0);
++    const __m256i v  = _mm256_loadu_si256((const __m256i *) src);
++    const __m256i lo = _mm256_shuffle_epi8(lut, _mm256_and_si256(v, _mm256_set1_epi8(0x0F)));
++    const __m256i hi = _mm256_shuffle_epi8(lut, _mm256_and_si256(_mm256_srli_epi16(v, 4), _mm256_set1_epi8(0x0F)));
++    // x = [s0 | s2], y = [s1 | s3] (128-bit halves)
++    const __m256i x = _mm256_unpacklo_epi64(lo, hi);
++    const __m256i y = _mm256_unpackhi_epi64(lo, hi);
++    _mm256_storeu_si256((__m256i *) dst,        _mm256_permute2x128_si256(x, y, 0x20));
++    _mm256_storeu_si256((__m256i *) (dst + 32), _mm256_permute2x128_si256(x, y, 0x31));
++}
 +// 16-byte source (one q4_1 block: lo -> elems 0..15, hi -> elems 16..31)
 +inline void tiled_unpk_nib4_16(const uint8_t * src, uint8_t * lo, uint8_t * hi) {
 +    const __m128i v = _mm_loadu_si128((const __m128i *) src);
@@ -985,17 +1125,25 @@ index 1a97f406a..4bd9642cc 100644
  // 2-bit values at bit offset S
  template <int S> inline void tiled_unpk_2bit(const uint8_t * src, uint8_t * dst) {
      _mm256_storeu_si256((__m256i *) dst, _mm256_and_si256(
-@@ -129,6 +141,9 @@ inline void tiled_unpk_tern8(const uint8_t * src, int8_t delta, uint8_t * dst) {
+@@ -129,6 +157,17 @@ inline void tiled_unpk_tern8(const uint8_t * src, int8_t delta, uint8_t * dst) {
  inline void tiled_unpk_nib4(const uint8_t * src, uint8_t * lo, uint8_t * hi) {
      for (int l = 0; l < 32; l++) { lo[l] = (uint8_t) (src[l] & 0xF); hi[l] = (uint8_t) (src[l] >> 4); }
  }
++inline void tiled_unpk_fp4_64(const uint8_t * src, uint8_t * dst) {
++    static const uint8_t lut[16] = { 12, 13, 14, 15, 16, 18, 20, 24, 12, 11, 10, 9, 8, 6, 4, 0 };
++    for (int B = 0; B < 32; B++) {
++        const int s = B / 8, j = B % 8;
++        dst[16 * s + j]     = lut[src[B] & 0xF];
++        dst[16 * s + 8 + j] = lut[src[B] >> 4];
++    }
++}
 +inline void tiled_unpk_nib4_16(const uint8_t * src, uint8_t * lo, uint8_t * hi) {
 +    for (int l = 0; l < 16; l++) { lo[l] = (uint8_t) (src[l] & 0xF); hi[l] = (uint8_t) (src[l] >> 4); }
 +}
  template <int S>
  inline void tiled_unpk_2bit(const uint8_t * src, uint8_t * dst) {
      for (int l = 0; l < 32; l++) { dst[l] = (uint8_t) ((src[l] >> S) & 3); }
-@@ -164,7 +179,9 @@ inline void tiled_unpk_tern8(const uint8_t * src, int8_t delta, uint8_t * dst) {
+@@ -164,7 +203,9 @@ inline void tiled_unpk_tern8(const uint8_t * src, int8_t delta, uint8_t * dst) {
  // ACTBIAS (AVX2 only): the activation is pre-biased +128 by tiled_repack_src1,
  // num_k = K-blocks per row: the tile holds num_k slabs at row stride num_k*256
  // (Default case is num_k=1, 256x256 tiles, we go to longer num_k to improve memory bandwidth when num_rows is small)
@@ -1007,10 +1155,10 @@ index 1a97f406a..4bd9642cc 100644
                           int i0, int j0, int num_k, int slab, float * buf, int buf_stride);
  
 diff --git a/ggml/src/ggml-cpu/tiled/tiled.cpp b/ggml/src/ggml-cpu/tiled/tiled.cpp
-index 77e1a9945..225ec0e78 100644
+index 77e1a9945..96640b50d 100644
 --- a/ggml/src/ggml-cpu/tiled/tiled.cpp
 +++ b/ggml/src/ggml-cpu/tiled/tiled.cpp
-@@ -534,11 +534,33 @@ static void tiled_unpack_src0(const block_iq1_m * rows, int64_t row_stride, int
+@@ -534,11 +534,57 @@ static void tiled_unpack_src0(const block_iq1_m * rows, int64_t row_stride, int
      }
  }
  
@@ -1035,6 +1183,30 @@ index 77e1a9945..225ec0e78 100644
 +        }
 +    }
 +}
++
++// nvfp4: 4 x 64-element blocks per 256-slab, each block 4 UE4M3 subblock scales (16 elems each) + 32 nibble bytes.
++// Scales go to sd16 (one per 16-wide subblock); codes are LUT + 12 (see tiled_unpk_fp4_64).
++static void tiled_unpack_src0(const block_nvfp4 * rows, int64_t row_stride, int n_rows, tiled_tile_src0 * tile, int num_k) {
++    GGML_ASSERT(n_rows <= TILED_TILE_ROWS);
++    constexpr int NBLK = TILED_TILE_K / 64; // 4 nvfp4 blocks per 256-slab
++    constexpr int NB16 = TILED_TILE_K / 16; // 16 subblocks per 256-slab
++
++    const int qk_stride = num_k * TILED_TILE_K;
++    const int nb_stride = NB16 * num_k;
++    for (int slab = 0; slab < num_k; slab++) {
++        for (int r = 0; r < n_rows; r++) {
++            const int q_off = r * qk_stride + slab * TILED_TILE_K;
++            const int s_off = r * nb_stride + slab * NB16;
++            const block_nvfp4 * x = rows + r * row_stride + slab * NBLK;
++            for (int b = 0; b < NBLK; b++) {
++                for (int s = 0; s < 4; s++) {
++                    tile->sd16[s_off + 4 * b + s] = ggml_ue4m3_to_fp32(x[b].d[s]);
++                }
++                tiled_unpk_fp4_64(x[b].qs, &tile->q[q_off + b * 64]);
++            }
++        }
++    }
++}
 +
  // unpack src1 tile from q8_K rows. num_k = K-blocks per row: each row decodes num_k consecutive
  // source blocks (rows[r][kblk + slab]) into the tile at row stride num_k*256, so each activation
@@ -1046,7 +1218,7 @@ index 77e1a9945..225ec0e78 100644
      GGML_ASSERT(n_rows <= TILED_TILE_ROWS);
      const int n_padded = (n_rows + TILED_MICRO - 1) & ~(TILED_MICRO - 1);
      const int bs_stride = (num_k == 1) ? TILED_TILE_ROWS : TILED_MICRO;
-@@ -569,6 +591,39 @@ static void tiled_unpack_src1_q8_K(const block_q8_K * const * rows, int n_rows,
+@@ -569,6 +615,71 @@ static void tiled_unpack_src1_q8_K(const block_q8_K * const * rows, int n_rows,
      }
  }
  
@@ -1082,19 +1254,52 @@ index 77e1a9945..225ec0e78 100644
 +        }
 +    }
 +}
++
++// nvfp4 src1: q8_0 activation rows (one f32 scale per 32-wide block, per-16 sums recomputed from the codes).
++static void tiled_unpack_src1(const block_q8_0 * const * rows, int n_rows, tiled_tile_src1 * tile,
++                              int kblk, int num_k) {
++    GGML_ASSERT(n_rows <= TILED_TILE_ROWS);
++    const int n_padded = (n_rows + TILED_MICRO - 1) & ~(TILED_MICRO - 1);
++    const int bs_stride = (num_k == 1) ? TILED_TILE_ROWS : TILED_MICRO;
++    constexpr int NBLK = TILED_TILE_K / 32; // 8 q8_0 blocks per 256-slab
++    constexpr int n16 = TILED_TILE_K / TILED_MICRO;
++    for (int slab = 0; slab < num_k; slab++) {
++        const int q_off = slab * TILED_TILE_K;
++        for (int r = 0; r < n_padded; r++) {
++            if (r < n_rows) {
++                const block_q8_0 * xb = rows[r] + (kblk + slab) * NBLK;
++                for (int b = 0; b < NBLK; b++) {
++                    memcpy(&tile->q[r * (num_k * TILED_TILE_K) + q_off + b * 32], xb[b].qs, 32);
++                    int32_t s0 = 0, s1 = 0;
++                    for (int e = 0; e < 16; e++) { s0 += xb[b].qs[e]; s1 += xb[b].qs[16 + e]; }
++                    tile->bsums[(slab * n16 + 2 * b + 0) * bs_stride + r] = s0;
++                    tile->bsums[(slab * n16 + 2 * b + 1) * bs_stride + r] = s1;
++                    tile->sd[(slab * NBLK + b) * bs_stride + r] = ggml_fp16_to_fp32(xb[b].d);
++                }
++                tile->d[slab * TILED_MICRO + r] = 1.0f; // per-32 scales are applied in the kernel
++            } else {
++                memset(&tile->q[r * (num_k * TILED_TILE_K) + q_off], 0, TILED_TILE_K);
++                for (int i = 0; i < n16; i++) { tile->bsums[(slab * n16 + i) * bs_stride + r] = 0; }
++                for (int b = 0; b < NBLK; b++) { tile->sd[(slab * NBLK + b) * bs_stride + r] = 0.0f; }
++                tile->d[slab * TILED_MICRO + r] = 1.0f;
++            }
++        }
++    }
++}
 +
  // GGML_CPU_TILED_MM: master switch, on by default. If off, we fast return false and normal vec_dot mul_mat resumes
  static bool ggml_tiled_matmul_enabled(void) {
      static bool enabled = true;
-@@ -617,6 +672,7 @@ static bool ggml_tiled_supported(const struct ggml_tensor * src0,
+@@ -617,6 +728,8 @@ static bool ggml_tiled_supported(const struct ggml_tensor * src0,
          case GGML_TYPE_Q6_K:
          case GGML_TYPE_Q5_K:
          case GGML_TYPE_Q4_K:
 +        case GGML_TYPE_Q4_1:
++        case GGML_TYPE_NVFP4:
          case GGML_TYPE_Q3_K:
          case GGML_TYPE_Q2_K:
          case GGML_TYPE_IQ4_XS:
-@@ -716,7 +772,9 @@ static void tiled_store_window(const float * buf, int n_src0, int n_src1, int bu
+@@ -716,7 +829,9 @@ static void tiled_store_window(const float * buf, int n_src0, int n_src1, int bu
  // unpack and microtile are reused unchanged from mul_mat.
  
  // src0 rows per g group; finer than the TILED_TILE_K kernel tile so all threads stay busy at small ne01
@@ -1105,7 +1310,7 @@ index 77e1a9945..225ec0e78 100644
  
  // like tiled_store_window, but the dst columns are not contiguous: col_ptrs[j] points at the
  // start of dst column j, whose rows are contiguous (dim 0, nb0 == 4 bytes).
-@@ -756,11 +814,11 @@ static void tiled_store_window_scatter(const float * buf, int n_src0, int n_src1
+@@ -756,11 +871,11 @@ static void tiled_store_window_scatter(const float * buf, int n_src0, int n_src1
  
  // one (g, k) macrotile: zero the acc window, sweep K in 256 element slabs, scatter the result rows
  // rows points at this k window's routed src1 rows, row r at its base block
@@ -1119,7 +1324,7 @@ index 77e1a9945..225ec0e78 100644
                                     tiled_ws * ws) {
      const int64_t ne00 = src0->ne[0];
      const int64_t ne01 = src0->ne[1];
-@@ -770,6 +828,8 @@ static void tiled_mmid_gemm_window(struct ggml_tensor * dst, const struct ggml_t
+@@ -770,6 +885,8 @@ static void tiled_mmid_gemm_window(struct ggml_tensor * dst, const struct ggml_t
  
      const size_t src0_bs = ggml_type_size(src0->type);
      const int64_t src0_stride = src0->nb[1] / src0_bs;
@@ -1128,7 +1333,7 @@ index 77e1a9945..225ec0e78 100644
  
      // the window is at most TILED_MMID_GROUP x TILED_TILE_K, so zero only that region of acc
      for (int64_t i = 0; i < n_src0; i++) {
-@@ -792,19 +852,19 @@ static void tiled_mmid_gemm_window(struct ggml_tensor * dst, const struct ggml_t
+@@ -792,19 +909,19 @@ static void tiled_mmid_gemm_window(struct ggml_tensor * dst, const struct ggml_t
              const int kstart = (int) (k0 / TILED_TILE_K);
              // activation: one 16-row band (nrows <= 16). Unpack the whole k_extent chunk (all
              // slabs) so each activation row is a long stream, then repack each slab in place
@@ -1151,7 +1356,7 @@ index 77e1a9945..225ec0e78 100644
                  }
              }
          }
-@@ -812,15 +872,15 @@ static void tiled_mmid_gemm_window(struct ggml_tensor * dst, const struct ggml_t
+@@ -812,15 +929,15 @@ static void tiled_mmid_gemm_window(struct ggml_tensor * dst, const struct ggml_t
          // standard per-slab path: K is stepped in 256-element slabs
          for (int64_t ib = 0; ib < ne00; ib += TILED_TILE_K) {
              const int kblk = (int) (ib / TILED_TILE_K);
@@ -1170,7 +1375,7 @@ index 77e1a9945..225ec0e78 100644
                          (int) (ir0 - r), (int) ir1, 1, 0,
                          ws->acc, TILED_TILE_ROWS);
                  }
-@@ -834,7 +894,7 @@ static void tiled_mmid_gemm_window(struct ggml_tensor * dst, const struct ggml_t
+@@ -834,7 +951,7 @@ static void tiled_mmid_gemm_window(struct ggml_tensor * dst, const struct ggml_t
  // one expert of MUL_MAT_ID: each k window's dispatched rows are pointed at by a per-window
  // row pointer list and swept over the thread's row windows, the dst rows are scattered back
  // per window
@@ -1179,7 +1384,7 @@ index 77e1a9945..225ec0e78 100644
  static void ggml_compute_forward_mul_mat_id_tiled_one_expert(
          const struct ggml_compute_params * params,
                struct ggml_tensor *         dst,
-@@ -877,22 +937,22 @@ static void ggml_compute_forward_mul_mat_id_tiled_one_expert(
+@@ -877,22 +994,22 @@ static void ggml_compute_forward_mul_mat_id_tiled_one_expert(
      for (int64_t k = 0; k < cne1; k += TILED_TILE_K) {
          const int64_t nrows = MIN(TILED_TILE_K, cne1 - k);
  
@@ -1206,7 +1411,7 @@ index 77e1a9945..225ec0e78 100644
  static void ggml_compute_forward_mul_mat_tiled_one_chunk(
      const struct ggml_compute_params * params,
      struct ggml_tensor * dst,
-@@ -928,6 +988,8 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
+@@ -928,6 +1045,8 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
  
      const int64_t src0_stride = nb01 / src0_bs;  // blocks between src0 rows
      const int64_t src1_stride = (src1->type == vec_dot_type ? src1->nb[1] : row_size) / src1_bs;
@@ -1215,7 +1420,7 @@ index 77e1a9945..225ec0e78 100644
  
      const int64_t TILE = 256;
      const int64_t MICRO = 16;
-@@ -960,9 +1022,9 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
+@@ -960,9 +1079,9 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
          char * dst_col = (char *) dst->data + i12 * nb2 + i13 * nb3;
  
          // rows[r] is the window's row r base block; the k-slab offset is applied in the unpack
@@ -1227,7 +1432,7 @@ index 77e1a9945..225ec0e78 100644
          }
  
          for (int64_t iir0 = ir0_start; iir0 < ir0_end; iir0 += TILE) {
-@@ -981,19 +1043,19 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
+@@ -981,19 +1100,19 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
                      const int kstart = (int) (k0 / TILED_TILE_K);
                      // activation: one 16-row band (n_src1 <= 16). Unpack the whole k_extent chunk
                      // (all slabs) so each activation row is a long stream, then repack each slab in place
@@ -1250,7 +1455,7 @@ index 77e1a9945..225ec0e78 100644
                          }
                      }
                  }
-@@ -1001,15 +1063,15 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
+@@ -1001,15 +1120,15 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
                  // standard per-slab path: K is stepped in 256-element chunks
                  for (int64_t ib = 0; ib < ne00; ib += TILE) {
                      const int kblk = (int) (ib / TILE);
@@ -1269,7 +1474,7 @@ index 77e1a9945..225ec0e78 100644
                                  (int) (ir0 - iir0), (int) (ir1 - iir1), 1, 0,
                                  ws->acc, TILED_TILE_ROWS);
                          }
-@@ -1024,7 +1086,7 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
+@@ -1024,7 +1143,7 @@ static void ggml_compute_forward_mul_mat_tiled_one_chunk(
      }
  }
  
@@ -1278,7 +1483,7 @@ index 77e1a9945..225ec0e78 100644
  static void ggml_compute_forward_mul_mat_tiled_driver(
          const struct ggml_compute_params * params,
                struct ggml_tensor * dst) {
-@@ -1133,7 +1195,7 @@ static void ggml_compute_forward_mul_mat_tiled_driver(
+@@ -1133,7 +1252,7 @@ static void ggml_compute_forward_mul_mat_tiled_driver(
          const int64_t ir1_start = dr1 * ith1;
          const int64_t ir1_end = MIN(ir1_start + dr1, nr1);
  
@@ -1287,7 +1492,7 @@ index 77e1a9945..225ec0e78 100644
  
          if (nth >= nchunk0 * nchunk1) {
              break;
-@@ -1145,7 +1207,7 @@ static void ggml_compute_forward_mul_mat_tiled_driver(
+@@ -1145,7 +1264,7 @@ static void ggml_compute_forward_mul_mat_tiled_driver(
  
  // src0 type dispatch, shared by the MUL_MAT and MUL_MAT_ID entries: one expert for
  // MUL_MAT_ID (expert_rows != NULL), the full op for MUL_MAT
@@ -1296,7 +1501,7 @@ index 77e1a9945..225ec0e78 100644
  static bool tiled_matmul_dispatch(const struct ggml_compute_params * params,
                                    struct ggml_tensor * dst,
                                    const int32_t * expert_rows,
-@@ -1153,9 +1215,9 @@ static bool tiled_matmul_dispatch(const struct ggml_compute_params * params,
+@@ -1153,9 +1272,9 @@ static bool tiled_matmul_dispatch(const struct ggml_compute_params * params,
                                    int64_t cne1,
                                    char * scratch) {
      if (expert_rows == NULL) {
@@ -1308,12 +1513,16 @@ index 77e1a9945..225ec0e78 100644
      }
      return true;
  }
-@@ -1174,6 +1236,8 @@ static bool ggml_tiled_matmul_type_dispatch(const struct ggml_compute_params * p
+@@ -1174,6 +1293,12 @@ static bool ggml_tiled_matmul_type_dispatch(const struct ggml_compute_params * p
              return tiled_matmul_dispatch<block_q5_K, 32, true,  0, false>(params, dst, expert_rows, cur_a, cne1, scratch);
          case GGML_TYPE_Q4_K:
              return tiled_matmul_dispatch<block_q4_K, 32, true,  0, false>(params, dst, expert_rows, cur_a, cne1, scratch);
 +        case GGML_TYPE_Q4_1:
 +            return tiled_matmul_dispatch<block_q4_1, 32, false, 0, false, block_q8_1, true>(params, dst, expert_rows, cur_a, cne1, scratch);
++#if defined(__AVX512VNNI__) && defined(__AVX512VL__) && defined(__AVX512DQ__)
++        case GGML_TYPE_NVFP4:
++            return tiled_matmul_dispatch<block_nvfp4, 16, false, 0, false, block_q8_0, true>(params, dst, expert_rows, cur_a, cne1, scratch);
++#endif
          case GGML_TYPE_Q3_K:
              return tiled_matmul_dispatch<block_q3_K, 16, false,  4, true>(params, dst, expert_rows, cur_a, cne1, scratch);
          case GGML_TYPE_Q2_K: