Skip to content

fix(ds4): correct layer-split HC handoff and per-device state#507

Draft
Graffioh wants to merge 9 commits into
Luce-Org:codex/ds4-rocmfpx-serverfrom
Graffioh:codex/fix-ds4-hc-shard-handoff
Draft

fix(ds4): correct layer-split HC handoff and per-device state#507
Graffioh wants to merge 9 commits into
Luce-Org:codex/ds4-rocmfpx-serverfrom
Graffioh:codex/fix-ds4-hc-shard-handoff

Conversation

@Graffioh

@Graffioh Graffioh commented Jul 11, 2026 •
edited
Loading

Copy link
Copy Markdown
Contributor

Summary

Fixes correctness and device-isolation issues in the DeepSeek4 HC layer-split path:

  • hand off the updated full HC state between local shards
  • scope HC weights and cached decode graphs to each shard
  • keep independent CUDA/HIP HC scratch and F16 mirrors per GPU device
  • select the correct device at layer-range shard entry
  • add coverage for HC handoff, per-shard runtime state, and per-device scratch

Stack

Stacked on #503. The HC-complete layer-range implementation fixed here is introduced by #503. Retarget this PR to the branch where #503 lands once merged.

Implementation

Local layer split now follows the full HC flow:

token embeddings -> shard 0 -> full HC state -> shard 1 -> ... -> final shard

Runtime caches are owned by each shard and keyed by weights, context, backend, device, layer range, and output ownership. HC scratch is stored in independently locked slots indexed by CUDA/HIP device.

Validation

  • git diff --check and C++ syntax checks passed
  • Release CUDA H200 builds passed for test_deepseek4_unit and dflash_server
  • tested on one H200 in both single-shard and virtual two-shard configurations
  • virtual split [0,22) / [22,43) on cuda:0 completed sequential requests without HC, CUDA, crash, or abort errors
  • the virtual split validates HC handoff and per-shard cache reuse
  • real multi-device device 0 -> device 1 -> device 0 validation on lucebox3 is still required
  • a barrier-based concurrent scratch test and complete shape-capacity coverage for all scratch buffers are still required before marking ready

davide221 and others added 4 commits July 8, 2026 17:23
@davide221
feat(ggml): vendor ROCmFPX quant support for DeepSeek V4 Flash
b9c0069 
Ports the ROCmFP4/ROCmFPX ggml work (previously lucebox-ggml Luce-Org#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, Luce-Org#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 Luce-Org#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).
@davide221
fix(ggml): make ROCmFP4 codebook self-contained on CUDA
a5f1146 
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.
@davide221
fix(deepseek4): ROCmFPX byte-safe decode + fused-decode path
42f1e10 
Server-side DeepSeek V4 Flash changes from Luce-Org#494, rebased onto the vendored
ROCmFPX ggml tree (no submodule pointer). Byte-identical to the validated
Luce-Org#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.
feat(deepseek4): add opt-in expert top-k knobs
6c7ee19
@Graffioh Graffioh changed the title fix(deepseek4): correct HC shard handoff and scope runtime caches per shard fix(ds4): correct HC shard handoff and scope runtime caches per shard Jul 11, 2026
@davide221

davide221 commented Jul 11, 2026

Copy link
Copy Markdown
Contributor

We plan to run two AMD GPUs in the same HIP process. g_scratch is currently shared across devices, so the second shard may reuse scratch memory allocated on the first GPU. Please keep one scratch buffer per HIP device, set the correct device in HC worker threads

@davide221
perf(ggml): specialize ROCmFPX AR decode on gfx1151
e943f2a 
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.
@Graffioh
Graffioh force-pushed the codex/fix-ds4-hc-shard-handoff branch 2 times, most recently from bf8dc45 to ecf55b1 Compare July 13, 2026 15:13
@Graffioh Graffioh changed the title fix(ds4): correct HC shard handoff and scope runtime caches per shard fix(ds4): correct layer-split HC handoff, per-device state, and greedy decode in local path Jul 13, 2026
@Graffioh Graffioh mentioned this pull request Jul 13, 2026
Closed
@Graffioh
Graffioh force-pushed the codex/fix-ds4-hc-shard-handoff branch 2 times, most recently from 581817d to b6a14e6 Compare July 14, 2026 19:48
This was referenced Jul 15, 2026
Open
Open
@Graffioh
Graffioh force-pushed the codex/fix-ds4-hc-shard-handoff branch from b6a14e6 to 0441ac2 Compare July 16, 2026 09:16
@Graffioh
Graffioh force-pushed the codex/fix-ds4-hc-shard-handoff branch from 0441ac2 to 4ae83dd Compare July 16, 2026 12:50
@Graffioh Graffioh mentioned this pull request Jul 16, 2026
Closed
@Graffioh Graffioh changed the title fix(ds4): correct layer-split HC handoff, per-device state, and greedy decode in local path fix(ds4): correct layer-split HC handoff and per-device state Jul 16, 2026
@davide221
davide221 force-pushed the codex/ds4-rocmfpx-server branch from e943f2a to d5fae56 Compare July 17, 2026 05:33
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

ds4flash

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@Graffioh @davide221