handoff/HANDOFF.md

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

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):

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.