feat(dflash): bridge DeepSeek4 DSpark aux heads#496
Draft
davide221 wants to merge 15 commits into
Draft
Conversation
howard0su
approved these changes
Jul 7, 2026
The previous pin lacked ROCMFPX (type 101) support, so the DS4-Flash ROCMFPX GGUF failed to load from a clean checkout of this branch.
deepseek4_step's non-hybrid branch skips hyper-connections entirely and generates garbage on all-hot boxes (Strix). Dispatch to deepseek4_step_layer_range when no MoE hybrid engine is present.
… fused decode Byte-identical, default ON (per-row accumulation order unchanged): - Ds4HcMatvecPool: persistent spin pool for HC fn matvecs and hc_post - f16c dot kernels with scalar-order adds (bit-exact vs scalar) - Ds4DecodeSharedInputs: per-step scalar inputs uploaded in 2 tensor_sets instead of ~430 tiny ones Opt-in, default OFF (reorder fp math, deterministic but not bit-identical to the reference path): - DFLASH_DS4_FFN_RAW_MMID=1 / DFLASH_DS4_FFN_FUSED_COMBINE=1: expert-FFN accumulation reordering - DFLASH_DS4_ROCMFPX_HC_GPU=1: GPU HC pre-mix - DFLASH_DS4_FUSED_DECODE=1: single-graph decode with GGML_OP_DS4_HC GPU hyper-connections, padded compressed-KV spans and an additive row mask (-1e30 underflows to exactly 0 in softmax) so the graph topology is stable across steps Steady-state decode, DS4-Flash ROCMFP2 on Strix gfx1151 (pinned iGPU perf level): 14.2 -> 17.6-18.5 tok/s on the byte-identical default path; ~19 tok/s with fused decode. Requires the iGPU perf-level pin for stable numbers (sclk sags 2900->1473 MHz under sustained decode otherwise).
The fused graph froze the attention window of layers without a mask segment (ratio-128 layers before token 128): they took the non-masked KV-assembly branch, whose shape depends on the live raw-row count, while the slot key assumed shape stability. On slot reuse the model never saw tokens generated after the graph was built, degenerating into repetition. Every layer now gets a mask-bundle segment (n_swa + padded comp rows, padded may be 0) and takes the masked full-ring attention branch. Also: - DFLASH_DS4_FUSED_DEBUG=1: dual-path mode that runs fused + per-layer reference on the same token and prints per-step argmax and max logit diff (this is what localized the bug) - DFLASH_DS4_FUSED_STABLE_GRAPH=1 (default OFF): single topology across flush/non-flush steps via an input-redirected state flush; prerequisite for CUDA/HIP graph replay, which measured as a net loss on gfx1151 Fused decode now: 18.6-19.0 tok/s coherent, per-step argmax matches the reference path; 18.9-19.15 with DFLASH_DS4_FFN_RAW_MMID=1 and DFLASH_DS4_FFN_FUSED_COMBINE=1 stacked.
Compressor state double-buffers and compressed-KV rows retained the previous request's values (counters reset, contents did not), so the first flush windows of each request pooled over leftover state and outputs could drift by a token or two from the 2nd request on. Clearing the cache buffer at kv_offset==0 makes every subsequent request byte-stable (requests 2..N now produce identical output).
The generic sinkhorn kernel spilled its c[] matrix to scratch memory, costing 95us x 86 launches per decoded token (~8ms/token). The unrolled template<NHC=4> variant keeps it in registers (5.3us measured). DS4-Flash fused decode on Strix gfx1151: 19.1 -> 23.0 tok/s.
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.
Contributor
Author
|
Validation update (2026-07-09): the converter/loader bridge in this PR was validated on lucebox2 with the converted DSpark drafter artifact. In the follow-on DeepSeek4 runtime worktree, DSpark loads successfully, prefill feature capture and generation are wired, the batched-verify causality bug is fixed, and causal batched output matched exact sequential verification in paired probes. Adaptive EWMA width and fused batched verify are working; fused verify reached approximately 20 tok/s (peak 20.1 tok/s), versus roughly 13–15 tok/s for the earlier dynamic causal path. Exact sequential verification measured 1.78–1.91 accepted candidates per 4-candidate step. These runtime changes are follow-on work in the hub-494 experiment worktree and are not claimed as already included in PR 496. |
…to dspark-publish
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.
Contributor
Author
|
Updated DSpark adaptive width in
Validation on lucebox2 / gfx1151 with the exact committed source:
The confidence policy therefore retains essentially all high-acceptance throughput while adapting verification width on harder prompts. |
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
2 participants
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
convert_dflash_to_gguf.pyto read DSpark aux heads from DeepSeek4/MTP safetensors shardsmtp.2.markov_head.*andmtp.2.confidence_head.proj.*into the existingdflash.dspark.*GGUF contractNotes
This is a draft bridge PR. It does not yet make the production DeepSeek4 layer-split backend consume DSpark during decode. It makes the DS4/MTP aux artifact consumable by the existing GGUF loader/runtime so the next PR can wire the runtime path cleanly.
The DS4 artifact observed on lucebox4 has no confidence bias tensor, so the converter writes a zero bias when only
mtp.2.confidence_head.proj.weightis present.Validation
python3 -m py_compile server/scripts/convert_dflash_to_gguf.pygit diff --checkCould not run the functional Python micro-test locally because this macOS temp clone lacks numpy, and
uv runcannot resolve the repo CUDA Torch wheel on macOS arm64.