Skip to content

perf(gguf): sm_120 dispatch thresholds + upstream int8-MMA MMQ port (Ornith) - #196

Open
lucaspirola wants to merge 17 commits into
FlashML-org:mainfrom
lucaspirola:ornith-sm120-gguf-mmq-phase2
Open

perf(gguf): sm_120 dispatch thresholds + upstream int8-MMA MMQ port (Ornith)#196
lucaspirola wants to merge 17 commits into
FlashML-org:mainfrom
lucaspirola:ornith-sm120-gguf-mmq-phase2

Conversation

@lucaspirola

Copy link
Copy Markdown

What

Two related optimizations for Ornith-1.5-35B-Q4_K_M on RTX 5080 (sm_120,
Blackwell), on top of the existing Ornith INT4 262K-context serving path:

  1. Arch-aware GGUF dispatch thresholds — the dense/MoE dequant-vs-MMQ
    crossover points measured on Ada don't hold on sm_120; make them
    architecture-gated instead of hardcoded.
  2. int8-tensor-core MMQ port — the vendored GGUF kernels (llama.cpp
    b2899 via vLLM/sgl-kernel) are DP4A-only with zero tensor-core use.
    Upstream llama.cpp master has since rewritten Q4_K/Q6_K MMQ around
    int8 tensor-core MMA tiles (turing_mma, sm_75+). This PR vendors
    that upstream kernel verbatim (kernel/csrc/gguf_mmq/, llama.cpp
    eab8ee41) and wires it in for both dense projections and grouped
    MoE experts, gated to sm_120 for now.

