perf(ggml): reuse ROCmFP4 weights across q=4 columns#534
Draft
davide221 wants to merge 11 commits into
Draft
Conversation
Ports the ROCmFP4/ROCmFPX ggml work (previously lucebox-ggml #36, on the old submodule pointer) into the vendored server/deps/llama.cpp tree. - ROCmFP4/ROCmFPX quant types + CPU reference conversions (ggml/rocmfp4, ggml/rocmfpx) and the ggml type-trait registrations. - CUDA/HIP dequant, copy, getrows, MMVQ vecdot, MMVF, unary and FA paths for the new types. - Fused DS4 hyper-connection op (GGML_OP_DS4_HC) with the register-resident sinkhorn kernel; inert unless emitted by the DS4 fused-decode path. - DS4 SwiGLU split op plumbing for the fused FFN matvec paths. Layered on current main, so the main-side ggml work is preserved (fp64 RoPE reduction, #497 RDNA MMQ tile, LUCE_MMQ_DP_MAX_NE1, MMVQ_MAX_MOE_BATCH_SIZE, fused dual set_rows, raw-span guard). Review fixes on top of #36: - ggml_ftype_to_ggml_type: handle the 11 new ROCmFPX ftypes (dominant-type mapping) so the enum switch is -Wswitch/-Werror clean. - FP6 MMVQ vecdot: pad qs[] to avoid a stack over-read of the last window (bit-identical; the over-read bits were already masked out).
The non-HIP fallback of rocmfp4_get_int_from_codebook_16 / rocmfp4_get_low_int_from_codebook_16 called get_int_from_table_16, which is defined in vecdotq.cuh. TUs that pull in this header without vecdotq.cuh (fattn-chunked.cu reaches it via the fattn dequant chain) failed to compile under nvcc: rocmfp4_hip_codebook.cuh: error: identifier "get_int_from_table_16" is undefined The HIP path never hit this (it uses __builtin_amdgcn_perm), so the ROCm CI and the Strix build stayed green while the sm_86 CUDA build broke. Fix: inline the generic table expander (the generic branch of get_int_from_table_16, verbatim) as a static helper in this header, so the fallback no longer depends on include order. Bit-identical; the HIP hot path is unchanged.
Server-side DeepSeek V4 Flash changes from #494, rebased onto the vendored ROCmFPX ggml tree (no submodule pointer). Byte-identical to the validated #494 server files. Correctness (default ON): - Token-by-token prefill by default; chunked prefill only fits inside the raw SWA ring, so long prompts / multi-turn degraded. Chunked stays available via DFLASH_DS4_CHUNKED_PREFILL for short-prompt benchmarking. - Clear the cache buffer at new-sequence prefill (kv_offset==0) so requests 2..N are byte-stable instead of pooling over leftover compressor state. - Route non-hybrid (all-hot) placement through the HC-complete layer-range path; deepseek4_step's non-hybrid branch is HC-less and generates garbage. - Key the cached decode-attn graph on the flush pattern; the old key collided once the compressed-KV ring filled, reusing a stale-topology graph. - Compressor decode graphs read state/comp-cache through the ggml_set_rows result tensors so the current-step writes are explicit graph dependencies. - Default DeepSeek4 chat prefix when the request has no system message, so the ROCmFPX "Src" GGUF stops behaving like a base model under bare prompts (explicit caller system prompts are preserved). Perf, byte-identical default: - Persistent HC matvec pool (row-split preserves per-row accumulation order) and f16c dot kernels with scalar-order adds. - Per-step decode scalar inputs uploaded in 2 tensor_sets instead of ~430. Opt-in, default OFF (documented as not bit-identical): - DFLASH_DS4_FUSED_DECODE single-graph decode with GGML_OP_DS4_HC. - DFLASH_DS4_FFN_RAW_MMID / DFLASH_DS4_FFN_FUSED_COMBINE accumulation reorder. - DFLASH_DS4_ROCMFPX_HC_GPU GPU HC pre-mix.
Return confidence scores from the fused Markov graph in the existing token-id synchronization, then select q=2/q=3/q=4 from cumulative prefix confidence. Compatible DSpark artifacts enable the policy automatically; artifacts without a confidence head retain the existing EWMA controller. No new deployment configuration is introduced.\n\nValidated on gfx1151: GSM+Math 10/10 at 29.25 tok/s weighted, within 0.8% of fixed q=4, with adaptive behavior retained on low-acceptance prompts.
Remove the no-op single-wave MMVQ shared reduction, add opt-in fixed-K FP2/FP3 and partially-unrolled FP4FAST q=1 kernels, and use AMD v_perm_b32 for packed FP3 decode. A fully build-matched Lucebox3 A/B used ROCm 7.2.4, gfx1151, wave32, -O3, FA_ALL_QUANTS=ON, fused DS4 decode, expert top-k 4, and HC-GPU disabled. Five balanced GSM8K runs per binary improved mean AR throughput from 24.270 +/- 0.018 to 25.664 +/- 0.014 tok/s: +5.744% +/- 0.122% at 95% confidence. All ten measured runs remained 5/5 with exactly 583 output tokens. Generated fused/plain specializations report zero private segment, dynamic stack, SGPR spills, and VGPR spills.
…odex/ds4-q4-weight-reuse
davide221
force-pushed
the
codex/deepseek4-dspark-draft
branch
from
July 17, 2026 20:02
726cc05 to
415432f
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
ROCmFP4_FASTweight fragment once for dense four-column MMVQgfx1151, with no new runtime settingThe dispatch is shape- and backend-based rather than DS4-specific, so other
models using the same quantized dense four-column path can reuse it. Routed
MUL_MAT_ID, AR (ncols_dst == 1), other widths, and other architectures keepthe existing kernels.
Performance
Measured on lucebox2 (
gfx1151, ROCm 7.2.2), pinned to the same 2.9 GHzperformance/highstate for both sides. The workload used fixed DSparkq=4,LUCE_MMVQ_MAX_NCOLS=4, top-k 4, and the accepted ROCmFPX AR specializations.Stable matched window: the first seven identical GSM+Math requests, totaling
896 generated tokens and 263 verification steps:
The candidate's ten requests ranged from 28.7 to 35.0 tok/s. The control
encountered a host-wide latency stall on requests 9-10, so those requests are
not included in the comparison above.
AR remains unchanged at 26.88 tok/s, 5/5 GSM, and 583 output tokens.
Correctness
Kernel checks
A full Release HIP/gfx1151
dflash_serverbuild succeeds. The dedicatedspecializations have no register spills, private segment, or dynamic stack:
The existing generic q=4 kernel metadata is unchanged from control.
Stack
This is intentionally based on
codex/deepseek4-dspark-draft/ #496 becausethat branch owns the fixed-q=4 DSpark verification path used for validation.