NVLink fix for PR625 - #637
Draft
jkyamog wants to merge 14 commits into
Draft
Conversation
Wire the DSpark drafter heads (low-rank Markov bigram correction + confidence head) into the qwen35 spec-decode loop, so Qwen3.8-27B DSpark drafters (e.g. RadixArk/Qwen3.8-27B-DSpark) run with full head support: - spec loop: markov-corrected greedy chain (fused single-graph variant with non-fused fallback) replaces plain argmax projection when the drafter ships DSpark heads; DDTree candidate top-k gets the markov bias too. Env-gated: DFLASH_QWEN35_DSPARK, DFLASH_QWEN35_FUSED_DSPARK, DFLASH_QWEN35_DSPARK_TREE (all default on). - target capture layers now follow the drafter GGUF's dflash.target_layer_ids instead of the evenly-spaced derivation; the Qwen3.8 drafter is trained on layers 4/16/28/40/52, not 1/16/31/46/61. - draft loader: dflash.mask_token_id from the drafter GGUF wins over the family default (Qwen3.8 drafter uses 248077, default was 248070), and optional YaRN rope scaling keys are parsed into DraftWeights. - draft graph: rope calls honor the drafter's YaRN config (previously hardcoded plain NEOX rope). - Qwen35DFlashTarget exposes lm_head for the fused head path. - convert_dflash_to_gguf.py: handle single-file DSpark releases (markov/ confidence heads inline in model.safetensors), transformers>=5 nested rope_parameters and dflash_config.mask_token_id, and emit YaRN scaling metadata. The confidence-gate adaptive block length is not wired yet (q_len sizes the per-request step buffers); the chain runs with the gate off.
Verify/accept now run over v_len (the drafted chain's actual length) instead of the buffer-sizing q_len, so the DSpark confidence gate's adaptive block truncation is structurally supported. The gate itself stays off by default (DFLASH_QWEN35_DSPARK_CONFIDENCE_THRESHOLD=0): with the RadixArk Qwen3.8 drafter, any threshold in 0.1-0.5 truncates to the same short chain regardless of value, so the confidence scores coming out of the shared head path look mis-scaled and need a separate investigation before the gate can help. threshold=0 is bench-verified regression-free.
- ggml_ssm_conv_step: one kernel for the causal-conv decode/verify step (history window + silu(conv) + in-place history write-back + optional rollback window copy) replacing transpose/concat/ssm_conv/silu/cpy. - ggml_gated_delta_net_set_raw_gates: the GDN kernel applies sigmoid(beta) and softplus(alpha + dt_bias) * A itself. - ADD + RMS_NORM + MUL fusion (residual add materialized alongside the normalized output) in the CUDA/HIP graph evaluator. - legacy pool MAX_BUFFERS 256 -> 1024: LUCE_Q8_MEMO holds ~300 pooled buffers per evaluation; a full pool freed in-flight buffers with cudaFree and produced illegal memory accesses on long prefills.
Rename the RDNA small-tile macro to GGML_CUDA_MMQ_SMALL_TILE and apply it to IQ4_XS/Q4_K/Q5_K/Q6_K/Q8_0 in addition to the ROCmFPX formats. At spec-decode verify widths (N<=16) the 128-row tile leaves a 5120-row projection with only 40 blocks on a 64-CU gfx1201; 64x64/4-warp tiles measured +12-23% on those shapes (verify step 43.8 -> 39.7 ms on Qwen3.8-27B) at ~8% prefill cost.
- loader places attn_gate|attn_qkv and ssm_beta|ssm_alpha back to back and exposes zero-copy stacked aliases (L.wqkv_z, L.ssm_ba): one GEMV each instead of two (DFLASH_QWEN35_NO_STACK=1 disables). - FFN uses ggml_swiglu_split so the backend fuses gate/up/GLU into one vector kernel at decode. - DeltaNet block: single l2_norm over the q|k slab, ggml_ssm_conv_step, raw-gate gated_delta_net (in place, no state copy), no q/k head repeat (the kernel broadcasts). DFLASH_QWEN35_NO_FUSED_KERNELS=1 keeps the op-by-op graph for A/B. - DFLASH_KV_ROTATE=0 skips the FWHT K/Q rotation (precision-neutral with q8_0/f16 caches, two fewer launches per attention layer). Qwen3.8-27B IQ4_XS on R9700: plain decode 30.4 -> 33.8 tok/s with identical greedy output.
- Qwen35AdaptiveSpecPolicy: EMA of accepted draft tokens per step; below 0.8*(spec_step_ratio-1) the loop runs a burst of plain-decode steps (seed-only verify, no drafter/heads/snapshot/rollback, features still captured) and probes again afterwards. Env DFLASH_QWEN35_SPEC_STEP_RATIO (default 1.7, 0 disables) and DFLASH_QWEN35_AR_BURST (default 40). Low-acceptance prose 28.1 -> 32.4 tok/s, code/mixed unchanged. - Confidence gate now uses the fused Markov graph and truncates on the host; DFLASH_QWEN35_DSPARK_CONF_DEBUG=1 prints per-position scores. - spec-profile hooks for the chain path (project/snapshot/verify/ rollback/feature).
launch_fattn was told the vec kernel consumes D keys per step; it walks nthreads (128) per step, so a 256-key window at head_dim 256 ran as one block per head. Passing nthreads lets it use two blocks per head plus the combine pass: Qwen3.8-27B plain decode 34.3 -> 34.6 tok/s on R9700, identical output.
The first spec step after a plain-decode burst updates the acceptance EMA with alpha 0.5 so a stream that became predictable leaves plain decode immediately; step ratio and start value keep the measured best balance (45.7 / 31.8 / 40.4 tok/s code / prose / mixed).
The break-even acceptance now follows live EMAs of the spec-step and plain-step wall times (default 1.9 until both are measured), so it is right for any drafter block size (width-8 DSpark and width-16 DFlash measure ~1.8 on gfx1201).
DFlash 2 (z-lab/inco, e.g. z-lab/Qwen3.8-27B-DFlash2) is the DFlash backbone plus a grouped dynamic causal conv around attention and MLP in every layer and a candidate selector head (top-k lm_head candidates per block position, one path scored by a low-rank bigram form). - converter: maps attention_conv/mlp_conv (base kernels F32, kernel projections) and candidate_selector tensors, emits dflash2.* metadata, reads block_size from dflash_config, emits SWA pattern for drafters with causal sliding layers. - loader: DraftConvWeights per layer, DraftSelectorWeights, shape checks. - draft graph: conv prepare/finish (two taps over the block, per-element base + per-group dynamic coefficient) in both the stateless and the cached-KV builders. - selector chain: top-k via the target's GPU top-k (kMaxK 8 -> 16), one cached graph for hproj + codebook row gathers, host path search. - spec loop uses the selector before the DSpark/argmax paths. Qwen3.8-27B IQ4_XS on R9700, q8_0 drafter, greedy: 109.9 code / 50.7 prose / 111.8 mixed tok/s (DSpark drafter: 45.6 / 32.4 / 38.6); avg 5.9-6.0 accepted tokens per 8-token block on code, ~2.7 on prose.
With the 64-row/4-warp tile the mmq_x=32 instantiation runs at 180 GB/s on gfx1201 (17408x5120 IQ4_XS) against 443 GB/s at mmq_x=16 and 315 at 48, so N=17..32 batches (DDTree budgets, prefill remainders) took 2.4x longer than N=16 or N=40. Choose the next tile instead.
Resolves the conflicts with SpecLA (Luce-Org#606) and the packed concurrent prefill work (Luce-Org#595): - SSM_CONV op_params[0]: 1 stays the SpecLA heavy-light conv, the dflash fused conv step now uses 2; the CUDA dispatcher and supports_op handle both. - GDN raw-gate mode no longer uses src[8]/op_params[2] (compact-decode slot ids and the SpecLA marker): the loader builds one f32 [dt_bias | A] tensor per DeltaNet layer (TargetLayer::ssm_gate_ba, own small buffer) and ggml_gated_delta_net_set_raw_gates() attaches it as src[9] with op_params[10] = 1. - gated_delta_net.cu launchers carry both the active-slot arguments and the raw-gate parameters. - build_delta_net_block keeps main's token-axis segment structure; the stacked (z|qkv) and (beta|alpha) projections, the fused conv step, raw gates, single l2_norm over q|k and the head-broadcast shortcut apply on the plain single-sequence chain path only (ragged, compact-decode, SpecLA and chunked paths take main's materialized ops). - MMQ tiles: Q4_K keeps main's 128x64 (LUCEBOX_RDNA_MMQ_Y); IQ4_XS, Q5_K, Q6_K, Q8_0 and the ROCmFPX formats keep the 64x64 small tile. - Spec loop: DSpark Markov top-k stays available inside the non-conditional DDTree branch; SpecLA's conditional-draft path and draft-KV flag are preserved. Verified on lucebox8 (R9700): AR 34.6 tok/s, DFlash2 110/54/115 (code/prose/mixed), DSpark 48/32/40.
The fused conv step (ggml_ssm_conv_step) and the SpecLA heavy-light conv (ggml_ssm_conv_specla) both keep the channel axis on axis 0 for x [C,T,S] and out [C,T,S], while the weight [K,C] and conv_state [K-1,C,S] carry the same channel partition on axis 1. The old handler only handled the equal-axis case, so tensor-split AR decode fell through to the generic handler and hit assert 556 (split state UNKNOWN, op=SSM_CONV). Key the branch on the axis layout (x axis0, weight/conv_state axis1) rather than op_params[0]: the flag's encoding differs between the step and SpecLA variants (and changed across heads), while the axis layout is the stable invariant. Guarded on src[2] existing so the plain/tree variants (2-3 srcs, channel on axis 1) are untouched. Verified on the 3090 pair (tensor split, peer access ON) at 480e60a: AR decode runs clean at 37.3 tok/s (greedy, 24538 in / 32 out), no SPLIT_AXIS_UNKNOWN assert.
Under tensor split the target lm_head is a meta tensor whose data pointer is a placeholder (0x2000000000000000). The fused DSpark paths pass it to graphs computed on the draft's plain CUDA backend, so the mul_mat kernel reads the fake pointer and faults (illegal memory access at GET_ROWS). Port the resolution from the deployed line (fd7b294): when lm_head is a meta tensor, read the mirrored simple tensor of the rank matching the draft GPU so the fused graph does a local read; when the draft GPU is not a target rank, fall back to the host chain. Guard both the fused greedy chain and the DDTree topk path, which had the same leak.
jkyamog
marked this pull request as draft
August 20, 2026 10:11
![cubic-dev-ai[bot]](https://avatars.githubusercontent.com/in/1082092?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
23 issues found across 34 files
Prompt for AI agents (unresolved issues)
Check if these issues are valid — if so, understand the root cause of each and fix them. If appropriate, use sub-agents to investigate and fix each issue separately.
<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py:109">
P3: `Q4_K` never receives the 64×64 guard described by this new group because the later special case overwrites `guard`. Remove it from this set, or change the special case if 64×64 is actually intended, so the grouping remains accurate.</violation>
</file>
<file name="server/src/qwen35/gguf_target_loader.cpp">
<violation number="1" location="server/src/qwen35/gguf_target_loader.cpp:1053">
P2: When a post-gate validation fails, the loader frees only `out.buf` and leaks the newly allocated gate buffer and context. Extend the failure cleanup to release `gate_buf` and `gate_ctx` along with the main target buffer.</violation>
</file>
<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:4385">
P2: When `MUL` uses the same RMS tensor for both operands, this branch assigns `w = rms` and the fused kernel reads the unmaterialized RMS output as its weight. Reject equal MUL operands before enabling this fusion, or explicitly reject `w == rms`.</violation>
</file>
<file name="server/deps/llama.cpp/ggml/src/ggml.c">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml.c:5894">
P3: Update this comment to say `op_params[0] = 2`; `1` selects SpecLA while `2` selects the fused step kernel, so the current explanation can mislead future callers or maintenance changes.</violation>
<violation number="2" location="server/deps/llama.cpp/ggml/src/ggml.c:5894">
P3: This comment states `op_params[0] = 1` marks step mode, but the code that follows sets `ggml_set_op_params_i32(result, 0, 2)` and the CUDA dispatcher routes step mode on `ggml_get_op_params_i32(dst, 0) == 2`, while `1` belongs to the SpecLA variant (also set by `ggml_ssm_conv_specla`). The comment contradicts both the implementation and the dispatch, which will mislead anyone reading or extending this path.</violation>
</file>
<file name="server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp:9336">
P3: The guard correctly rejects every non-standard SSM_CONV mode on CPU, but its comment and abort message name only `ggml_ssm_conv_step`. The op-param check `== 0` also rejects the SpecLA heavy-light mode (`ggml_ssm_conv_specla` sets op_params[0]=1 via ggml_set_op_params_i32(result, 0, 1)), which trips the same assert and aborts with a message only about step mode. Broaden the message so a developer hitting this assert for mode 1 understands the actual condition.</violation>
</file>
<file name="server/scripts/convert_dflash_to_gguf.py">
<violation number="1" location="server/scripts/convert_dflash_to_gguf.py:120">
P1: When the config uses the legacy `rope_scaling` format, YaRN is stored under `type`, so this check skips the scaling metadata and the drafter runs plain RoPE. Accept `type` as a fallback when reading `rope_type`.</violation>
</file>
<file name="server/src/common/dflash2_head.cpp">
<violation number="1" location="server/src/common/dflash2_head.cpp:82">
P1: After `park()` frees `dw_` and `unpark()` reloads it into the same object, this cache key still matches. The graph then dereferences freed selector tensors; invalidate it on reload or key it by selector tensor identity or generation.</violation>
<violation number="2" location="server/src/common/dflash2_head.cpp:82">
P2: The cached selector graph is keyed only on (dw, backend, n_cand, K), so a rebuilt graph is reused for any hidden/rank shape as long as the draft pointer and backend are unchanged. If a DraftWeights object is reloaded into the same address (free + reload at the same allocation) or its selector rank/hidden dims change, the stale graph's inp_hidden ([hdim,n_cand]) and hproj/succ/pred spans no longer match, and ggml_backend_tensor_set will write a buffer sized for the old dims into a graph allocated for new dims, producing wrong or out-of-range h_succ/h_pred reads in the path search.</violation>
<violation number="3" location="server/src/common/dflash2_head.cpp:121">
P1: When a selector codebook vocabulary differs from the target vocabulary, `pred_ids` or `cand_ids` can index past `pred_cb`/`succ_cb` in `ggml_get_rows`. Validate every ID against the codebook row count, and reject incompatible selector/target vocabularies before graph execution.</violation>
</file>
<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh:4909">
P2: For every small-tile format on every RDNA3/4 device, this skips the 32-wide candidate even though the cited measurement covers only IQ4_XS on gfx1201. A 25–32-column Q8_0, Q5_K, Q6_K, or ROCmFPX call therefore selects a 40-wide tile and computes eight unused outputs; restrict this heuristic to the measured type/device or validate each affected format first.</violation>
</file>
<file name="server/src/draft/draft_gguf_loader.cpp">
<violation number="1" location="server/src/draft/draft_gguf_loader.cpp:233">
P1: When a drafter GGUF uses a mask ID different from the target default, this assignment does not affect drafting: the backend still embeds `target->mask_token_id()`. Pass `draft_weights.mask_token_id` into noise construction, including the remote-drafter path, so the metadata override actually takes effect.</violation>
<violation number="2" location="server/src/draft/draft_gguf_loader.cpp:518">
P1: When `dflash.dflash2.conv_group_size` is zero or does not evenly divide `n_embd`, the loader divides or reshapes with invalid dimensions instead of rejecting the GGUF. Check that the group size is positive and divides `n_embd` before computing `groups`.</violation>
<violation number="3" location="server/src/draft/draft_gguf_loader.cpp:520">
P2: A DFlash 2 GGUF with a bad later-layer or MLP convolution shape can pass loading even though the graph uses those tensors in every layer. Validate both base/projection tensors for both convolutions across all layers before enabling the feature.</violation>
<violation number="4" location="server/src/draft/draft_gguf_loader.cpp:520">
P2: A malformed GGUF with a wrong mlp_conv shape (or a conv_group_size that doesn't divide n_embd) loads without error, then builds a corrupt dynamic-conv graph. Mirror the attention shape checks for mlp_conv.base/proj on every layer and verify n_embd % conv_group_size == 0.</violation>
<violation number="5" location="server/src/draft/draft_gguf_loader.cpp:546">
P3: check_shape_2d(pred_cb, R, pred_cb->ne[1]) validates ne1 against pred_cb's own value, so a pred_cb with a nonsense vocabulary dimension loads without error. Drop the self-reference and validate pred_cb's ne1 against a real bound (e.g. the draft vocab_size / target vocab) and require succ_cb to match it.</violation>
</file>
<file name="server/scripts/requant_target_rocmfp.py">
<violation number="1" location="server/scripts/requant_target_rocmfp.py:223">
P1: When the input declares a non-default `general.alignment`, this writer emits inconsistent offsets and produces an unreadable or corrupted GGUF. Preserve the input alignment in `GGUFWriter` before copying metadata.</violation>
</file>
<file name="server/src/qwen35/qwen35_backend.cpp">
<violation number="1" location="server/src/qwen35/qwen35_backend.cpp:2394">
P1: When the draft runs on a different GPU without peer access, the fused DSpark graph uses the target GPU's simple `lm_head` directly on `draft_backend_`. Verify the tensor belongs to the draft backend or copy it to that GPU, and return `nullptr` to use the host fallback otherwise.</violation>
<violation number="2" location="server/src/qwen35/qwen35_backend.cpp:2850">
P2: When confidence truncation or an adaptive plain-decode burst shortens a step, the acceptance denominator still counts the full configured `q_len`. Accumulate the actual offered length (`v_len`, or one for AR steps) per step and use that total for `out_accept_rate` and telemetry.
(Based on your team's feedback about offered-candidate accept-rate denominators.)</violation>
<violation number="3" location="server/src/qwen35/qwen35_backend.cpp:2916">
P1: When the fused DSpark head has no compatible confidence head, enabling `DFLASH_QWEN35_DSPARK_CONFIDENCE_THRESHOLD` truncates every chain to one token instead of disabling the gate. Apply the gate only when confidence scores cover the complete candidate chain, or make the fused helper fail when scores are requested but unavailable.</violation>
</file>
<file name="server/src/internal.h">
<violation number="1" location="server/src/internal.h:160">
P2: When target loading fails after alias or raw-gate allocation, `release_out_buffer()` leaves `stack_ctx`, `gate_ctx`, and `gate_buf` allocated. Extend the failure cleanup to release all three before returning, otherwise repeated invalid or partial loads leak GPU memory and ggml contexts.</violation>
</file>
<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q8_0.cu">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/mmq-instance-q8_0.cu:3">
P2: On RDNA3/4, this forces every Q8_0 MMQ launch to use the small tile, including large prefill batches. Gate it to short verify widths or select the tile at runtime so prefill does not incur the documented throughput loss.</violation>
</file>
<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu:404">
P3: The raw-gate gate/beta math (softplus with the 20.0f clip, gate_A scaling, sigmoid) is duplicated between `gated_delta_net_cuda` and `gated_delta_net_cuda_grouped_cols`. Extract it into a shared `__device__ __forceinline__` helper (or compute it in `launch_gated_delta_net`) so the two kernels cannot diverge on the threshold or formula.</violation>
</file>
Tip: cubic used a learning from your PR history. Let your coding agent read cubic learnings directly with the cubic MCP.
Re-trigger cubic
| if isinstance(rp, dict): | ||
| if rp.get("rope_theta") is not None: | ||
| a["rope_theta"] = float(rp["rope_theta"]) | ||
| if str(rp.get("rope_type", "")).lower() == "yarn": |
Contributor
There was a problem hiding this comment.
P1: When the config uses the legacy rope_scaling format, YaRN is stored under type, so this check skips the scaling metadata and the drafter runs plain RoPE. Accept type as a fallback when reading rope_type.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/scripts/convert_dflash_to_gguf.py, line 120:
<comment>When the config uses the legacy `rope_scaling` format, YaRN is stored under `type`, so this check skips the scaling metadata and the drafter runs plain RoPE. Accept `type` as a fallback when reading `rope_type`.</comment>
<file context>
@@ -110,6 +110,37 @@ def pick(*keys):
+ if isinstance(rp, dict):
+ if rp.get("rope_theta") is not None:
+ a["rope_theta"] = float(rp["rope_theta"])
+ if str(rp.get("rope_type", "")).lower() == "yarn":
+ a["yarn_factor"] = float(rp.get("factor", 0.0))
+ a["yarn_orig_ctx"] = int(rp.get("original_max_position_embeddings", 0))
</file context>
Suggested change
| if str(rp.get("rope_type", "")).lower() == "yarn": | |
| if str(rp.get("rope_type") or rp.get("type") or "").lower() == "yarn": |
| pred_ids[0] = last_tok; | ||
| std::memcpy(pred_ids.data() + 1, cand_ids.data(), sizeof(int32_t) * (size_t)n_cand * K); | ||
| ggml_backend_tensor_set(g.inp_hidden, local_hidden + (size_t)hdim, 0, sizeof(float) * (size_t)hdim * n_cand); | ||
| ggml_backend_tensor_set(g.inp_succ, cand_ids.data(), 0, sizeof(int32_t) * (size_t)n_cand * K); |
Contributor
There was a problem hiding this comment.
P1: When a selector codebook vocabulary differs from the target vocabulary, pred_ids or cand_ids can index past pred_cb/succ_cb in ggml_get_rows. Validate every ID against the codebook row count, and reject incompatible selector/target vocabularies before graph execution.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/common/dflash2_head.cpp, line 121:
<comment>When a selector codebook vocabulary differs from the target vocabulary, `pred_ids` or `cand_ids` can index past `pred_cb`/`succ_cb` in `ggml_get_rows`. Validate every ID against the codebook row count, and reject incompatible selector/target vocabularies before graph execution.</comment>
<file context>
@@ -0,0 +1,156 @@
+ pred_ids[0] = last_tok;
+ std::memcpy(pred_ids.data() + 1, cand_ids.data(), sizeof(int32_t) * (size_t)n_cand * K);
+ ggml_backend_tensor_set(g.inp_hidden, local_hidden + (size_t)hdim, 0, sizeof(float) * (size_t)hdim * n_cand);
+ ggml_backend_tensor_set(g.inp_succ, cand_ids.data(), 0, sizeof(int32_t) * (size_t)n_cand * K);
+ ggml_backend_tensor_set(g.inp_pred, pred_ids.data(), 0, sizeof(int32_t) * (size_t)n_rows_pred);
+ if (ggml_backend_graph_compute(backend, g.gf) != GGML_STATUS_SUCCESS) {
</file context>
| // reused across steps. | ||
| const int n_rows_pred = 1 + n_cand * K; | ||
| SelectorGraph & g = selector_graph(); | ||
| if (!g.ctx || g.dw != &dw || g.backend != backend || g.n_cand != n_cand || g.K != K) { |
Contributor
There was a problem hiding this comment.
P1: After park() frees dw_ and unpark() reloads it into the same object, this cache key still matches. The graph then dereferences freed selector tensors; invalidate it on reload or key it by selector tensor identity or generation.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/common/dflash2_head.cpp, line 82:
<comment>After `park()` frees `dw_` and `unpark()` reloads it into the same object, this cache key still matches. The graph then dereferences freed selector tensors; invalidate it on reload or key it by selector tensor identity or generation.</comment>
<file context>
@@ -0,0 +1,156 @@
+ // reused across steps.
+ const int n_rows_pred = 1 + n_cand * K;
+ SelectorGraph & g = selector_graph();
+ if (!g.ctx || g.dw != &dw || g.backend != backend || g.n_cand != n_cand || g.K != K) {
+ selector_graph_free(g);
+ const size_t arena_size = ggml_tensor_overhead() * 32 + ggml_graph_overhead() + 4096;
</file context>
| out.conv_kernel_size = conv_k; | ||
| out.conv_group_size = (int)read_u32("dflash.dflash2.conv_group_size", 16); | ||
| const DraftLayer & L0 = out.layers[0]; | ||
| const int64_t groups = out.n_embd / out.conv_group_size; |
Contributor
There was a problem hiding this comment.
P1: When dflash.dflash2.conv_group_size is zero or does not evenly divide n_embd, the loader divides or reshapes with invalid dimensions instead of rejecting the GGUF. Check that the group size is positive and divides n_embd before computing groups.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/draft/draft_gguf_loader.cpp, line 518:
<comment>When `dflash.dflash2.conv_group_size` is zero or does not evenly divide `n_embd`, the loader divides or reshapes with invalid dimensions instead of rejecting the GGUF. Check that the group size is positive and divides `n_embd` before computing `groups`.</comment>
<file context>
@@ -451,6 +497,64 @@ bool load_draft_gguf(const std::string & path,
+ out.conv_kernel_size = conv_k;
+ out.conv_group_size = (int)read_u32("dflash.dflash2.conv_group_size", 16);
+ const DraftLayer & L0 = out.layers[0];
+ const int64_t groups = out.n_embd / out.conv_group_size;
+ char shape_err[192];
+ if (!check_shape_3d(L0.attn_conv.base, out.n_embd, conv_k, 2, "attn_conv.base", shape_err, sizeof(shape_err)) ||
</file context>
Suggested change
| const int64_t groups = out.n_embd / out.conv_group_size; | |
| if (out.conv_group_size <= 0 || (out.n_embd % out.conv_group_size) != 0) { | |
| set_last_error("draft GGUF: invalid DFlash 2 conv_group_size"); | |
| ggml_free(meta_ctx); out.ctx = nullptr; gguf_free(gctx); | |
| return false; | |
| } | |
| const int64_t groups = out.n_embd / out.conv_group_size; |
| // embedding silently destroys acceptance. | ||
| { | ||
| const uint32_t mask_meta = read_u32("dflash.mask_token_id", 0); | ||
| if (mask_meta != 0) out.mask_token_id = (int32_t)mask_meta; |
Contributor
There was a problem hiding this comment.
P1: When a drafter GGUF uses a mask ID different from the target default, this assignment does not affect drafting: the backend still embeds target->mask_token_id(). Pass draft_weights.mask_token_id into noise construction, including the remote-drafter path, so the metadata override actually takes effect.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/draft/draft_gguf_loader.cpp, line 233:
<comment>When a drafter GGUF uses a mask ID different from the target default, this assignment does not affect drafting: the backend still embeds `target->mask_token_id()`. Pass `draft_weights.mask_token_id` into noise construction, including the remote-drafter path, so the metadata override actually takes effect.</comment>
<file context>
@@ -209,6 +224,14 @@ bool load_draft_gguf(const std::string & path,
+ // embedding silently destroys acceptance.
+ {
+ const uint32_t mask_meta = read_u32("dflash.mask_token_id", 0);
+ if (mask_meta != 0) out.mask_token_id = (int32_t)mask_meta;
+ }
</file context>
| return result; | ||
| } | ||
|
|
||
| // dflash: fused conv step. Same op id as ggml_ssm_conv; op_params[0] = 1 |
Contributor
There was a problem hiding this comment.
P3: Update this comment to say op_params[0] = 2; 1 selects SpecLA while 2 selects the fused step kernel, so the current explanation can mislead future callers or maintenance changes.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml.c, line 5894:
<comment>Update this comment to say `op_params[0] = 2`; `1` selects SpecLA while `2` selects the fused step kernel, so the current explanation can mislead future callers or maintenance changes.</comment>
<file context>
@@ -5891,6 +5891,53 @@ struct ggml_tensor * ggml_ssm_conv_tree(
return result;
}
+// dflash: fused conv step. Same op id as ggml_ssm_conv; op_params[0] = 1
+// marks step mode, srcs are (x, c, conv_state, conv_input_out).
+struct ggml_tensor * ggml_ssm_conv_step(
</file context>
Suggested change
| // dflash: fused conv step. Same op id as ggml_ssm_conv; op_params[0] = 1 | |
| // dflash: fused conv step. Same op id as ggml_ssm_conv; op_params[0] = 2 |
| const ggml_compute_params * params, | ||
| ggml_tensor * dst) { | ||
| // dflash: the fused step mode (ggml_ssm_conv_step) is CUDA/HIP only | ||
| GGML_ASSERT(ggml_get_op_params_i32(dst, 0) == 0 && "ggml_ssm_conv_step is not supported on CPU"); |
Contributor
There was a problem hiding this comment.
P3: The guard correctly rejects every non-standard SSM_CONV mode on CPU, but its comment and abort message name only ggml_ssm_conv_step. The op-param check == 0 also rejects the SpecLA heavy-light mode (ggml_ssm_conv_specla sets op_params[0]=1 via ggml_set_op_params_i32(result, 0, 1)), which trips the same assert and aborts with a message only about step mode. Broaden the message so a developer hitting this assert for mode 1 understands the actual condition.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp, line 9336:
<comment>The guard correctly rejects every non-standard SSM_CONV mode on CPU, but its comment and abort message name only `ggml_ssm_conv_step`. The op-param check `== 0` also rejects the SpecLA heavy-light mode (`ggml_ssm_conv_specla` sets op_params[0]=1 via ggml_set_op_params_i32(result, 0, 1)), which trips the same assert and aborts with a message only about step mode. Broaden the message so a developer hitting this assert for mode 1 understands the actual condition.</comment>
<file context>
@@ -9332,6 +9332,8 @@ void ggml_compute_forward_flash_attn_back(
const ggml_compute_params * params,
ggml_tensor * dst) {
+ // dflash: the fused step mode (ggml_ssm_conv_step) is CUDA/HIP only
+ GGML_ASSERT(ggml_get_op_params_i32(dst, 0) == 0 && "ggml_ssm_conv_step is not supported on CPU");
const ggml_tensor * src0 = dst->src[0]; // conv_x
const ggml_tensor * src1 = dst->src[1]; // conv1d.weight
</file context>
Suggested change
| GGML_ASSERT(ggml_get_op_params_i32(dst, 0) == 0 && "ggml_ssm_conv_step is not supported on CPU"); | |
| GGML_ASSERT(ggml_get_op_params_i32(dst, 0) == 0 && "SpecLA/step (ggml_ssm_conv_specla/ggml_ssm_conv_step) is not supported on CPU"); |
| if (gate_bias != nullptr) { | ||
| // raw-gate mode: g = exp(softplus(alpha_raw + bias) * A), beta = sigmoid(beta_raw) | ||
| const float a = g[gb_offset] + gate_bias[h_idx]; | ||
| const float sp = (a > 20.0f) ? a : logf(1.0f + expf(a)); |
Contributor
There was a problem hiding this comment.
P3: The raw-gate gate/beta math (softplus with the 20.0f clip, gate_A scaling, sigmoid) is duplicated between gated_delta_net_cuda and gated_delta_net_cuda_grouped_cols. Extract it into a shared __device__ __forceinline__ helper (or compute it in launch_gated_delta_net) so the two kernels cannot diverge on the threshold or formula.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cuda/gated_delta_net.cu, line 404:
<comment>The raw-gate gate/beta math (softplus with the 20.0f clip, gate_A scaling, sigmoid) is duplicated between `gated_delta_net_cuda` and `gated_delta_net_cuda_grouped_cols`. Extract it into a shared `__device__ __forceinline__` helper (or compute it in `launch_gated_delta_net`) so the two kernels cannot diverge on the threshold or formula.</comment>
<file context>
@@ -387,8 +398,16 @@ gated_delta_net_cuda_grouped_cols(const float * q,
+ if (gate_bias != nullptr) {
+ // raw-gate mode: g = exp(softplus(alpha_raw + bias) * A), beta = sigmoid(beta_raw)
+ const float a = g[gb_offset] + gate_bias[h_idx];
+ const float sp = (a > 20.0f) ? a : logf(1.0f + expf(a));
+ g_val = expf(sp * gate_A[h_idx]);
+ beta_val = 1.0f / (1.0f + expf(-beta[gb_offset]));
</file context>
Comment on lines
+5894
to
+5895
| // dflash: fused conv step. Same op id as ggml_ssm_conv; op_params[0] = 1 | ||
| // marks step mode, srcs are (x, c, conv_state, conv_input_out). |
Contributor
There was a problem hiding this comment.
P3: This comment states op_params[0] = 1 marks step mode, but the code that follows sets ggml_set_op_params_i32(result, 0, 2) and the CUDA dispatcher routes step mode on ggml_get_op_params_i32(dst, 0) == 2, while 1 belongs to the SpecLA variant (also set by ggml_ssm_conv_specla). The comment contradicts both the implementation and the dispatch, which will mislead anyone reading or extending this path.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml.c, line 5894:
<comment>This comment states `op_params[0] = 1` marks step mode, but the code that follows sets `ggml_set_op_params_i32(result, 0, 2)` and the CUDA dispatcher routes step mode on `ggml_get_op_params_i32(dst, 0) == 2`, while `1` belongs to the SpecLA variant (also set by `ggml_ssm_conv_specla`). The comment contradicts both the implementation and the dispatch, which will mislead anyone reading or extending this path.</comment>
<file context>
@@ -5891,6 +5891,53 @@ struct ggml_tensor * ggml_ssm_conv_tree(
return result;
}
+// dflash: fused conv step. Same op id as ggml_ssm_conv; op_params[0] = 1
+// marks step mode, srcs are (x, c, conv_state, conv_input_out).
+struct ggml_tensor * ggml_ssm_conv_step(
</file context>
Suggested change
| // dflash: fused conv step. Same op id as ggml_ssm_conv; op_params[0] = 1 | |
| // marks step mode, srcs are (x, c, conv_state, conv_input_out). | |
| // dflash: fused conv step. Same op id as ggml_ssm_conv; op_params[0] = 2 | |
| // marks step mode (1 = SpecLA heavy-light conv), srcs are (x, c, conv_state, conv_input_out). |
| char shape_err[192]; | ||
| const int64_t R = out.selector.rank; | ||
| if (!check_shape_2d(out.selector.hproj, out.n_embd, R, "selector.hproj", shape_err, sizeof(shape_err)) || | ||
| !check_shape_2d(out.selector.pred_cb, R, out.selector.pred_cb->ne[1], "selector.pred_cb", shape_err, sizeof(shape_err)) || |
Contributor
There was a problem hiding this comment.
P3: check_shape_2d(pred_cb, R, pred_cb->ne[1]) validates ne1 against pred_cb's own value, so a pred_cb with a nonsense vocabulary dimension loads without error. Drop the self-reference and validate pred_cb's ne1 against a real bound (e.g. the draft vocab_size / target vocab) and require succ_cb to match it.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/draft/draft_gguf_loader.cpp, line 546:
<comment>check_shape_2d(pred_cb, R, pred_cb->ne[1]) validates ne1 against pred_cb's own value, so a pred_cb with a nonsense vocabulary dimension loads without error. Drop the self-reference and validate pred_cb's ne1 against a real bound (e.g. the draft vocab_size / target vocab) and require succ_cb to match it.</comment>
<file context>
@@ -451,6 +497,64 @@ bool load_draft_gguf(const std::string & path,
+ char shape_err[192];
+ const int64_t R = out.selector.rank;
+ if (!check_shape_2d(out.selector.hproj, out.n_embd, R, "selector.hproj", shape_err, sizeof(shape_err)) ||
+ !check_shape_2d(out.selector.pred_cb, R, out.selector.pred_cb->ne[1], "selector.pred_cb", shape_err, sizeof(shape_err)) ||
+ !check_shape_2d(out.selector.succ_cb, R, out.selector.pred_cb->ne[1], "selector.succ_cb", shape_err, sizeof(shape_err))) {
+ set_last_error(shape_err);
</file context>
Contributor
Author
|
Wow Dflash2 makes the NVLinked 3090s faster than a 5090 in TG.
|
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.
This is tracking the current #625. When it was tested there was a regression. This is very much WIP. The work is done using just Qwen 3.8 on Lucebox and DeepSeek V4 Flash.
Date: 2026-08-20 · branch:
work/625-nvlink-fix-clean@f55de5a(base480e60a,"Merge origin/main into qwen38-dspark") · pushed to a fork of
Luce-Org/lucebox.What was broken
Tensor-split on the 2x3090 NVLink bench bed crashed on first generate, blocking
every spec path. Two independent crashes, two commits:
GGML_ASSERT(ret.axis != GGML_BACKEND_SPLIT_AXIS_UNKNOWN)atggml-backend-meta.cpp:570/556, op=SSM_CONV, during tensor-split graph allocf949e6a0x2000000000000000) in fused spec decode under tensor splitf55de5aCommit 1 —
f949e6aggml-meta: split SSM_CONV step/SpecLA by axis layout, not op_paramsserver/deps/llama.cpp/ggml/src/ggml-backend-meta.cpp(+12).handle_ssm_convdispatches on the actual src axis layout. Fused step mode (src0=AXIS_0 x [C,T,S],
src1=AXIS_1 weight [K,C], src2=AXIS_1 conv_state [K-1,C,S]) now returns
{AXIS_0,{0},1,{1}}— the same conv_out state layout as the proven non-fused path.Previously the handler demanded equal src axes, so the fused graph fell into
handle_generic→SPLIT_AXIS_UNKNOWN→ assert.Root-cause context: the crash is triggered by the stacked-projection target graph
(single target-graph file, present since
fa345d6); the meta split planning has norule for the fused-step layout. The fix mirrors what the non-fused/plain path
already did, keeping the conv_out state layout consistent.
Commit 2 —
f55de5afix(qwen35): resolve fused DSpark lm_head to draft-GPU simple tensorserver/src/qwen35/qwen35_backend.cpp(+51/-9). Under tensor split the targetlm_head is a meta tensor (data pointer is a
0x2000...placeholder). The fusedDSpark graphs run on the draft's plain CUDA backend, so
mul_mat(lm_head, hidden)read the fake pointer → illegal memory access (was reported at GET_ROWS node_2).
resolve_fused_lm_head()returnsggml_backend_meta_simple_tensor(lm_head, j)for the rank whose
layer_split_gpus[j] == draft_gpu(lm_head is a full-copymirror on every rank → local read is valid), or nullptr → host-chain fallback when
the draft GPU is not a target rank. Applied at both draft-backend call sites:
fused greedy chain (≈line 2907) and DDTree topk (≈line 3071) — the same latent
leak existed in both.
Partially a port of the deployed perf-line fix
fd7b294; the main line's DSparkimplementation (
cc75a80) predates it and merge480e60adid not carry it over.Verification (2x RTX 3090 + NVLink,
--target-split-mode tensor --peer-access, Q4_0 KV, greedy)Image baked from
f55de5a(sm_86). Prompt: 24538 prompt tokens, 1024 generated,temp 0, verify-width 8.
DFLASH_QWEN35_NO_FUSED_KERNELS=1producebyte-identical completions.
target_split=tensor peer_access=ON, 4x NVLink links.Post-fix: all cells complete, zero crash lines, fused spec paths active
(DSpark Markov head + DFlash2 selector both log active).
Open items (not blocking this gate)
quant/prompt mix). Speed-parity work item, tracked separately.
available on the bench host. (KV q8_0 loses to q4_0 on these 3090s, but
target-quant is a separate question.)
reviewer; result to follow.