Repository navigation
feat(ds41): two-box expert-parallel serving - #826
Conversation
CI failed on
|
b34c688 to
80e5ec2
Compare
CI failed on
|
80e5ec2 to
d5ba225
Compare
CI failed on
|
The verify step of DS4.1 runs many short kernels per layer; these cut the launch count without changing a bit of the output: - one launch for an HC sub-block boundary (norm, mix, Sinkhorn split, post, collapse and the next norm), and for the router's select and weights over the same logits; - RMS norm, rope and the F16 copy in one launch (HIP only: nvcc's fast math contracts rope_norm's rotation in a way another kernel cannot be guaranteed to repeat), and a rope that also writes the q8_1 input of the grouped output projection; - the masked sink softmax in one launch, and one split-KV flash launch for all single-token verify lanes of a layer; - a skinny-K single-column F16 matvec, and wide Q4_K products over many rows routed to MMQ on RDNA4; - one expert grouping per MoE layer for gate, up and down, and a q8_1 activation memo shared across reshapes; - column-range Q8_0 products that keep the whole row's lane mapping, so a slice of a weight matrix gives the same sums as the full product; - a split's batched peer copies as one copy kernel, an in-place cluster all-reduce node, and an argmax pair/pick that finds the greedy token of a vocabulary split across ranks exactly. Each fused kernel repeats the replaced kernels' arithmetic in the same order; test_ds4_fused_ops_cuda checks them bit for bit. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
… owners - The owner route mask, the route weighting and the route sum run in one launch, and the clamped routed SwiGLU in one launch; the masked combine rounds its products and its sums separately, as the unfused ops do. - The shared-expert slice is queued ahead of the owner join. - Opt-in (LUCE_MOE_Q8_HANDOFF=1): the secondary expert owner receives its FFN input as the q8_1 rows its MMVQ reads, quantized once on the primary device. That moves about 3.6x fewer bytes over PCIe and the secondary skips its own quantize; the products are unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
AdaptiveSpecWidth seeded every width with one device's cost curve and replaced an entry only when it ran that width, so a narrow width measured on faster hardware could price every wider one out for good. Unmeasured widths are now priced from the seed shape scaled by the nearest measured width, never below a measured narrower one; graph-build outliers are clipped out of the samples, and a width's first live sample is held to 1.25x its unmeasured price once another width has been measured. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
A port of maikzz32's lucebox-halo-cluster runtime onto the in-graph collective of #700 and the cold-owner-None hybrid path of #701: - rank bootstrap over TCP (build, model and placement identity checked), a control channel on which rank 0 decides and the workers follow, an RCCL communicator for prefill-sized reductions, a worker main, and the expert placement JSON with its uniform and balanced builders (scripts/cluster/build_expert_placement.py); - fast_reduce, the decode all-reduce for discrete GPUs: raw RDMA verbs between two ranks with stream write/wait flags instead of a host round trip, a zero-copy bf16 wire, and exchange kernels that publish their own completion flags. Its switches treat 0 as off, so cluster mode can turn them on by default; - the RCCL environment of a rank, now with NCCL_PROTO=Simple, the only protocol that runs reliably on the R9700; - server flags --cluster-rank, --cluster-size, --cluster-head, --cluster-ifname, --cluster-ib-hca, --cluster-gid-index, --cluster-expert-placement, --cluster-shared-expert, --cluster-allreduce-dtype, --cluster-timeout-ms, --cluster-verify-hash. Built with -DLUCE_CLUSTER=ON (RCCL and libibverbs). Without it the RCCL communicator and fast_reduce compile as unavailable stubs, so the default build links and a rank falls back to RCCL. test_cluster_unit covers the protocol, placement and control channel in both builds. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Each rank keeps the single-box hybrid tiers (hot experts on the discrete GPU, the rest on the integrated GPU) for its share of the routed experts, and the two ranks reduce every MoE layer over the cluster exchange: - hybrid FFN exchange through the integrated GPU (it sums the cold and imported rows and sends one export), a lean exchange, and one export per layer; - attention heads split across the ranks, and a rank that loads only its attention slice; - greedy verify head and DSpark draft head with their vocabulary split across the ranks (exact argmax over both halves), and the DSpark drafter split across both ranks; - workers that skip the drafter they never use, a drafter swap for long prompts with chunks agreed across ranks, and a pipelined prefill only when every rank pipelines; - prefill that keeps the device hot and cold join and reduces in the HC post graph, with a bf16 wire for prefill-sized reductions; - the drafter's context window and the verify's window rows kept on the device, batched verify uploads, one wait per step for small transfers, and the Engram projection run inside the FFN waits. Cluster mode (--cluster-size > 1) installs the measured configuration as environment defaults (cluster_launch_env in launch_profiles.h) before a launch profile's, so a two-box run needs only the operator's inputs: expert budget, placement and router bias. Each default is a switch that 0 turns off, and an explicitly set variable always wins. A cluster rank never retries a failed prefill or a failed drafter swap-in on its own: the ranks run every chunk together. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
d5ba225 to
9f350f0
Compare
…U-DS4 A shared expert's gate and up products (Q8_0, several columns in a verify) run as two batch-invariant RDNA4 launches and SwiGLU-DS4 as a third. mul_mat_vec_q8_0_rdna4 gets a GLU variant that reads both matrices with one load of each activation fragment, keeps each product's partials and reduction exactly as the single kernel does, and writes ggml_cuda_op_swiglu_ds4_single(gate, up) in its epilogue: the value the two products and the GLU kernel give, in one launch. The graph matcher also finds the pair with weight views between the products (the cluster shared expert's row slices; views launch nothing). ggml_cuda_mmvq_rdna4_glu_pair gates it: RDNA4, batch-invariant products, 2..8 columns, a Q8_0 pair on the same activations, a non-swapped SwiGLU-DS4 that reads exactly their outputs, no biases. LUCE_CUDA_MMVQ_Q8_RDNA4_GLU=0 keeps three launches. The pair has its own launcher on one block shape (two rows per wave, four waves) whose partials of both products fit the LDS at eight columns: seven new kernel instances, the dispatch table unchanged. test_ds4_fused_ops_cuda: the pair against the products kept apart, and a whole-row-lane column slice against the zero-padded product, bit for bit at 1..8 columns. 0 failures on gfx1201 and gfx1151; a kernel trace on gfx1201 shows the fused launch at 2, 4, 5 and 8 columns. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…loop The grouped MoE MMVQ (LUCE_MMID_GROUPED) ran ROCmFP2 experts on the generic loop, where every iteration waits for a row's scale byte before the next loads leave: on the Strix Halo iGPU the cold experts read at ~100 GB/s. The per-expert kernel already has a prefetching loop for this (mul_mat_vec_q_moe_fp2_prefetch: the loads of several K iterations first, then the dot products in the same kbx order into the same accumulators, so the result is bit-identical). The grouped kernel now uses it too, and both kernels read one switch: on by default on gfx1151, LUCE_CUDA_MMVQ_MOE_FP2_PREFETCH=0/1 overrides (cluster mode sets 1). Two boxes, DS4.1, with the gate/up launch (previous commit), 2 x 3 runs against both off, same text on both prompts: step 42.1-42.4 -> 41.5-41.8 ms (code), 42.1-43.9 -> 41.4-41.7 ms (2K); code 100.7-101.2 -> 102.1-102.7 tok/s, 2K 55.8-58.8 -> 59.3-60.0. test_deepseek4_mmid_grouped_cuda passes on gfx1201 and gfx1151 with the switch at its default and forced on. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The two commits above built on HIP only. ggml_cuda_mmvq_rdna4_glu_pair sat
in the HIP-only RDNA4 block while ggml-cuda.cu calls it on every build: it
is now defined everywhere and answers false outside HIP. nvcc rejects the
grouped kernel's unbraced else before the K loop ("expected a }"): the loop
now sits in braces, the way the per-expert kernel writes it. No change on
HIP.
Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
CI failed on
|
ggml_backend_cuda_mul_mat_whole_row_lanes promises the padded product's bits only where the lanes are implemented (HIP); other kernels compute the range's product as usual. The column-slice check now compares bits on HIP and the value to 1e-5 relative elsewhere: on the DGX GB10 (sm_121) the slice and the padded product differ in the last bit (-1.77656889 vs -1.77656901), which failed the CUDA GPU job. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
- A cluster worker runs to completion inside load_model and has no HTTP server; main() went on to call run() on the null server, so every clean shutdown of rank 1 crashed. It now returns the worker's exit code. - --cluster-selftest and --cluster-replicate-hot were documented, wired to the head, the worker and the placement, but rejected by the parser. - The cluster lines of --help sat inside --ds4-prefill's description. - LUCE_CLUSTER_FAST_REDUCE_STALL_S=0 made every wait a stall; 0 now turns the watchdog off, and the defaults comment says which entries are sizes. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
davide221
left a comment
davide221
left a comment
There was a problem hiding this comment.
Code review
Found 1 issue:
- HeadHooks fallback RNG is seeded once and ignores later request seeds · likely (
server/src/cluster/cluster_decision_hooks.cpp:62-64)
Note: the diff of server/src/cluster/fast_reduce.cpp, server/src/deepseek4/deepseek4_dspark_spec.cpp, server/src/deepseek4/deepseek4_fused_verify.inc, server/src/deepseek4/deepseek4_graph.cpp, server/test/unit/test_adaptive_spec_width.cpp, server/test/unit/test_cluster_unit.cpp, server/test/unit/test_deepseek4_unit.cpp, server/test/unit/test_ds4_fused_ops_cuda.cpp, server/test/unit/test_launch_policy.cpp was too large to include; the reviewers could read those files but not see what changed in them; tick the box in the bot's comment below to review them.
Reviewed a49bc3ef · 2 candidate issues checked, kept at ≥70% confidence · add the code-review label to review again
| static thread_local std::mt19937_64 rng{cfg.seed != 0 ? cfg.seed : 0x5eedULL}; | ||
| std::vector<int32_t> no_history; | ||
| token = common::sample_logits(logits, n_vocab, cfg, no_history, rng); |
There was a problem hiding this comment.
HeadHooks fallback RNG is seeded once and ignores later request seeds · likely
In HeadHooks::decide_next_token, when the caller passes token < 0 and cfg.temp > 0, the sampler RNG is a static thread_local std::mt19937_64 initialized from cfg.seed on the first call only. Every subsequent request — including ones with a different seed, or the same seed that should reproduce a stream — reuses the first request's generator state, so sampled tokens are wrong relative to the requested seed and are not reproducible across requests/restarts. Construct the RNG per request (or reseed when cfg.seed changes).
Checked (70%): cluster_decision_hooks.cpp:62 uses static thread_local std::mt19937_64 rng{cfg.seed ...}, so the seed from cfg is applied only on the first call and every later request's cfg.seed is ignored, exactly as claimed. However, the only current callers (deepseek4_backend.cpp:4860/4927/5198) all pass an already-sampled token >= 0, so the path is a documented fallback (cluster_decision_hooks.h:84-86) not exercised in normal use today, which lowers severity but the defect is present in the new code.
|
9 files left out of the code review of
|
|
The code review bot stopped. An unexpected error occurred. Maintainers: see the lucebox_bots run log. |
davide221
left a comment
davide221
left a comment
There was a problem hiding this comment.
Code review
Found 9 issues:
-
bf16 wire length is chosen from env only, but the receive/publish path depends on cfg.size and KERNEL_FLAGS: silent corr · confirmed (
server/src/cluster/fast_reduce.cpp:680-692) -
Lean hybrid exchange never writes the slot-done flag, but submit's slot-reuse guard still waits on it when ZC is off · confirmed (
server/src/cluster/fast_reduce.cpp:1006-1007) -
Hybrid-exchange call counters are never reset, so one rank-local failure permanently desynchronizes the k indices · likely (
server/src/deepseek4/deepseek4_cluster.h:62) -
Async Engram read starts even when the graph has no cluster broadcast node, so the keys are never uploaded · confirmed (
server/src/deepseek4/deepseek4_fused_verify.inc:2326-2332) -
Whole-row-lanes flag is read as a single byte, so it never matches the magic and the feature is dead · confirmed (
server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu:4016-4019) -
HC-boundary spin barrier bakes a stale target under CUDA-graph replay · likely (
server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu:1127) -
moe_combine_masked_kernel indexes
validwith an unchecked (possibly -1) route id · likely (server/deps/llama.cpp/ggml/src/ggml-cuda/moe-fused.cu:409-411) -
Worker death mid-request calls fail_cluster, which _Exit(3)s immediately and discards the error response the code just b · likely (
server/src/cluster/cluster_head_backend.cpp:429-434) -
Boundary kernel spin-barrier can deadlock when its blocks are not all co-resident · likely (
server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu:848-858)
Reviewed a49bc3ef · 13 candidate issues checked, kept at ≥70% confidence · add the code-review label to review again
| static const bool bf16_wire = fr_env_on("LUCE_CLUSTER_FAST_REDUCE_BF16") && | ||
| fr_env_on("LUCE_CLUSTER_FAST_REDUCE_ZC"); | ||
| sge.length = (uint32_t) ((size_t) n * (bf16_wire && !exact_row ? 2u : sizeof(float))); | ||
| sge.lkey = s.mr_data->lkey; | ||
|
|
||
| ibv_send_wr w{}; | ||
| w.wr_id = next; | ||
| w.sg_list = &sge; | ||
| w.num_sge = 1; | ||
| w.opcode = IBV_WR_RDMA_WRITE; | ||
| w.send_flags = 0; | ||
| w.wr.rdma.remote_addr = pe.data_addr + | ||
| ((size_t) slot * s.slot_stride() + (size_t) s.cfg.rank * s.cfg.max_elems) * sizeof(float); |
There was a problem hiding this comment.
bf16 wire length is chosen from env only, but the receive/publish path depends on cfg.size and KERNEL_FLAGS: silent corr · confirmed
The progress thread's bf16_wire is LUCE_CLUSTER_FAST_REDUCE_BF16 && LUCE_CLUSTER_FAST_REDUCE_ZC, with no check of cfg.size or LUCE_CLUSTER_FAST_REDUCE_KERNEL_FLAGS. But FastReduce::submit only takes the zero-copy bf16 path when zc && !kernel_flags && cfg.size == 2; otherwise it publishes the row as f32 (hipMemcpyAsync of n * sizeof(float), or the kernel-flags path with no packing). With the cluster-mode defaults (which install _ZC and _BF16 for any --cluster-size > 1) and --cluster-size 3, every reduction sends only 2*n bytes of an f32 row while the receiving rank copies and adds 4*n bytes as float — silently garbage output. The same mismatch happens at size 2 with LUCE_CLUSTER_FAST_REDUCE_KERNEL_FLAGS=1. The bf16_wire decision must match the submit-side condition (gate it on cfg.size == 2 and !kernel_flags, or export the decision from one place).
Checked (90%): The progress thread's bf16_wire (fast_reduce.cpp:680-682) is only LUCE_CLUSTER_FAST_REDUCE_BF16 && LUCE_CLUSTER_FAST_REDUCE_ZC, while the submit-side zero-copy/bf16 path (fast_reduce.cpp:833) requires zc && !kernel_flags && s.cfg.size == 2; otherwise submit publishes an f32 row via hipMemcpyAsync (line 856, nsizeof(float)). cluster_launch_env (launch_profiles.h:244-247) installs ZC=1 and BF16=1 for any --cluster-size > 1, and cluster_config allows size up to 8, so with size 3 (or KERNEL_FLAGS=1 at size 2) the wire carries 2n bytes while the peer's stream ops read the row as n floats — a silent length/format mismatch on every reduction. Only slight assumption: that a size>2 run is a configuration operators would hit, which the code explicitly supports.
| static const bool bf16_wire = fr_env_on("LUCE_CLUSTER_FAST_REDUCE_BF16") && | |
| fr_env_on("LUCE_CLUSTER_FAST_REDUCE_ZC"); | |
| sge.length = (uint32_t) ((size_t) n * (bf16_wire && !exact_row ? 2u : sizeof(float))); | |
| sge.lkey = s.mr_data->lkey; | |
| ibv_send_wr w{}; | |
| w.wr_id = next; | |
| w.sg_list = &sge; | |
| w.num_sge = 1; | |
| w.opcode = IBV_WR_RDMA_WRITE; | |
| w.send_flags = 0; | |
| w.wr.rdma.remote_addr = pe.data_addr + | |
| ((size_t) slot * s.slot_stride() + (size_t) s.cfg.rank * s.cfg.max_elems) * sizeof(float); |
| if (fast_reduce_lean()) return true; | ||
| return hipStreamWriteValue32(stm, s.flag_dev + s.done_off(slot), (uint32_t) seq, 0) == hipSuccess; |
There was a problem hiding this comment.
Lean hybrid exchange never writes the slot-done flag, but submit's slot-reuse guard still waits on it when ZC is off · confirmed
Sequence numbers are shared by submit and the hybrid exchange (hyb_seq_for draws from the same s.seq), and a ring slot's previous user can be either kind. With LUCE_CLUSTER_FAST_REDUCE_LEAN=1 (a cluster-mode default) hyb_import returns before hipStreamWriteValue32(done_off) (line 1006), so hybrid sequences never publish done. submit's slot-reuse guard (lines 818-826) is skipped only when zc_lean (lean and LUCE_CLUSTER_FAST_REDUCE_ZC and size==2). The documented operator kill switch LUCE_CLUSTER_FAST_REDUCE_ZC=0 leaves lean and LUCE_CLUSTER_HYBRID_EXCHANGE on: submits then take the guard, and once a submit lands on a slot whose previous user (seq - 4096) was a hybrid exchange, hipStreamWaitValue32(done >= seq-slots) can never be satisfied. The stream parks forever and the stall watchdog _Exit(4)s the whole server mid-request. hyb_export's guard (line 935) has the same asymmetry (!fast_reduce_lean() only). The skip conditions must match: only omit the done write when the submit guard is also omitted.
_Checked (90%): With LUCE_CLUSTER_FAST_REDUCE_LEAN=1 and ZC=0 (the documented kill switch), hyb_import returns before the done write (fast_reduce.cpp:1006-1007) and hyb_combine never writes done either, while submit's slot-reuse guard is skipped only for zc_lean (lean && ZC && size==2, line 818), so a submit whose slot's previous user (seq-slots) was a hybrid exchange waits at line 822 on hipStreamWaitValueGte for a done value that is never written — the stream parks and the watchdog Exit(4)s. hyb_export's guard at line 935 checks only fast_reduce_lean(), confirming the skip conditions are asymmetric.
| if (fast_reduce_lean()) return true; | |
| return hipStreamWriteValue32(stm, s.flag_dev + s.done_off(slot), (uint32_t) seq, 0) == hipSuccess; | |
| static const bool zc_lean = fast_reduce_lean() && fr_env_on("LUCE_CLUSTER_FAST_REDUCE_ZC") && | |
| s.cfg.size == 2; | |
| if (zc_lean || std::getenv("LUCE_CLUSTER_FAST_REDUCE_KERNEL_FLAGS")) return true; | |
| return hipStreamWriteValue32(stm, s.flag_dev + s.done_off(slot), (uint32_t) seq, 0) == hipSuccess; |
| // whatever it declines (prefill-sized payloads). | ||
| std::unique_ptr<cluster::FastReduce> fast; | ||
| // Hybrid exchange call counters: the k-th export/combine/import call of each kind belongs to exchange k. | ||
| uint64_t hyb_export0 = 0, hyb_export1 = 0, hyb_combine = 0, hyb_import = 0; |
There was a problem hiding this comment.
Hybrid-exchange call counters are never reset, so one rank-local failure permanently desynchronizes the k indices · likely
hyb_export0/1, hyb_combine and hyb_import are monotonic counters used as the exchange index k passed to FastReduce::hyb_export/combine/import (deepseek4_cluster.cpp:648-664). Nothing in the runtime ever resets them — not at request start, not when node_error is cleared before a new graph, not on error recovery. As long as both ranks execute exactly the same number of callbacks over the process lifetime the pairing holds, but any rank-local failure that aborts a step on one rank only (a fast_reduce timeout, an OOM in verify_batch, a run_split_head failure at deepseek4_dspark_spec.cpp:1372 which does ok = false; break on that rank alone) leaves the two ranks' counters different. Every subsequent hybrid exchange then pairs mismatched slots: the peer is already waiting on a different row, so the next decode step hangs until the 30 s watchdog or silently exchanges wrong buffers and corrupts the model state. Reset the counters at a point both ranks reach by agreement (e.g. when the request ends, alongside the node_error clear at deepseek4_graph.cpp:11173 / deepseek4_dspark_spec.cpp:1368), or make the slot index positional rather than a running count.
Checked (80%): The counters are only ever initialized (deepseek4_cluster.h:62) and incremented in the callbacks (deepseek4_cluster.cpp:650-663); grep finds no reset anywhere. FastReduce's hyb_seq_for/fast_reduce.cpp:919-926 derives the wire sequence number from these per-rank call counts, and rank-local failures exist that clear node_error and fail only one rank's step (deepseek4_graph.cpp:11170-11174, deepseek4_dspark_spec.cpp:1368-1375 with ok=false; break), so a recovered rank then pairs a different seq than its peer's flag waits expect, hanging or corrupting subsequent exchanges. Remaining uncertainty is only whether every such failure is treated as process-fatal in practice, which the clearing code suggests it is not.
| if (pinned_io && engram_rt && engram && token_ids && ds4_engram_async_enabled() && | ||
| ex->engram_keys->buffer) { | ||
| if (!vc.engram_async) vc.engram_async = std::make_unique<Ds4EngramAsync>(); | ||
| engram_async = vc.engram_async->start( | ||
| backend, w, cache.engram_tokens, token_ids, kv_start, q, | ||
| (size_t) engram->n_layers() * (size_t) q * engram->key_floats()); | ||
| } |
There was a problem hiding this comment.
Async Engram read starts even when the graph has no cluster broadcast node, so the keys are never uploaded · confirmed
The async reader is started whenever pinned_io && engram_rt && engram && token_ids holds, where engram_rt only checks cache.cluster_rt && size > 1. But the graph only contains the ds4_cluster_broadcast_node that consumes the reader's pinned buffer (via the before_broadcast hook installed at line 2610) when cluster_hybrid is true, i.e. when the graph was built with a non-null hybrid. In deepseek4_step_layer_range the fused verify is invoked with fused_hybrid_ready ? moe_hybrid : nullptr, so whenever the cluster is attached but the fused hybrid graph is unavailable (ds4_cluster_fused_graph_available false, or the hybrid storage/cold backend not ready), the graph still allocates ex.engram_keys as a plain input while the step takes the async path and skips the in-line ds4_fv_set(ex->engram_keys, ...) upload. The before_broadcast hook never fires, the device keys tensor is never written, and every Engram layer reads uninitialized device memory — silent output corruption on a normal decode step. Gate the async start on the same condition that adds the broadcast node (a non-null hybrid for this step), e.g. add hybrid && to the condition at line 2326.
Checked (92%): The async start at deepseek4_fused_verify.inc:2326-2332 gates only on pinned_io/engram_rt/engram/token_ids, but engram_rt is set whenever cache.cluster_rt->size() > 1 regardless of hybrid; the ds4_cluster_broadcast_node that consumes the reader's buffer via the before_broadcast hook (lines 592, 2610) is added only when cluster_hybrid = hybrid && cluster_rt && size>1, and deepseek4_graph.cpp:11168 passes fused_hybrid_ready ? moe_hybrid : nullptr. When the cluster is attached but the fused hybrid graph is unavailable (fused_hybrid_ready false, e.g. ds4_cluster_fused_graph_available false or cold backend not ready), the graph keeps ex.engram_keys as a plain input, the in-line ds4_fv_set upload is skipped, the hook never fires, and the Engram layers read uninitialized device memory.
| if (!ids && src0->type == GGML_TYPE_Q8_0 && dst->op == GGML_OP_MUL_MAT && | ||
| dst->op_params[1] == GGML_CUDA_WHOLE_ROW_LANES && | ||
| src0->view_src != nullptr && src0->nb[1] == src0->view_src->nb[1]) { | ||
| kb_off = (int) ((src0->view_offs % src0->nb[1]) / sizeof(block_q8_0)); |
There was a problem hiding this comment.
Whole-row-lanes flag is read as a single byte, so it never matches the magic and the feature is dead · confirmed
ggml_backend_cuda_mul_mat_whole_row_lanes stores GGML_CUDA_WHOLE_ROW_LANES (0x574C524B) with ggml_set_op_params_i32(mul_mat, 1, ...), i.e. as a 32-bit value in int32 slot 1 (bytes 4..7 of op_params). But ggml_tensor::op_params is a byte array, and the check here reads dst->op_params[1] — a single byte (part of the precision slot) — and compares it to the 4-byte magic. The comparison is always false, so kb_off is always 0 and the whole-row-lane block mapping never engages. The column-range product then uses the range-only lane order, which is not bit-identical to the whole-row product this PR promises (the slice bit-exactness test would also pass vacuously, since both sides take the same unflagged path). Read the parameter with ggml_get_op_params_i32(dst, 1) instead.
Checked (95%): ggml-cuda.cu:159 sets the magic as a 32-bit value at int32 slot 1 (bytes 4..7 of the byte-array op_params, ggml-impl.h:163-166), but mmvq.cu:4017 reads dst->op_params[1], a single byte that is part of int32 slot 0, so the comparison to 0x574C524B can never be true and kb_off stays 0; the whole-row-lane path used by deepseek4_graph.cpp:4760 is dead.
| if (!ids && src0->type == GGML_TYPE_Q8_0 && dst->op == GGML_OP_MUL_MAT && | |
| dst->op_params[1] == GGML_CUDA_WHOLE_ROW_LANES && | |
| src0->view_src != nullptr && src0->nb[1] == src0->view_src->nb[1]) { | |
| kb_off = (int) ((src0->view_offs % src0->nb[1]) / sizeof(block_q8_0)); | |
| if (!ids && src0->type == GGML_TYPE_Q8_0 && dst->op == GGML_OP_MUL_MAT && | |
| ggml_get_op_params_i32(dst, 1) == GGML_CUDA_WHOLE_ROW_LANES && | |
| src0->view_src != nullptr && src0->nb[1] == src0->view_src->nb[1]) { |
| CUDA_CHECK(cudaMemsetAsync(ctx.ds4_hc_barrier, 0, sizeof(unsigned long long), ctx.stream())); | ||
| ctx.ds4_hc_barrier_blocks = 0; | ||
| } | ||
| ctx.ds4_hc_barrier_blocks += (unsigned long long) n_tokens; |
There was a problem hiding this comment.
HC-boundary spin barrier bakes a stale target under CUDA-graph replay · likely
The cross-block barrier uses a monotonically growing device counter (ctx.ds4_hc_barrier) with barrier_target taken from a host-side running count (ctx.ds4_hc_barrier_blocks += n_tokens) at launch time. That is only correct while every launch executes exactly once in order. The DS4.1 decode graph is run under ggml's CUDA-graph capture in this serving path (the PR even ships capture-related cluster defaults), and during capture the kernel is recorded with the current cumulative count baked in while the device counter does not advance. On the second and later replays the device counter already meets or exceeds every baked target, so while (atomicAdd(barrier, 0ull) < barrier_target) returns immediately and blocks no longer wait for their peers. The barrier exists precisely because out may share memory with R/block_out (the kernel comment says so), so on replays a block can write out over R while another block is still reading it, silently corrupting the HC streams. Fix by making the target self-contained per launch (e.g. a per-launch target slot/phase, or resetting the counter via a captured cudaMemsetAsync to 0 paired with a per-launch base), or by refusing the fusion when the op will be captured in a CUDA graph.
Checked (88%): The fusion runs in the graph-eval path (ggml-cuda.cu:6078) which executes under cudaStreamBeginCapture (ggml-cuda.cu:7090) when CUDA graphs are enabled. The kernel's barrier_target is a host running count baked in at launch (ds4-hc.cu:1127, 1151/1168) while the device counter is only advanced by the recorded kernel; on graph replays the kernel re-executes with the stale constant while the counter has already passed it, so the spin at ds4-hc.cu:1064-1065 returns immediately and the write-to-out-after-all-reads ordering the kernel comment requires (ds4-hc.cu:1024-1028) is not enforced, corrupting out when it aliases R/block_out. Only slight uncertainty is whether some guard disables this fusion under capture, which I did not find.
| const int32_t id = *(const int32_t *)(ids + (size_t)e * sizeof(int32_t) + (size_t)t * ids_nb1); | ||
| const float valid_w = *(const float *)(valid + (size_t)id * valid_nb1 + (size_t)t * valid_nb2); | ||
| const float w = *(const float *)(weights + (size_t)e * sizeof(float) + (size_t)t * weights_nb1) * valid_w; |
There was a problem hiding this comment.
moe_combine_masked_kernel indexes valid with an unchecked (possibly -1) route id · likely
In the cluster expert-parallel path (and the cold-owner-None hybrid path), routes not owned by this rank arrive masked to expert id -1 (see the new foreign_by_global/foreign_routes handling in moe_hybrid_storage.h and the comment in moe_hybrid_ffn_eval.cpp: "routes masked to -1 by the cluster runtime are evaluated elsewhere"). build_moe_hybrid_ffn_graph now defaults to routing those ids straight into ggml_moe_combine_masked (combine_ids = global_ids, mask_in_combine is on by default via moe_masked_combine_enabled). The kernel then computes valid + (size_t)id * valid_nb1 with id = -1, reading four bytes before the valid LUT tensor — an out-of-bounds device read. If the word read is non-zero, w = router_weight * garbage is non-zero and the kernel adds the expert slot's stale (possibly NaN, since masked owner ids leave the slot unwritten — exactly the hazard the sibling laguna_moe_combine_kernel documents) value into the output, corrupting the layer. Whether it fires depends on allocator layout, so it can pass validation runs and corrupt later ones. Guard the lookup: treat id < 0 (or id >= n_expert) as a masked route and continue.
Checked (85%): moe_combine_masked_kernel reads id at moe-fused.cu:409 and indexes valid with (size_t)id * valid_nb1 at :410 with no id<0 guard, so a -1 route reads before the valid tensor. The default path sets combine_ids = global_ids (moe_hybrid_ffn_eval.cpp:1024) and moe_masked_combine_enabled() defaults on (:1081-1087), while the codebase explicitly routes foreign cluster experts as ids masked to -1 (moe_hybrid_storage.h:110-113, moe_hybrid_ffn_eval.cpp:1883-1886, mmvq.cu:1604), so masked -1 ids reach the kernel in cluster mode. Only residual uncertainty is whether some unverified upstream step filters -1 out of global_ids before the graph combine.
| const int32_t id = *(const int32_t *)(ids + (size_t)e * sizeof(int32_t) + (size_t)t * ids_nb1); | |
| const float valid_w = *(const float *)(valid + (size_t)id * valid_nb1 + (size_t)t * valid_nb2); | |
| const float w = *(const float *)(weights + (size_t)e * sizeof(float) + (size_t)t * weights_nb1) * valid_w; | |
| const int32_t id = *(const int32_t *)(ids + (size_t)e * sizeof(int32_t) + (size_t)t * ids_nb1); | |
| if (id < 0) { | |
| continue; | |
| } | |
| const float valid_w = *(const float *)(valid + (size_t)id * valid_nb1 + (size_t)t * valid_nb2); |
| if (!result.ok() && (!control_.all_alive() || (comm_ && comm_->async_error(&err)))) { | ||
| fail_cluster("request " + std::to_string(request_id) + " failed: " + | ||
| std::string(result.error_detail()) + (err.empty() ? "" : " / " + err), | ||
| kAbortCodeRequest, request_id); | ||
| return result; | ||
| } |
There was a problem hiding this comment.
Worker death mid-request calls fail_cluster, which _Exit(3)s immediately and discards the error response the code just b · likely
When a request fails because a worker died, run_request calls fail_cluster(...), which (with the default exit_on_abort_ = true) ends in std::_Exit(3) on the request thread. The result carrying the DecodeFailed error detail is never returned to the HTTP layer, so the client gets a reset connection instead of the error response. The class has request_fatal_exit specifically to give the request thread kFatalExitGraceMs to unwind and write the HTTP response before exiting; this call site bypasses it. Use request_fatal_exit(...) here (or route fail_cluster's exit through the grace path when a request is in flight).
Checked (85%): fail_cluster (cluster_head_backend.cpp:291-295) calls std::Exit(3) directly on the calling thread, and run_request's worker-death path (lines 429-434) calls fail_cluster on the request thread before returning result, so the HTTP error response is discarded. The class provides request_fatal_exit precisely to give kFatalExitGraceMs for the response to be written, and this call site bypasses it (the watchdog path can't save it either, since request_in_flight is cleared at line 420 before the check).
| if (!result.ok() && (!control_.all_alive() || (comm_ && comm_->async_error(&err)))) { | |
| fail_cluster("request " + std::to_string(request_id) + " failed: " + | |
| std::string(result.error_detail()) + (err.empty() ? "" : " / " + err), | |
| kAbortCodeRequest, request_id); | |
| return result; | |
| } |
| if (tid == 0) { | ||
| __threadfence(); | ||
| atomicAdd(barrier, 1ull); | ||
| while (atomicAdd(barrier, 0ull) < barrier_target) { | ||
| #if defined(GGML_USE_HIP) | ||
| __builtin_amdgcn_s_sleep(1); | ||
| #endif | ||
| } | ||
| __threadfence(); | ||
| } | ||
| __syncthreads(); |
There was a problem hiding this comment.
Boundary kernel spin-barrier can deadlock when its blocks are not all co-resident · likely
Both ds4_hc_boundary_kernel and ds4_hc_boundary_split_kernel end with a block-wide spin barrier (atomicAdd/spin until the cumulative target) that every launched block must reach before any may exit. Each block is 1024 threads with up to 48 KB of dynamic shared memory, so at most one block can be resident per CU; the launch therefore deadlocks whenever the GPU cannot host all n_tokens (or n_tokens * NSPLIT) blocks simultaneously — the unlaunched blocks can never arrive and the resident ones spin forever. ggml_cuda_ds4_hc_boundary_supported checks only shapes and mmv_block, not the device's CU count or occupancy; the code asserts n_tokens <= 64, and any RDNA part with fewer than 64 CUs (e.g. the gfx1151 Strix Halo iGPU with ~40 CUs, which this project explicitly supports in mixed-backend setups) hangs as soon as a boundary runs with more tokens than resident workgroups. The split path's n_tokens * 6 <= 48 gate shows residency was considered there, but the non-split path has no equivalent guard. Fix: compute the maximum co-resident block count (CU count × blocks-per-CU given 1024 threads and the shmem size) in ggml_cuda_ds4_hc_boundary_supported and fall back to the unfused launches when the launch would not fully fit.
Checked (82%): Confirmed in the code: ds4-hc.cu:848-858/1062-1071 has every launched block spin on a cumulative atomicAdd barrier before exiting, the non-split launch at ds4-hc.cu:1164 issues n_tokens (asserted only <= 64, line 1118) 1024-thread blocks with up to 48KB dynamic shmem (so at most one block per CU/SM), and ggml_cuda_ds4_hc_boundary_supported (1099-1106) checks only shapes/mmv_block/shmem, never occupancy — so on any GPU with fewer co-resident blocks than n_tokens the resident blocks spin forever. The split path's own comment (1130-1133, gate n_tokens*6 <= 48) shows residency was a known requirement that the non-split path omits; only uncertainty is whether decode batches large enough to exceed CU count occur on a given device, hence not 95.
davide221
left a comment
davide221
left a comment
There was a problem hiding this comment.
Code review
Found 3 issues:
-
sink-col masked softmax relies on a clang-only fp-contract pragma that nvcc ignores · confirmed (
server/deps/llama.cpp/ggml/src/ggml-cuda/softmax.cu:61-64) -
--cluster-replicate-hot and --cluster-selftest are documented but rejected by the flag parser · confirmed (
server/src/server/server_main.cpp:657-659) -
Head-split output_b view applies the group offset twice when the loader already sliced the tensor · confirmed (
server/src/deepseek4/deepseek4_graph.cpp:2342-2346)
Reviewed a49bc3ef · 5 candidate issues checked, kept at ≥70% confidence · add the code-review label to review again
| static __device__ __forceinline__ float soft_max_sink_col_value(float x, float scale, float m) { | ||
| #pragma clang fp contract(off) | ||
| return x*scale + m; | ||
| } |
There was a problem hiding this comment.
sink-col masked softmax relies on a clang-only fp-contract pragma that nvcc ignores · confirmed
soft_max_sink_col_value must round the scale multiply and the mask add separately (the op's contract is bit-exactness with a scale launch followed by an add launch, and run_sink_mask checks bit equality on every GPU backend, including CUDA). It only uses #pragma clang fp contract(off). This PR itself notes in moe_combine_masked_kernel (moe-fused.cu) that nvcc ignores the clang pragma and builds with fast math, so on CUDA builds x*scale + m will be contracted to an FMA and the result differs in the last bits from the separate mul+add. Use the same belt-and-braces form as moe_combine_masked_kernel: keep the pragma for HIP and force the multiply with __fmul_rn.
Checked (90%): The new soft_max_sink_col_value at softmax.cu:61-64 guards x*scale + m only with #pragma clang fp contract(off), but CUDA builds with -use_fast_math (CMakeLists.txt:226) and nvcc ignores the clang pragma — a fact this same PR documents in moe-fused.cu:396-399, where it uses __fmul_rn for exactly this reason. Since run_sink_mask (test_ds4_fused_ops_cuda.cpp:732-779) compares the fused path bit-exactly (memcmp) against separate scale-then-add launches, an FMA contraction on CUDA would produce last-bit mismatches.
| static __device__ __forceinline__ float soft_max_sink_col_value(float x, float scale, float m) { | |
| #pragma clang fp contract(off) | |
| return x*scale + m; | |
| } | |
| static __device__ __forceinline__ float soft_max_sink_col_value(float x, float scale, float m) { | |
| #pragma clang fp contract(off) | |
| return __fmul_rn(x, scale) + m; | |
| } |
| } else { | ||
| std::fprintf(stderr, "[server] unknown or incomplete cluster flag %s\n", flag); | ||
| return 2; |
There was a problem hiding this comment.
--cluster-replicate-hot and --cluster-selftest are documented but rejected by the flag parser · confirmed
ClusterConfig in cluster_config.h documents --cluster-replicate-hot <k> and --cluster-selftest as operator flags, and the runtime consumes both (deepseek4_cluster.cpp builds the placement from cfg.replicate_hot; ClusterHeadBackend::init and run_cluster_worker branch on cfg.selftest). But the --cluster-* parsing block added in server_main.cpp has no branch for either flag, so passing them falls into the catch-all and the server exits with "unknown or incomplete cluster flag". An operator following the header documentation (or trying to replicate the k hottest experts, a knob the placement builders and the Python script both support) cannot start the server. Add the two branches, e.g.:
} else if (std::strcmp(flag, "--cluster-replicate-hot") == 0) {
if (!int_value(bargs.cluster.replicate_hot) || bargs.cluster.replicate_hot < 0) {
std::fprintf(stderr, "[server] --cluster-replicate-hot must be non-negative\n");
return 2;
}
} else if (std::strcmp(flag, "--cluster-selftest") == 0) {
bargs.cluster.selftest = true;
}
Checked (95%): ClusterConfig documents --cluster-replicate-hot (cluster_config.h:71) and --cluster-selftest (cluster_config.h:81-82), and the runtime consumes both (deepseek4_cluster.cpp:266,285; cluster_head_backend.cpp:224; cluster_worker_main.cpp:389), but the --cluster-* parsing block in server_main.cpp:595-660 has branches for neither flag, so both fall into the catch-all at lines 657-659 and the server exits with 'unknown or incomplete cluster flag'. A grep confirms no other code in server_main.cpp sets replicate_hot or selftest.
| ggml_tensor * out_b_w = head_split | ||
| ? ggml_view_2d(ctx, L.attn_output_b, (int64_t) w.n_lora_o * n_groups, | ||
| L.attn_output_b->ne[1], L.attn_output_b->nb[1], | ||
| ggml_row_size(L.attn_output_b->type, (int64_t) group_first * w.n_lora_o)) | ||
| : L.attn_output_b; |
There was a problem hiding this comment.
Head-split output_b view applies the group offset twice when the loader already sliced the tensor · confirmed
When LUCE_CLUSTER_SLICE_WEIGHTS makes the loader store only this rank's output groups, L.attn_output_b is loaded with ne[0] == n_lora_o * n_groups (deepseek4_loader.cpp slices .attn_output_b.weight columns to [slice_group_first*n_lora_o, +slice_groups*n_lora_o)). The first stage handles this correctly: stored_groups detects the sliced attn_output_a and skips the group_first view. But out_b_w unconditionally applies ggml_row_size(type, group_first * n_lora_o) as the view offset whenever head_split is true. On rank 1 of 2 that offset equals the entire sliced tensor's width, so the view starts at (or past) the end of the buffer and the output projection reads out of bounds / garbage on every attention layer of the prefill and uncached-decode paths. The offset must be zero when the stored tensor is already the rank's slice (mirror the stored_groups logic, e.g. use L.attn_output_b->ne[0] == (int64_t) w.n_lora_o * n_groups ? 0 : group_first * w.n_lora_o for the offset), matching how out_a_3d is handled.
Checked (90%): The loader slices attn_output_b columns under LUCE_CLUSTER_SLICE_WEIGHTS (deepseek4_loader.cpp:2320-2325, slice_group_first = rankslice_groups at 2288), so on rank 1 ne[0] = n_groupsn_lora_o, yet build_mla_output_projection applies the group_first*n_lora_o offset unconditionally when head_split (deepseek4_graph.cpp:2342-2345) while only attn_output_a gets the stored_groups already-sliced handling (2330-2338). With head split active whenever w.cluster_heads_local > 0 (graph 11804-11819), rank 1's view offset equals the full sliced tensor width, so the output projection reads at/past the buffer end on every attention layer.
…release - The bootstrap exchange waits at most timeout_ms for each accept, read and write (cluster_control's deadline helpers), closes every socket it opened on failure and refuses a duplicate rank: a worker that died during its load fails rank 0's set-up instead of parking it. - The unpublished-reduction watchdog starts its clock when the host hands the reduction out, not when the progress thread starts waiting, so an idle server no longer trips it between requests. - A failed ibv_post_send or a failed completion now releases every reduction still in flight and exits, like a stall, instead of stopping the progress thread with the GPU parked on a flag that never comes. Both use one release_and_exit, which says whether the GPU drained or was only released. - The zero-copy path allocates no device scratch (80 MiB at the cluster sizes); the per-device publish counters are freed at shutdown; the submit trace and the unused system-scope flag kernel are gone; the kernel-flag and import switches read "0" as off like the others; comments match the code. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…s fail the load - Every rank keeps the fast reduce only if all ranks brought it up: one rank on RCCL and the other on the raw-RDMA path desynced the collectives. - After the comm is up the ranks compare drafters (none, whole, split). Rank 0 decides whether a request drafts, so without a drafter there every rank decodes AR; otherwise every rank must draft the way rank 0 does, and a mismatch fails the load with a message naming both sides. - A cluster slice the quant blocks cannot follow fails the load. Keeping that one tensor whole left it mismatched with its sliced partner. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…asured slower - ds4_cluster_checksum and ds4_cluster_env_prefill_single_token had no callers. - LUCE_CLUSTER_PREFILL_ATTN_SPLIT forced the prefill head split on unsliced weights and measured slower at 2K and 30K; the split now follows the sliced weights alone, as it did by default. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The fused boundary kernel synchronized its blocks on a running device count against a target the host computed at each launch. A CUDA graph replays the launch with the target it captured while the count keeps growing, so from the second replay on the barrier let every block through. The barrier now keeps its whole state on the device: the last block to arrive resets the count and advances a generation the others wait on. HIP builds, which do not capture graphs, run the same barrier. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
davide221
left a comment
davide221
left a comment
There was a problem hiding this comment.
Code review
Found 4 issues:
-
RDMA write failure kills the progress thread and parks the GPU forever · confirmed (
server/src/cluster/fast_reduce.cpp:735-742) -
Head-split wo_b view applies the global group offset to an already-sliced weight · confirmed (
server/src/deepseek4/deepseek4_graph.cpp:2342-2346) -
HC boundary barrier breaks on CUDA graph replay (target baked at capture, device counter keeps growing) · confirmed (
server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu:850-856) -
ggml_soft_max_ext_sink_col_mask allows a mask the CPU backend refuses (GGML_ASSERT abort) · likely (
server/deps/llama.cpp/ggml/src/ggml.c:4429-4446)
Reviewed a49bc3ef · 7 candidate issues checked, kept at ≥70% confidence · add the code-review label to review again
| if (wc[(size_t) i].status != IBV_WC_SUCCESS) { | ||
| std::fprintf(stderr, | ||
| "[fast-reduce] write failed: %s (seq %llu)\n", | ||
| ibv_wc_status_str(wc[(size_t) i].status), | ||
| (unsigned long long) wc[(size_t) i].wr_id); | ||
| s.stop.store(true); | ||
| return; | ||
| } |
There was a problem hiding this comment.
RDMA write failure kills the progress thread and parks the GPU forever · confirmed
In the default configuration (LUCE_CLUSTER_FAST_REDUCE_KERNEL_FLAGS unset), the GPU waits for peer data with hipStreamWaitValue32, which has no timeout. When a send fails — either ibv_post_send returning an error at line 722 or a non-success completion here — the progress thread just sets stop and returns. It does not set failed (so submit() keeps returning true and enqueueing more stream waits), does not release any flags, and does not exit. The stall watchdog that handles the peer-gone case lives inside this same progress thread, so once it returns nothing ever fires: the next reduction's hipStreamWaitValue32 blocks the GPU queue indefinitely, which is exactly the "parked queue costs a GPU reset" outcome the stall path at lines 647-668 was built to avoid. The comment here ("the waiting kernels time out into a counted, visible failure") only describes the kernel-flag form, which is not the default. Handle it like the stall path: set failed, release the visible/hybrid flags for every queued reduction so the GPU drains, and _Exit (or otherwise fail the run loudly) instead of silently returning.
_Checked (90%): Confirmed in the diff: at fast_reduce.cpp:722-725 and 735-742 a send failure or bad completion only stores stop and returns from the progress thread, never setting failed, never calling release_one, and never exiting. submit() (line 764) checks only failed, so it keeps returning true and enqueueing hipStreamWaitValue32 waits (lines 846, 882) that no one will ever satisfy in the default stream-flag mode, and the stall watchdog that would have released flags and Exit'd lives inside the same returned thread (lines 647-668) — the GPU queue parks indefinitely, exactly the outcome the stall path avoids.
| ggml_tensor * out_b_w = head_split | ||
| ? ggml_view_2d(ctx, L.attn_output_b, (int64_t) w.n_lora_o * n_groups, | ||
| L.attn_output_b->ne[1], L.attn_output_b->nb[1], | ||
| ggml_row_size(L.attn_output_b->type, (int64_t) group_first * w.n_lora_o)) | ||
| : L.attn_output_b; |
There was a problem hiding this comment.
Head-split wo_b view applies the global group offset to an already-sliced weight · confirmed
When the cluster rank loads only its own output groups (w.cluster_heads_local > 0, the TargetLoadPlan-sliced case this function explicitly handles for out_a_3d), L.attn_output_b already holds just this rank's n_groups columns (stored_groups == n_groups). The out_b_w view is taken whenever head_split is true, with a byte offset of group_first * n_lora_o rows, where group_first is the global first group of the rank (e.g. half the groups on rank 1). Applied to the sliced tensor that offset lands past the end of the buffer, so the second stage of the output projection reads out-of-bounds/wrong weight rows on every attention layer of rank 1 whenever weights are sliced (the LUCE_CLUSTER_SLICE_WEIGHTS default). Mirror the out_a_3d handling: only view when the tensor still holds all groups.
Checked (95%): The loader slices attn_output_b by columns to only the rank's groups (deepseek4_loader.cpp:2320-2324, ne[0] becomes n_lora_oslice_groups), and the same function's assertion confirms stored_groups == n_groups in this case, yet out_b_w at deepseek4_graph.cpp:2342-2346 is taken whenever head_split is true and adds the global group_firstn_lora_o byte offset — on rank 1 that offset equals the sliced tensor's full width, reading past the buffer every attention layer under LUCE_CLUSTER_SLICE_WEIGHTS.
| ggml_tensor * out_b_w = head_split | |
| ? ggml_view_2d(ctx, L.attn_output_b, (int64_t) w.n_lora_o * n_groups, | |
| L.attn_output_b->ne[1], L.attn_output_b->nb[1], | |
| ggml_row_size(L.attn_output_b->type, (int64_t) group_first * w.n_lora_o)) | |
| : L.attn_output_b; | |
| ggml_tensor * out_b_w = head_split && stored_groups != n_groups | |
| ? ggml_view_2d(ctx, L.attn_output_b, (int64_t) w.n_lora_o * n_groups, | |
| L.attn_output_b->ne[1], L.attn_output_b->nb[1], | |
| ggml_row_size(L.attn_output_b->type, (int64_t) group_first * w.n_lora_o)) | |
| : L.attn_output_b; |
| atomicAdd(barrier, 1ull); | ||
| while (atomicAdd(barrier, 0ull) < barrier_target) { | ||
| #if defined(GGML_USE_HIP) | ||
| __builtin_amdgcn_s_sleep(1); | ||
| #endif | ||
| } | ||
| __threadfence(); |
There was a problem hiding this comment.
HC boundary barrier breaks on CUDA graph replay (target baked at capture, device counter keeps growing) · confirmed
ggml_cuda_ds4_hc_boundary implements the cross-block barrier as a running count: the host increments ctx.ds4_hc_barrier_blocks per launch and passes the cumulative value as barrier_target, while each block does atomicAdd(barrier, 1) and spins until the device counter reaches it (also in the split kernel at lines 1063-1069). This invariant — host count == device count — only holds when every kernel execution is preceded by the host incrementing the counter.
These fused launches happen inside ggml_cuda_graph_evaluate_and_capture (called from the graph evaluation loop in ggml-cuda.cu), so they are recorded into CUDA graphs. On capture the kernels do not execute, but the host counter advances; the recorded kernels carry the capture-time target by value. The first replay is consistent, but from the second replay onward the device counter already exceeds the baked target before the launch starts (previous replays added their blocks without any host increment), so while (atomicAdd(barrier, 0ull) < barrier_target) never blocks. The barrier exists because out may alias R/block_out (per the comment at lines 812-814): without the wait, one block can write out while another token's block still reads those buffers, silently corrupting results on the normal decode path, which replays captured graphs every step.
The counter/target scheme needs to be made replay-safe — e.g. reset the device counter (and host count) at the start of each evaluation/replay, or synchronize without depending on absolute cumulative targets (per-launch sense-reversing epoch, or a zeroed scratch counter with the launch writing its own target into device memory at replay time).
Checked (90%): ggml_cuda_ds4_hc_boundary (ds4-hc.cu:1121-1127) keeps a host-side running count and passes the cumulative value as a by-value kernel argument, while the kernels spin until the device counter reaches it (ds4-hc.cu:851, 1065). The launch is made from the node loop inside ggml_cuda_graph_evaluate_and_capture (ggml-cuda.cu:6078), which per line 5909 only walks nodes when !use_cuda_graph || cuda_graph_update_required — i.e. during capture the host counter advances (kernels only recorded with the capture-time target) and on later replays the host loop is skipped entirely, so the device counter already meets or exceeds the baked target and the cross-block wait never blocks, defeating the alias-protection the barrier exists for (comment at ds4-hc.cu:810-814). Nothing resets ds4_hc_barrier_blocks or the device counter between replays.
| struct ggml_tensor * ggml_soft_max_ext_sink_col_mask( | ||
| struct ggml_context * ctx, | ||
| struct ggml_tensor * a, | ||
| struct ggml_tensor * mask, | ||
| struct ggml_tensor * sinks, | ||
| float scale) { | ||
| struct ggml_tensor * result = ggml_soft_max_ext_sink_col(ctx, a, sinks, scale); | ||
| if (mask) { | ||
| GGML_ASSERT(mask->type == GGML_TYPE_F16 || mask->type == GGML_TYPE_F32); | ||
| GGML_ASSERT(mask->nb[0] == ggml_type_size(mask->type)); | ||
| GGML_ASSERT(mask->ne[0] == a->ne[0]); | ||
| GGML_ASSERT(mask->ne[1] >= a->ne[1]); | ||
| GGML_ASSERT(a->ne[2] % mask->ne[2] == 0); | ||
| GGML_ASSERT(a->ne[3] % mask->ne[3] == 0); | ||
| result->src[1] = mask; | ||
| } | ||
| return result; | ||
| } |
There was a problem hiding this comment.
ggml_soft_max_ext_sink_col_mask allows a mask the CPU backend refuses (GGML_ASSERT abort) · likely
The new public op builds a sink-column softmax with result->src[1] = mask, and the CUDA softmax kernel (softmax.cu, part 7) was extended to handle that combination. The CPU kernel in ggml-cpu/ops.cpp was not: in sink_col mode it has GGML_ASSERT(sk && !src1) (ops.cpp:5369), which rejects exactly the mask+sink_col combination this function creates. GGML_ASSERT aborts the process, and the CPU backend's supports_op still advertises SOFT_MAX unconditionally, so any graph that lands this op on the CPU backend — a CPU-fallback split in the scheduler, or a CPU run of the op outside the HIP/CUDA test path — crashes at compute time instead of being rejected at planning. Either teach the CPU kernel the masked sink-col path (scale, add mask row, append sink column, softmax over ne00+1) or make the CPU backend reject the op when sink_col is set and src[1] is non-NULL so the scheduler places it elsewhere.
Checked (75%): ggml.c:4443 sets result->src[1] = mask on a sink_col SOFT_MAX, and the CPU kernel asserts GGML_ASSERT(sk && !src1) in sink_col mode (ggml-cpu/ops.cpp:5369), while ggml_backend_cpu_device_supports_op has no SOFT_MAX case and so returns true by default (ggml-cpu/ggml-cpu.cpp:499-500) — unlike SSM_CONV/FLASH_ATTN_EXT, which are explicitly guarded. So the op is admitted by the CPU backend and aborts at compute time; the only assumption left is that a graph actually places this op on the CPU backend (scheduler fallback or CPU run), which I could not directly confirm.
Conflicts with #827 (Qwen3.8-Flash-Next split), resolved: - ggml-backend.cpp: a split input copy stays persistent for the parallel scheduler's copy events (main) and for dedicated copies (this branch). - ggml-cuda.cu: the DS4 fusions keep their place before main's qwen4exp fusions, which now run on every device of the graph. - adaptive_spec_width.h: main's shared controller replaces this branch's change, as planned: DS4 seeds it with CostSeed::kShape, and this branch's tests of the old change go with it. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…/up pair The generic gate/GLU fusion forms a one-column gate/up pair with SwiGLU-DS4 (a decode step without a drafter, or the last tokens of a reply), and the RDNA4 Q8_0 dispatch took it to the pair launcher, which had no one-column case: "unreachable q8_0 rdna4 glu width" on a two-box rank. One column is now a case of the same launcher, so a decode step runs the kernel a verify's columns run and the two stay batch-invariant. test_ds4_fused_ops_cuda checks the pair at one column too, bit for bit against the products kept apart. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ough A cluster rank kept prefill bands whole unless LUCE_CLUSTER_PREFILL_PIPELINE was set, because at 512-row bands the pipeline lost. With bands of 1536 rows or more, one part's cold FFN running beside the next part's attention wins: two boxes, a 2K prompt in one band, 649-669 -> 710-727 tok/s (2 runs each). The pipeline is now on whenever every band of the pass is that long; the switch still forces it either way. Every rank sees the same bands, so the decision needs no extra agreement. Each prefill chunk reads all of a rank's experts once, so the rank with the least free memory sets the cluster's chunk for every prompt. Rank 0 now names that rank when it holds the chunk below the pipeline's length (a 2K prompt then runs as two chunks: 558-598 tok/s): a smaller expert budget there leaves room for longer chunks. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ernels too The ds41-gorgon profile sends only the cold owner's experts through the grouped MUL_MAT_ID prefill graph (LUCE_DS4_GROUPED_HYBRID_PREFILL=1); the R9700's hot experts took the per-expert path. On a cluster rank the R9700 does about 70% of the routed work in prefill, so cluster mode now sets 2 (both owners). Two boxes, 2K prompt in one band: the hot owner's prefill time 776 -> 582-599 ms, prefill 688 -> 711-731 tok/s; decode unchanged (code 85.5-88.1, 2K 50.8-51.5). The 2K text differs at near-ties, as any change of the prefill kernels does. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The F16 WMMA routed-expert GEMM for large mul_mat_id batches (ROCmFP2/FP3 MIX and plain ROCmFP2) ran on RDNA3.5 only; an R9700 (gfx1201) took MMQ. gfx12 has the same 16x16x16 F16 WMMA with another fragment layout (a lane holds 8 of the K values and C rows 8*hi + e instead of all 16 and 2*e + hi), so the kernel now builds for both and picks the layout per arch. Two boxes, DS4.1, 2K prompt: the R9700's expert time in prefill 600 -> 463 ms, prefill 715-733 -> 744-756 tok/s; decode unchanged. The kernel decodes the same F16 weights, so results stay within F16 rounding of MMQ: test_rocmfp_mix_gateup_glu passes on gfx1201 through path=mix_wmma, and test_deepseek4_mmid_grouped_cuda on both GPUs. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
LUCE_CLUSTER_PREFILL_Q8=1 sends prefill-sized partials as blocks of 32 int8 values with one f32 scale, all-gathered (1.125 bytes a value against the bf16 all-reduce's 2 per direction), and every rank sums the ranks' dequantized values in rank order, so the result is the same on all ranks. On a link held to about 1.8 GB/s by the NIC's x2 slot, one exchange of a 1K-row band took 5.95 ms in bf16 and 4.36 ms in 8 bits (waits for the other rank included). Two boxes, 2K prompt: 744-756 -> 789-807 tok/s. Off by default: it rounds every prefill partial to 8 bits, so outputs change at near-ties; it stays opt-in until a quality check against bf16. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The quality check the 8-bit exchange waited for, two boxes, harness suites, greedy: Math500 9/10 (bf16: 8/10), GSM8K 9/10 (bf16: 10/10), one answer flipping each way, 18/20 both. Small, but no sign that rounding the prefill partials to 8-bit blocks costs answers, and 2K prefill goes 744-756 -> 789-807 tok/s. LUCE_CLUSTER_PREFILL_Q8=0 returns to the bf16 all-reduce. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The HC pre-mix of a prefill band is RMS_NORM of the [n_hc * n_embd, T] residual read only by a product against 24 F16 rows. It ran as a norm, an F16 copy of the normed rows and a GEMM 24 rows tall: three passes over a band's ~83 MB residual for 1 GFLOP. ggml_cuda_ds4_rms_norm_mix computes the 24 dot products and the sum of squares in one pass, four tokens a block, F32 throughout. Only for bands of 64 tokens or more: decode and verify keep their batch-invariant products. Two boxes, DS4.1, 2K prompt: the HC pre-mix 235 -> 106 ms per prefill. test_ds4_fused_ops_cuda checks it against a double-precision reference on gfx1201 and gfx1151 (the unfused path rounds its input to 16 bits on some GPUs, so it is no reference for an F32 kernel). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
A layer-major prefill band holds every Engram layer's rows, and rank 0 broadcast all of them at each Engram layer the band reached: twice per band on V4.1 (layers 1 and 14), about 100 MB a time for a 1K-row band. The band now remembers that its rows reached every rank. The prefill timing line counts the broadcast (engram_read) and the pipelined joins (hc_post). Two boxes, 2K prompt, with the one-pass HC mix: prefill 788-809 -> 877-902 tok/s; decode unchanged (code 86.8-88.9, 2K 53.8-53.9). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Main's shared thinking budget (#816) replaces the DSpark budget hook this branch carried, so the cluster paths now drive the one ThinkingBudget: - AR decode: rank 0 applies the budget and sends each token with its flags (the budget flag is forced_close()); workers adopt the token and never run a controller of their own. - DSpark seed: rank 0 applies the budget before the seed exchange. - DSpark blocks: rank 0's accept and bonus still cross before the emit loop. A cluster rank that drafts has a settled budget and no local cancel check (rank 0 stops through an empty draft block); the decoder refuses anything else. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Rank 0 picks each request's decode mode for every rank and sends it in the request message: sampled and budgeted requests decode AR. The workers followed it, but rank 0 handed the original request to its own backend, which speculates on both now that main runs sampling and the thinking budget inside DSpark. The ranks then sat in two different decode loops. Rank 0 now decodes the same forced-AR request the workers rebuild, through one rule both sides share (request_decodes_ar). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…ents Both were off by default and measured no better: LUCE_DS4_STAGE_BATCH (the step's pinned uploads as one copy kernel) gained nothing, and LUCE_DS4_VERIFY_PREBUILD (building the next shapes while a verify runs) filled the slot cache and made the two-box verify 42 -> 82 ms. The pinned stage loses its batch mode and ds4_fused_graph_compute is main's again. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Each rank now computes half of the shared expert's intermediate axis and sends its partial in the routed all-reduce, the configuration the measured numbers and the documented command use. Replicate computed the whole expert on both ranks: two boxes, code 89.7-91.9 -> 93.2-97.4 tok/s, 2K prose 52.5 -> 56.3-56.5, at 44.2 -> 41.7 ms a code step. The rank0 mode, accepted by the parser but refused at load, is gone. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
DSpark's width controller started every request from one set of step costs, whatever the context. A step costs more as the context grows (on two boxes q5 is 43 ms on a short prompt and 47-53 ms after a 2K one), so a request that starts from another context's costs picks its widths from the wrong prices until its own steps outweigh them. The costs now live in the shared SpecWidthCostMemory, one set per context range (under 2K, 2-8K, 8-32K, longer), as Qwen's MTP keeps them. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… path - The progress thread sized each RDMA write as bf16 whenever _BF16 and _ZC were set, while submit packs bf16 only on the zero-copy path (two ranks, stream flags). With _KERNEL_FLAGS=1 or more than two ranks it sent half of an f32 row. fr_bf16_wire() is now the one rule: the progress thread, submit and the hybrid combine all follow it. - _LEAN drops the slot-reuse wait and the done write on every stream-flag path. Submit kept its wait unless zero-copy was also on, so with _ZC=0 a submit landing on a slot a hybrid exchange last used waited for a done flag nothing writes. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
- A worker lost mid-request, or a failed request/RequestEnd/report/backend op broadcast, called fail_cluster, which exits on the spot: the client got a reset connection instead of the error. These request-thread paths now use request_fatal_exit, which exits after a grace period, as the watchdog already did. fail_cluster stays for start-up failures. - The head hook no longer samples when the caller passes no token: every caller samples in its decode loop with the request's sampler and RNG, and the unused fallback seeded its RNG once per process. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
- The verify step read the Engram keys on a worker whenever the cluster was attached, but only a graph built for the hybrid cluster path has the broadcast whose callback uploads them. A cluster step without the hybrid graph would have read stale keys; it now reads them in line. - The head-split wo_b column view is taken only from unsliced storage, like wo_a's group view (a sliced rank's head range starts at 0, so the view was a no-op there). - The hybrid exchange counters say what they assume: every rank runs every exchange, and a rank-local step failure ends both ranks. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…lumn - moe_combine_masked skips a route masked to -1 instead of reading the validity row before its tensor. - The split HC boundary (LUCE_DS4_HC_BOUNDARY_SPLIT) runs only when the device holds its whole grid at once (occupancy x CUs): its blocks wait for each other at the output barrier, so a wider grid would never end. The unsplit grid stays at most MMVF_MAX_BATCH_SIZE blocks. - The sink-column softmax rounds the scaled score before the mask add on CUDA too (__fmul_rn; nvcc ignores the clang pragma). - The CPU softmax runs the masked sink-column form it used to assert on. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
It kept a per-device stream and pinned host word in statics that were never freed, and its only caller is an opt-in switch that goes away. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The CONT-into-MMVQ fusion ran MMVQ directly and the rope q8_1 producer assumed MMVQ, each with its own copy of a partial eligibility check. ggml_cuda_mul_mat sends some of those products elsewhere: RDNA4 Q4_K at two or more columns over 2048 or more rows goes to MMQ, MXFP8 to its own matvec, and the qwen4exp BF16 shadow to cuBLAS. The fused graph then ran another kernel than the unfused one (or MMVQ on MXFP8, which it has no kernel for). ggml_cuda_mul_mat now takes its MMVQ eligibility, the RDNA4 Q4_K MMQ preference and the BF16 shadow test from shared helpers, and both fusions decline unless ggml_cuda_mul_mat_takes_mmvq agrees. Also drops the kill switches this PR added for its default paths: LUCE_CUDA_Q4K_NCOL_MMQ, LUCE_DS4_FUSE_HC_BOUNDARY, LUCE_MMID_PREP_MEMO and LUCE_CUDA_MMVQ_Q8_RDNA4_GLU. The tests compare fused against unfused by graph shape, not by these switches. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The skinny-K kernel accumulates in F32, so it changes the bits of the single-column F16 products it takes. It had no vendor guard and changed CUDA results too. It now runs only on RDNA4 under HIP, the GPU it was tuned on; CUDA and other AMD archs keep the generic MMVF kernel. Its LUCE_CUDA_MMVF_SKINNY switch and benchmark-log comment are gone. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…builder The per-verify input timers and the vocab split debug line were research aids. The vocab split is now a builder argument instead of a field kept across Extra::reset(), and the split head logs its rows once. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Pinned step I/O, the async Engram read, the lane batch compressor and selection, the masked MoE combine and the drafter's device context window are on wherever they run, so their switches and off paths go. The shared expert slice view, sliced dense weights and the prefill device join are cluster-only and on in cluster mode. The worker keepalive was never on. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The attribution to the upstream cluster project stays once in deepseek4_cluster.h. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Remove ds4_cluster_allreduce_layer and ds4_cluster_allreduce_fn, which had no callers, and the LUCE_CLUSTER_TRACE_COLL debug print. The plain and exact in-graph broadcasts share one callback body and one node builder. Name the prefill partial threshold. Comments lose benchmark numbers, plan labels, the protocol changelog and a stale note about mixed qtypes; one attribution stays in deepseek4_cluster.h. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Nothing outside the unit tests ever called hash_probe, so the flag only travelled in Welcome and was printed. Remove the flag, its validation, the Welcome field, HashProbeMsg, hash_floats/hash_tokens, the mismatch counters and the head hooks' unused config. Welcome changes on the wire, so the protocol version goes to 3. The gather timeout test now calls the control channel directly. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…tale switches The head backend's introspection accessors, its fault reason and its exit-on-abort toggle had no readers or callers, and common/cluster_view.h declared types nothing used. Cluster ranks always slice their dense weights and keep the device join in prefill now, so LUCE_CLUSTER_SLICE_WEIGHTS, LUCE_DS4_SHEXP_SLICE_VIEW and LUCE_CLUSTER_PREFILL_DEVICE_JOIN go from the cluster defaults and the code. Plan labels and a reference to a launcher script that does not exist are gone from the comments. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
Cluster mode always ran the zero-copy, bf16-wire, lean exchange with in-kernel publishes, one hybrid export and a DMA import. The other paths behind LUCE_CLUSTER_FAST_REDUCE_ZC, _BF16, _LEAN, _KPUB, _KERNEL_FLAGS, LUCE_CLUSTER_HYBRID_ONE_EXPORT and LUCE_CLUSTER_HYBRID_IMPORT_KERNEL were never taken by default, so they go with their switches: the copy-engine add and its scratch, the kernel-form flags and their words, the done flags and the second hybrid export. The fast reduce now refuses any size but two, so other sizes keep RCCL as documented. The MTU and stall settings are read once at init instead of in function-local statics, and the accessors nothing read are gone. The default path's arithmetic is unchanged. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
1 participant
What this does
Serves DeepSeek V4.1 Flash (ROCMFP2S) expert-parallel across two boxes, each an R9700 plus a Ryzen AI Max (Strix Halo) iGPU, linked by a direct ConnectX-5 cable.
fast_reduce, a raw RDMA exchange between the two discrete GPUs. It waits on stream flags, so there is no host round trip, and sends a zero-copy bf16 wire.test_ds4_fused_ops_cuda); the one exception is the drafter head's skinny F16 matvec, which accumulates in F32 and runs on RDNA4 only.Commits, by area:
3d44f159fused DS4.1 decode kernels,228b5bf2RDNA4 Q8_0 gate/up pair,40c4808egrouped ROCmFP2 prefetch loop,4f2c6933their CUDA builds,3a89de1cHC boundary barrier under graph replay,ab496e10single-token gate/up pair,e5f6405bmasked routes, resident split boundary, exact sink column35251926fused routed combine and q8_1 inputs for the hybrid ownersd8f5c547two-box runtime and raw-RDMA exchange,53d673ddworker exit and flags,bcbfc8cabounded handshake and stall clock,b38712518-bit prefill exchange,6a28b6a4rank 0 decodes in the mode it broadcast,3645375fshared expert sharded by default,190676b8one wire format and lean rule,e28cf0f6failed requests answer before exit9f350f02expert-parallel DS4.1 and the cluster defaults,4e6cdcb4start-up agreements,8108b661dead helpers,5306ca57pipelined prefill bands,11a92e02grouped kernels on both owners,5e9448748-bit prefill exchange by default,60364d60Engram rows once per band,9c537619async Engram keys only with the broadcastffc9649cROCmFP2 expert prefill on RDNA4 matrix cores,c8d12671one-pass HC mixaf2e85f4(superseded by #827 in a merge),fcfe8d59width costs per context range,cb630567two experiments removeda49bc3efwhole-row-lane slice46e9fbekeepalive removed,71c78c3MMVQ fusion follows mul_mat's kernel choice,617c75eskinny matvec on RDNA4 only,baac9392144c383e6bf26DS4.1 kill switches, timing and comments,0972e72de9b185effc0c8dead cluster code and the hash probe,1996c81one fast-reduce pathThe branch merges main up to #835. Main's shared thinking budget (#816) replaced the DSpark budget hook this branch carried; the merge commit says how the cluster paths use it.
Review findings on
a49bc3efThe 17 bot findings, against this head:
bcbfc8ca), the two unparsed flags (53d673dd), the HC boundary barrier under graph replay (3a89de1c, both reports), the bf16 wire rule and the lean slot guard (190676b8), the error response on a failed request and the unused sampling fallback (e28cf0f6), the async Engram keys (9c537619), the masked combine's -1 routes, the split boundary's residency, the CUDA sink-column rounding and the CPU masked sink column (e5f6405b).op_paramsisint32_t[]in this ggml, soop_params[1]is the second int32.Second review: cleanup
A structure and leftovers review of
e5f6405bfound no wrong results but too much weight. The branch is now 1,115 lines smaller (+17,378 / -728 against its merge base):--cluster-verify-hashprobe (never called outside tests), unused head-backend accessors and state,common/cluster_view.h, two cluster allreduce helpers with no callers, the worker keepalive (opt-in, never on) withggml_backend_cuda_keepalive, and the verify input timing.fast_reducekeeps its zero-copy, bf16-wire, lean exchange and drops the copy-engine, kernel-flag, done-flag and two-export variants (and 7 switches); DS4.1 drops 9 kill switches whose on path was the default everywhere, and the kernels drop 5. Cluster mode now installs 26 environment defaults instead of 32.ggml_cuda_mul_matwould not pick MMVQ, so fused and unfused take the same kernel; the skinny F16 matvec runs on RDNA4 only, so CUDA keeps upstream's kernel.Not done here: deduplicating the split HC boundary kernel and the five hand-written q8_1 quantizers, and splitting the largest graph functions.
Results
Two boxes, defaults only: the command under "How to run", no tuning variables, five server starts. One box: a Zero 495 with 192 GB,
ds41-gorgon, this branch's build and main's, same harness. One code prompt and one 2K-token prose prompt, 256 tokens, greedy, three requests each after a warm-up request.The shared expert's mode, back to back on the same build:
8270cf3174ec). Sharding changes the rounding of the shared expert's sum, so greedy text can differ from the replicated mode and from a single box at near-ties.How to run
Build with
-DLUCE_CLUSTER=ON, which needs RCCL and libibverbs. Run the same command on each rank, changing only--cluster-rank:The placement and router-bias files are operator inputs, not part of this PR.
scripts/cluster/build_expert_placement.pybuilds a uniform or balanced cluster placement from a routing hotness CSV.Cluster-mode defaults
With
--cluster-sizeabove 1, the server installs the measured configuration as environment defaults (cluster_launch_envinlaunch_profiles.h). They go in before a launch profile's defaults, so they win over the profile but never over a variable that is already set. Set a variable to0to turn its feature off; the log names the variables the operator set.LUCE_CLUSTER_FAST_REDUCE,_STALL_S=15,LUCE_CLUSTER_HYBRID_EXCHANGELUCE_CLUSTER_PREFILL_Q8,LUCE_CLUSTER_PREFILL_BF16(the wire when Q8 is off)LUCE_CLUSTER_ATTENTION_PARALLEL,LUCE_CLUSTER_VOCAB_SPLIT,LUCE_CLUSTER_DRAFT_VOCAB_SPLIT,LUCE_CLUSTER_DRAFT_SPLIT,LUCE_CLUSTER_WORKER_SKIP_DRAFTLUCE_DS4_GPU_ARGMAX_VERIFY,LUCE_DS4_LANE_BATCH,LUCE_DS4_DRAFT_CONTEXT_KV_CACHE,LUCE_DS4_DRAFT_SWAP,LUCE_DS4_SCHED_DEDICATED_COPIES,LUCE_DS4_HC_BOUNDARY_SPLIT,LUCE_DS4_ENGRAM_PRECOMPUTE=2560,LUCE_DS4_TP_CAPTURE_CACHE_SLOTS=64,LUCE_DS4_TP_FUSED_CACHE_SLOTS=64LUCE_MOE_Q8_HANDOFF,LUCE_MMID_GROUPED,LUCE_DS4_GROUPED_HYBRID_PREFILL=2,LUCE_MMID_GROUPED_TYPES=79,LUCE_CUDA_MMVQ_MOE_FP2_PREFETCH,GGML_CUDA_PEER_COPY_KERNELHSA_FORCE_FINE_GRAIN_PCIE--cluster-shared-expertnow defaults toshard;replicateremains for A/B runs.LUCE_DS4_DRAFT_SWAPstays opt-in on a single box, so that prefix-cache hits match a full prefill. In cluster mode it is on, because cluster ranks run without the prefix cache.NCCL_PROTO=Simple, next to the existing IB hardening defaults.LUCE_CLUSTER_TRACE.Known limits
LUCE_CLUSTER_FAST_REDUCE_STALL_S(15 s in cluster mode); restart both.fast_reduceis a two-rank exchange; more ranks fall back to RCCL.Validation
On the rank-0 box, built from this branch with
-DLUCE_CLUSTER=ON(gfx1151 + gfx1201):CPU tests:
test_cluster_unit26/26,test_thinking_budget,test_server_unit657 passed, 2 skipped,.github/ci/check_commit_messages.py(30 commits).GPU tests, on both GPUs:
test_ds4_fused_ops_cuda(0 failures, also withLUCE_DS4_HC_BOUNDARY_SPLIT=1),test_deepseek4_unit. With the split boundary on, the R9700 takes the split kernel for every launch; the Strix Halo iGPU (40 CUs) falls back to the unsplit kernel for the grids it cannot hold at once.Two boxes, defaults only: the Results tables, on two builds (before and after the review fixes; the same text, decode within 2%), one with the bench run twice on one server. A sampled request, a thinking-budget request (closed at its 48-token budget) and a greedy request after them all complete on both ranks; before
6a28b6a4, rank 0 would have decoded the first two speculatively while the worker decoded AR.After the cleanup (
1996c81): every unit and GPU test above passes again on both GPUs. Two boxes: the code and 2K answers are the same text as before (8270cf3174ec,020f5e2ba61b); code 93.5, 97.7, 102.9 tok/s on the first three requests and 103.4-103.8 after, 2K prose 55.3-58.4; sampled, budgeted and greedy requests complete. One box: the same text as before the cleanup on both prompts, at the same speed.Before merge
ds41-gorgonagainst main (55a2dd0), on a Zero 495 with 192 GB: same CMake config, model, drafter, placement and bias, one warm-up request, 256 tokens, greedy.--ds4-prefill exactand no drafter, both prompts give main's text (code6e212c26fc09, 2K prompt6eb664db61c2).ffc9649c, within F16 rounding of MMQ) and the HC pre-mix of bands of 64+ tokens in one F32 pass (c8d12671).LUCE_CLUSTER=ON.test_cluster_unitruns in the CPU job andtest_ggml_cluster_allreducein the AMD GPU job. ALUCE_CLUSTER=ONbuild there also needs the libibverbs development package on the AMD runner; until then the RCCL and RDMA paths are covered by two-box runs.🤖 Generated with Claude Code