FREETOKEN_GGUF_DISABLE_MMA=1 forces the old DP4A/dequant path (debugging,
or a toolchain that can't build the extension) — used below as the control
in every comparison.

Gains

Synthetic kernel benchmarks (real Ornith-1.5-35B-Q4_K_M tensors, sm_120, median of CUDA-event timings):

shape rows/tokens int8-MMA MMQ DP4A MMQ (old) dequant+cuBLAS speedup vs DP4A
Q4_K attn_q [8192x2048] (dense) 8192 rows 1.79 ms 22.9 ms 2.40 ms 12.8x
Q6_K lm_head [248320x2048] (dense) 2048 tokens 17.5 ms 262 ms 19.2 ms 15.0x
Q4_K gate_up (MoE, E=256 top-8) 8192 tokens 4.16 ms 23.2 ms n/a 5.6x
Q6_K down (MoE, E=256 top-8) 8192 tokens 4.90 ms 15.3 ms n/a 3.1x

MMA beats dequant+cuBLAS too (the prior best option at large batch), not just
the DP4A kernel it replaces — e.g. attn_q 1.79ms vs 2.40ms, lm_head 17.5ms vs
19.2ms.

Live end-to-end A/B, full 262K-context serve config, same box, same prompt, minutes apart:

Config: --attention-backend triton --kv-cache-dtype q4_0 --num-tokens 262144 --kv-reserve-tokens 262144 --max-seq-len-override 262144 --max-running-requests 1 --moe-backend offload --moe-cache-auto --max-prefill-length 8192. Prompt: 28K words of seeded non-repetitive text
(~50K tokens, 6+ chunked prefill batches) with three distinct needles at
10/50/90% depth, greedy decode.

int8-MMA (this PR) DP4A/dequant (FREETOKEN_GGUF_DISABLE_MMA=1)
wall time 13.88 s 21.94 s
effective prefill rate ~4,700 tok/s ~2,700 tok/s
needles recovered 3/3 exact 3/3 exact
tracebacks 0 0

~1.75x prefill, ~1.58x end-to-end wall time, identical answers on the
same hardware and prompt.

Dispatch thresholds (sm_120 vs the existing Ada-tuned defaults, unchanged
off sm_120):

  • dense dequant-vs-MMQ crossover: 24 rows (was 32) — Q4_K attention shapes
    cross at 24 (0.0645ms dequant vs 0.0778ms MMQ); 16 would regress the Q6_K
    lm_head, where MMQ still wins at 16.
  • MoE grouped-MMQ-vs-vec crossover: 16 tokens (was 32) — 0.314ms MMQ vs
    0.324ms vec at 16 tokens, widening to 0.382 vs 0.475 at 24.
  • Live-reverified on real Ornith tensors: 24 rows now 0.069ms via dequant vs
    0.081ms with the old threshold.

How

  • python/freetoken/kernel/csrc/gguf_mmq/ — llama.cpp master's CUDA MMQ
    vendored verbatim (mmq/mma/load-tiles/vec-dot/configs/quantize/mmid + the
    ggml headers it needs). mmq_ext.cu is the only hand-written file: backend
    shims (device info, a torch-allocator-backed pool, error/abort plumbing)
    and the torch bindings for the dense (ggml_mul_mat_a8_mma) and grouped-MoE
    (ggml_moe_a8_mma) entry points. Only Q4_K/Q6_K mul_mat_q cases are
    instantiated to keep JIT compile time down.
  • layers/gguf.pydequant_gemm_min_rows(cc) (arch-gated threshold),
    _use_mma_mmq dispatch gate (sm_120 + supported type + successful JIT
    build), wired into fused_mul_mat_gguf.
  • moe/fused_gguf.pymmq_min_tokens(cc) (arch-gated threshold),
    _use_mma_moe gate (adds a block-alignment check on the padded expert
    slot stride), wired into _moe_matmul for the 320–16384 token band
    (below 320, DP4A wins — per-expert tiles waste work at small batches;
    MMVQ still owns decode).
  • kernel/gguf.py_mma_module() lazy JIT loader, mma_mmq_supported().

Testing

  • Numerics: MMA output verified against the dequantized reference on real
    Ornith tensors (rel error <= 0.013, on par with the existing DP4A kernel)
    and against gguf-py on random-but-safe packed bytes
    (tests/kernels/test_gguf_mma.py). MoE broadcast (gate/up) and gather
    (down) forms verified against the per-expert dense reference and the
    existing vec kernel, end-to-end through fused_experts_gguf (rel ~1e-3).
  • tests/kernels/test_gguf_dispatch.py — pure threshold-function tests for
    both archs, plus CUDA dispatch-branch tests with faked device capability.
  • Kernel + benchmark suites: 336 passed, 1 skipped.
  • Full non-slow suite (-m "not slow", excluding e2e/server): failure set
    unchanged from clean main (same 7 pre-existing failures).
  • Live 262K-context A/B above, including a hostile (non-repetitive,
    multi-needle) prompt specifically to catch any tile/row-range corruption a
    compressible-filler prompt could hide.
  • ruff check clean on all changed files.

Falls back cleanly to the existing DP4A/dequant path off sm_120, on a build
failure, or with FREETOKEN_GGUF_DISABLE_MMA=1; existing sm_89/Ada behavior
is unchanged (verified via the A/B, not just by inspection).

probe and others added 17 commits August 24, 2026 05:52
…ated e2e

Adds the `laguna` GGUF architecture: hybrid full/SWA attention with per-layer
query-head counts (48 full / 72 SWA on S), QK-norm, per-layer-type rope (YaRN
partial-dim on full layers, plain on SWA), a per-head softplus attention output
gate, and sigmoid+correction-bias MoE routing with one always-on shared expert.
Reference semantics follow llama.cpp `src/models/laguna.cpp`.

Unsloth/poolside laguna checkpoints quantize per tensor, so this also generalizes
the GGUF plumbing:

- six new ggml types (Q4_K, Q5_K, IQ1_S, IQ2_XXS, IQ3_XXS, IQ4_XS) wired into the
  dequant tables and the mmvq/mmq/dequant dispatch sets
- a "gguf" expert-bank format whose per-layer quant types vary: flat padded
  [E, stride] host banks plus a new `expert_stride_bytes` argument threaded through
  the vendored moe_vec launchers (0 = previous dense behaviour)
- q/k/v kept as separate projections, since a layer may quantize attn_v
  differently from attn_q/k
- deferred GGUF linears materialized from the file's tensor table at conversion
  time, before the engine collects the state dict

Verified by unit tests (81 green) and by real-file probes: full tensor-name
coverage, 529-param weight iteration, and an expert-bank matmul within 0.5% of the
gguf-py reference. NOT yet verified end to end -- no forward pass on real weights
and no comparison against llama.cpp; the host used for development lacked the RAM
to hold S's expert banks. See tasks/laguna-handover.md.

Co-Authored-By: Claude <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Je2rjENB9qJiiRGmNcnAct
Stores the KV cache in 8 bits plus an fp16 scale per 32 elements along head_dim
(1.0625 bytes/element vs 2), freeing VRAM for the MoE expert cache. Two schemes
share the scale tensor, store kernel and dequant path -- q8_0 (int8, s = max/127)
and fp8_e4m3 (s = max/448) -- so comparing them is a flag change, not a port.

- kvcache/quant.py: KVQuantSpec (storage dtype, block 32, torch reference
  quantize/dequantize, effective bytes/element)
- kernel/triton/kv_quant.py: store kernel computing per-block max-abs and writing
  the quantized buffer + scales
- kernel/triton/attention.py: dequant inside the four attention kernels behind a
  QUANT constexpr (0 compiles the existing bf16 path unchanged); the scale varies
  along head_dim, the reduction dim, so K/V dequantize to bf16 before the dot
- kvcache pools: parallel scale buffers, k_scale()/v_scale(), rebuild() realloc,
  unit_bytes()/kv_cost() accounting on effective bytes
- server/args.py, engine: --kv-cache-dtype {auto,q8_0,fp8_e4m3} with gating
  (triton backend only, head_dim % 32 == 0, supported pool families)

Tests: 53 new (round-trip vs torch reference, quantized attention vs the bf16
reference, pool sizing and hot rebuild, flag gating); the existing 33 triton
attention tests still pass.

Step 9 of tasks/todo.md is NOT done: no end-to-end validation on this host --
needle-in-246k, perplexity vs bf16, and the real expert-slot / tok-s gain are
unmeasured, so q8_0 vs fp8_e4m3 as the default is still an open question.

Co-Authored-By: Claude <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Je2rjENB9qJiiRGmNcnAct
The handover was written before the KV-cache quantization commit landed and still
described it as uncommitted. Restates the relationship instead: the two workstreams
meet at --kv-cache-dtype, and both open validation questions want the same big-host
run (one loaded model, one long context).

Co-Authored-By: Claude <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Je2rjENB9qJiiRGmNcnAct
…idate XS

moe_vec_q indexes experts via blockIdx.z, so tokens*top_k was capped by CUDA's
65535-row grid-z limit and a 16k-token x top-8 prefill chunk overflowed it (reported
asynchronously as "device not ready"). fused_experts_gguf now chunks calls to
min(65535, 16384) rows; the 16384 tie also bounds transient VRAM.

Enables Q3_K and IQ2_S (block 256, 110/82 bytes), the two remaining types the
APEX-Mini XS build uses; both already had CUDA dispatch, table entries only.

Validation on Laguna-XS-2.1-APEX-I-Mini.gguf: NIAH 3/3 at 250k tokens, decode
157-162 tok/s at 64k ctx and 21-23 tok/s at 262k. Handover updated with the S-host
runbook, the offload-cache port hygiene, and the hybrid-backend fix path.

Co-Authored-By: Claude <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Je2rjENB9qJiiRGmNcnAct
The handover mis-stated the hybrid/cpu enablement as "~1 file". It is a C++
SIMD kernel port (vec-dot for every ggml type in cpu_moe_ext.cpp) plus a
Python resolver, and only pays off where CPU bandwidth beats PCIe -- which this
box's own bench-bw profile says it does not.

Co-Authored-By: Claude <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Je2rjENB9qJiiRGmNcnAct
Half the bytes of the 8-bit formats (0.5625 vs 1.0625 B/element incl.
the fp16 per-block scale), two signed nibbles per uint8 byte. Shares the
per-block scale geometry, store kernel and Triton dequant path with q8_0/
fp8_e4m3; only payload layout and max-magnitude divider differ.

- quant.py: INT4 spec + packed quantize/dequantize (low nibble = even
  element, high = odd), export in __all__.
- kernel/triton/kv_quant.py: EPB==2 nibble-pack store path.
- kernel/triton/attention.py: EPB constexpr through _load_kv, the three
  attention kernels and _kv_scale_args (logical head_dim stays element
  space; only byte addressing divides by epb).
- pools: storage slab last dim halves when packed; scales key off logical
  D // BLOCK; store routing preserves logical head_dim.
- args/config/base: CLI help, config comment, cost model (0.5625 B/E).
- tests: int4 in the kernel/pool/gating parametrizations, physical-layout
  and nibble-packing assertions.

E2E validated 2026-08-24 on Laguna-XS-2.1-APEX-I-Mini at 262144 tokens:
KV = 4.65 GiB (vs fp8 8.79), NIAH 3/3 exact passcode at 10/50/90% depth.
107 tests pass.

Co-Authored-By: Claude <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01NvYubTfxXxRYCmz6XapVRC
The native fp32 -> float8e4nv downcast does not round to nearest
everywhere: on sm_89 triton lowers it as a truncating fp32 -> fp16 ->
e4m3 double-round, so values just above a grid midpoint collapse
downward and disagree with the RNE torch reference (~0.4% of elements).
Round explicitly with round_e4m3 before clamping, mirroring the int
branch's rounding. Fixes the two fp8 reference-equality tests on sm_89.

Co-Authored-By: Claude <noreply@anthropic.com>
Stores the KV cache in 8 bits plus an fp16 scale per 32 elements along head_dim
(1.0625 bytes/element vs 2), freeing VRAM for the MoE expert cache. Two schemes
share the scale tensor, store kernel and dequant path -- q8_0 (int8, s = max/127)
and fp8_e4m3 (s = max/448) -- so comparing them is a flag change, not a port.

- kvcache/quant.py: KVQuantSpec (storage dtype, block 32, torch reference
  quantize/dequantize, effective bytes/element)
- kernel/triton/kv_quant.py: store kernel computing per-block max-abs and writing
  the quantized buffer + scales
- kernel/triton/attention.py: dequant inside the four attention kernels behind a
  QUANT constexpr (0 compiles the existing bf16 path unchanged); the scale varies
  along head_dim, the reduction dim, so K/V dequantize to bf16 before the dot
- kvcache pools: parallel scale buffers, k_scale()/v_scale(), rebuild() realloc,
  unit_bytes()/kv_cost() accounting on effective bytes
- server/args.py, engine: --kv-cache-dtype {auto,q8_0,fp8_e4m3} with gating
  (triton backend only, head_dim % 32 == 0, supported pool families)

Tests: 53 new (round-trip vs torch reference, quantized attention vs the bf16
reference, pool sizing and hot rebuild, flag gating); the existing 33 triton
attention tests still pass.

Step 9 of tasks/todo.md is NOT done: no end-to-end validation on this host --
needle-in-246k, perplexity vs bf16, and the real expert-slot / tok-s gain are
unmeasured, so q8_0 vs fp8_e4m3 as the default is still an open question.

Co-Authored-By: Claude <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Je2rjENB9qJiiRGmNcnAct
Half the bytes of the 8-bit formats (0.5625 vs 1.0625 B/element incl.
the fp16 per-block scale), two signed nibbles per uint8 byte. Shares the
per-block scale geometry, store kernel and Triton dequant path with q8_0/
fp8_e4m3; only payload layout and max-magnitude divider differ.

- quant.py: INT4 spec + packed quantize/dequantize (low nibble = even
  element, high = odd), export in __all__.
- kernel/triton/kv_quant.py: EPB==2 nibble-pack store path.
- kernel/triton/attention.py: EPB constexpr through _load_kv, the three
  attention kernels and _kv_scale_args (logical head_dim stays element
  space; only byte addressing divides by epb).
- pools: storage slab last dim halves when packed; scales key off logical
  D // BLOCK; store routing preserves logical head_dim.
- args/config/base: CLI help, config comment, cost model (0.5625 B/E).
- tests: int4 in the kernel/pool/gating parametrizations, physical-layout
  and nibble-packing assertions.

E2E validated 2026-08-24 on Laguna-XS-2.1-APEX-I-Mini at 262144 tokens:
KV = 4.65 GiB (vs fp8 8.79), NIAH 3/3 exact passcode at 10/50/90% depth.
107 tests pass.

Co-Authored-By: Claude <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01NvYubTfxXxRYCmz6XapVRC
The native fp32 -> float8e4nv downcast does not round to nearest
everywhere: on sm_89 triton lowers it as a truncating fp32 -> fp16 ->
e4m3 double-round, so values just above a grid midpoint collapse
downward and disagree with the RNE torch reference (~0.4% of elements).
Round explicitly with round_e4m3 before clamping, mirroring the int
branch's rounding. Fixes the two fp8 reference-equality tests on sm_89.

Co-Authored-By: Claude <noreply@anthropic.com>
Add native NVIDIA Nemotron 3 Super NVFP4 and Qwen3.5-MoE GGUF loading, including Ornith tokenizer/layout handling. Keep all expert compute on GPU under WSL pin limits through pageable miss staging. Speed GGUF prefill with grouped expert MMQ and transient dense dequantized GEMM, fix the grouped MMQ live-prefix bounds check, and make auto expert sizing reserve explicit KV geometry.
Load packed KV bytes once before nibble interleave, tune Ornith's split-K decode geometry for sm_89, pipeline D=256 prefill attention, and reuse the in-tree fused renormalized router. This raises measured 170K decode from 7.58 to 33.67 tok/s and cuts the progressive 140K-to-170K extension TTFT from 245.2s to 113.4s on RTX 2000 Ada WSL.
Phase 2a: make the dense/MoE GGUF dispatch crossovers architecture-aware.
sm_120 (RTX 5080) measured differently from the Ada-tuned defaults:
dequant_gemm_min_rows 24 (was 32), mmq_min_tokens 16 (was 32). Both keep
the existing constants on every other architecture.

Phase 2b: vendor llama.cpp master's int8-tensor-core MMQ (mul_mat_q,
turing_mma path) verbatim into kernel/csrc/gguf_mmq/, replacing the
vendored DP4A-only kernels for Q4_K/Q6_K on sm_120. mmq_ext.cu is the
only hand-written file: backend shims (device info, torch-allocator
pool, error/abort) plus torch bindings for the dense and grouped-MoE
entry points. Wired into fused_mul_mat_gguf (dense, rows > _MMVQ_SAFE)
and _moe_matmul (320-16384 tokens); FREETOKEN_GGUF_DISABLE_MMA=1 forces
the DP4A/dequant fallback for debugging or a toolchain that can't build
the extension.

Measured on real Ornith-1.5-35B-Q4_K_M tensors: Q4_K attn_q 8192 rows
1.79ms (MMA) vs 2.40ms (dequant+cuBLAS) vs 22.9ms (DP4A); MoE gate_up
@8192 tokens 4.16ms vs 23.2ms DP4A. Live A/B at the production 262K
serve config (hostile 50K-token 3-needle prompt): MMA 13.88s wall vs
21.94s with the port disabled, identical (3/3 exact) answers.

Numerically verified against the dequant reference and the existing
DP4A kernels on real tensors and gguf-py cross-checks; full non-slow
test suite failure set unchanged from clean main.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EULQTS73xnKEKmRqsWHKYg
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant