Blame for 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

```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.