vendor: grouped-expert MUL_MAT_ID for small MoE verify batches#501
Open
davide221 wants to merge 13 commits into
Open
vendor: grouped-expert MUL_MAT_ID for small MoE verify batches#501davide221 wants to merge 13 commits into
davide221 wants to merge 13 commits into
Conversation
![cubic-dev-ai[bot]](https://avatars.githubusercontent.com/in/1082092?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
All reported issues were addressed across 1 file
Reply with feedback, questions, or to request a fix.
Re-trigger cubic
howard0su
approved these changes
Jul 8, 2026
| // bit0 = Q4_K, bit1 = Q6_K, bit2 = Q4_0/Q8_0/Q5_K. Default Q4_K only: | ||
| // Q6_K keeps its tuned MMQ route above 5 tokens unless explicitly enabled. | ||
| static const int mask = []() { | ||
| const char * e = std::getenv("DFLASH_MMID_GROUPED_TYPES"); |
Contributor
There was a problem hiding this comment.
I suggest we don't keep experiement env if we believe this doesn't bring regression. too many env is so hard to maintain and keep track on them.
Contributor
There was a problem hiding this comment.
1 issue found across 5 files (changes from recent commits).
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/src/common/mmid_adaptive_k.h">
<violation number="1" location="server/src/common/mmid_adaptive_k.h:38">
P2: DFLASH_ADAPTIVE_K_DENSE has no effect for qwen35moe/gemma4 adaptive gating because this helper skips dense-list parsing whenever callers pass `il < 0`. Consider passing real layer indices for those architectures or avoiding advertising/wiring dense exclusions where they cannot be honored.</violation>
</file>
Reply with feedback, questions, or to request a fix.
Re-trigger cubic
| if (tau <= 0.0f || n_tokens < 2 || n_tokens > 16 || ids == nullptr || weights == nullptr) { | ||
| return; | ||
| } | ||
| if (il >= 0) { |
Contributor
There was a problem hiding this comment.
P2: DFLASH_ADAPTIVE_K_DENSE has no effect for qwen35moe/gemma4 adaptive gating because this helper skips dense-list parsing whenever callers pass il < 0. Consider passing real layer indices for those architectures or avoiding advertising/wiring dense exclusions where they cannot be honored.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/common/mmid_adaptive_k.h, line 38:
<comment>DFLASH_ADAPTIVE_K_DENSE has no effect for qwen35moe/gemma4 adaptive gating because this helper skips dense-list parsing whenever callers pass `il < 0`. Consider passing real layer indices for those architectures or avoiding advertising/wiring dense exclusions where they cannot be honored.</comment>
<file context>
@@ -0,0 +1,53 @@
+ if (tau <= 0.0f || n_tokens < 2 || n_tokens > 16 || ids == nullptr || weights == nullptr) {
+ return;
+ }
+ if (il >= 0) {
+ const char * e = std::getenv("DFLASH_ADAPTIVE_K_DENSE");
+ const std::string str = e ? e : (dense_default ? dense_default : "");
</file context>
Contributor
There was a problem hiding this comment.
1 issue found across 11 files (changes from recent commits).
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/topk-rows.cu">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/topk-rows.cu:132">
P1: GPU top-k can read stale or partially-written logits when the producing CUDA backend stream is not synchronized with the default stream. Consider launching this helper on the same backend stream or adding an explicit synchronization/event dependency before reading `logits->data`.</violation>
</file>
Tip: Review your code locally with the cubic CLI to iterate faster.
Re-trigger cubic
| d_dev = attr.device; | ||
| } | ||
|
|
||
| topk_rows_kernel<<<(int) nrows, TOPK_ROWS_BLOCK>>>( |
Contributor
There was a problem hiding this comment.
P1: GPU top-k can read stale or partially-written logits when the producing CUDA backend stream is not synchronized with the default stream. Consider launching this helper on the same backend stream or adding an explicit synchronization/event dependency before reading logits->data.
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/topk-rows.cu, line 132:
<comment>GPU top-k can read stale or partially-written logits when the producing CUDA backend stream is not synchronized with the default stream. Consider launching this helper on the same backend stream or adding an explicit synchronization/event dependency before reading `logits->data`.</comment>
<file context>
@@ -0,0 +1,145 @@
+ d_dev = attr.device;
+ }
+
+ topk_rows_kernel<<<(int) nrows, TOPK_ROWS_BLOCK>>>(
+ (const float *) logits->data, (int) ncols, d_vals, d_ids);
+ float h_vals[TOPK_ROWS_MAX_ROWS * TOPK_ROWS_K];
</file context>
Speculative-verify batches (2-16 tokens) route consecutive draft tokens to heavily overlapping expert sets, but the multi-token MoE MMVQ kernel reads each (token, slot) pair's expert weights independently. Measured on a 256-expert top-8 MoE, an 8-row verify batch reads 64 expert matrices per layer where only 39.4 distinct ones are needed (1.62x duplication); at 15 rows it is 120 vs 58.9 (2.04x). This adds an opt-in path (DFLASH_MMID_GROUPED=1) that rank-sorts the (token, slot) pairs by routed expert in a tiny capture-safe prep kernel, then runs a variant of mul_mat_vec_q_moe whose warp w processes sorted pair blockIdx.y*8+w instead of (slot, token). Same-expert pairs are adjacent after sorting, so the warps of a block share an expert and duplicate weight reads are served from L1/L2 instead of DRAM. Launch shape, per-warp structure, vec_dot sequence and warp reduction are identical to mul_mat_vec_q_moe, so outputs are bit-exact and CUDA-graph capture stays supported. DFLASH_MMID_GROUPED_TYPES bitmask selects quant types: 1 = Q4_K (default), 2 = Q6_K (also lifts the Q6_K MUL_MAT_ID MMVQ ceiling to 16, keeping CUDA graphs enabled for those batches), 4 = Q4_0/Q8_0/Q5_K. Model-agnostic for any MoE with n_expert_used <= 16. Measured on RTX 3090 (sm_86), Laguna XS 2.1 Q4_K_M target + DFlash drafter, greedy HumanEval prompts, outputs byte-identical: verify width 8: 212.0 -> 219.9 tok/s (+3.7%) verify width 15: ~157 -> 173.3 tok/s (~+10%) Gains grow with verify width (more expert overlap to dedup), which is what makes wide/hedged verify configurations cheaper.
…batches Speculative-verify tokens differ in difficulty: for most tokens the router mass concentrates in a few experts, and dropping the low-mass tail rarely changes the verified argmax. This adds DFLASH_ADAPTIVE_K_TAU=<0..1>: inside the grouped MUL_MAT_ID prep kernel, each token keeps its leading experts until their cumulative combine weight reaches tau; the remaining ids are sentineled to -1 (skipped by the grouped kernel, exact zero contribution) and the kept weights are renormalized in place. Requires DFLASH_MMID_GROUPED=1 with DFLASH_MMID_GROUPED_TYPES=7 so every routed quant type takes the grouped path. The graph-builder side tags the router ids tensor with the combine-weights tensor via ids->extra (server/src/common/mmid_adaptive_k.h). Wired for laguna (DFlash capture layers 1,13,25,33,39 kept dense by default), qwen35moe and gemma4 (dense layers via DFLASH_ADAPTIVE_K_DENSE). Prefill and single-token decode are untouched. Not bit-exact (near-lossless): tau trades a small per-token risk for speed. Measured on RTX 3090, Laguna XS 2.1 Q4_K_M + DFlash drafter, HumanEval-10 greedy, tau=0.80: 212.0 -> 224.5 tok/s (+6% combined with the grouped kernel), expected_pass 10/10 at every tau tested; tau=0.80 best. tau needs per-model validation; qwen35moe/gemma4 are mechanism-wired but not yet swept.
…decode Every verify row on an MoE target routes to its own top-k experts, so rows cost real weight bandwidth (~1ms/row for a Q4 target on a 3090). A fixed verify width pays that on every step; this trims the batch per step using the drafter's own confidence (keep slot j while the product of top-1 probs through j stays above DFLASH_ADAPTIVE_WIDTH_THETA). Output-exactness is preserved: only the number of speculated slots changes. - ggml-cuda: ggml_backend_cuda_topk_rows, a small top-k (k<=8) + softmax probability reduction over draft-head logits rows, ~KB readback instead of the full vocab logits. Also used to accelerate project_hidden_to_topk (draft-tree candidates) at unit temperature. - common: adaptive_verify_width.h helper + DFlashTarget hook returning per-slot candidate probs alongside draft tokens; domino fused head exposes its corrected logits for the same extraction. - laguna: reference integration (projection extraction, per-width verify graph slots so every width's CUDA graph captures once and replays, trim before verify_batch). Laguna XS 2.1 Q4_K_M + official DFlash drafter, 3090, HE-10 512-tok: theta=0.20 min=4 --verify-width 8 -> 236.0 tok/s vs 226.4 fixed w6 control (+4.2%), commit/step 3.58->3.79, avg rows 5.35, 10/10 pass. Off by default.
…lag for the rest Collapse the env sprawl from the previous three commits into a clean surface: - Grouped MUL_MAT_ID: bit-exact, so ON by default on CUDA (HIP stays opt-in, unvalidated). DFLASH_MMID_GROUPED=0 is the kill switch; the type mask defaults to all validated types and is a debug override. - Adaptive verify width: output-exactness-preserving, so ON by default (theta 0.20, min 4). Greedy chains now default to a base width of 8 rows trimmed per step by drafter confidence; the legacy accept-EWMA AUTO remains the fallback for theta=0, sampled verify and --ddtree. - Adaptive expert count (the only output-changing knob): promoted from env to --adaptive-experts [tau] (default 0.80). Explicit env still wins for sweeps. Net: a default server run gets the exact optimizations with zero configuration; one documented flag opts into the near-lossless expert gating.
… aux hidden norms, context-KV norms) The official Laguna XS 2.1 DFlash drafter carries per-layer attention gates (blk.<i>.attn_gate.weight), per-capture-layer aux hidden norms (dflash.aux_hidden_norm.<n>.weight) and a context-KV layer input norm flag. The draft loader on main silently ignored all three, so the drafter ran without its feature norms: the fc projection overflowed in f16 and every draft hidden state came out NaN, killing spec decode on any main build (laguna_verify_batch: embed failed on token id -1 from an all-NaN argmax). Port the loader + draft-graph support so main can run the official drafter.
davide221
force-pushed
the
perf/mmid-grouped-verify
branch
from
807f64f to
4a4d519
Compare
…ter requant tool - fused Domino head reads the draft graph's device hidden in place (hidden_dev param); the 64KB D2H hidden readback is now lazy and only paid by fallback heads / ddtree / sampled-verify paths - build_draft_step skips the per-step rebuild while ctx stays inside the same 64-aligned pad bucket (copy-mode builds only); laguna passes copy-mode by default, kill with DFLASH_DRAFT_PERSIST=0 - scripts/requant_dflash_draft.py: GGUF->GGUF drafter requant (q8/q4), metadata verbatim; q4 drafter measured +3% end-to-end on Laguna XS 2.1 (3090), acceptance unchanged - drafter is latency-bound, not bandwidth-bound - step-prof: commit/build laps close the loop-top blind spot
DFlash drafters re-encoded their whole feature window (up to draft_ctx_max
tokens of target hidden states) every step: ~10ms/step once the window
fills, on every model family, previously hidden behind slower targets.
common/dflash_draft_kv.{h,cpp}: per-layer head-major F16 ring caches
(slot = pos % cap) with absolute-position RoPE, so entries stay valid as
the window slides. One fixed-topology step graph (fold-in append of newly
committed rows via set_rows with a trash slot for pads, then the noise
forward flash-attending over the ring) built once, CUDA-graph replayed
forever; bulk append after prefill. All inputs live in a dedicated buffer
so gallocr never recycles them.
draft/draft_graph.cpp: build_draft_kv_append/build_draft_kv_step next to
the legacy one-shot builder, sharing draft_fuse_features. The per-position
k_norm and RoPE are separable, so append/step splitting matches the legacy
concat math exactly (f16 storage is the only numeric difference).
Any backend adopts it with ~20 lines: draft_kv_init once, then per step
draft_kv_begin_step + upload noise embed + compute.
Measured (laguna XS 2.1, RTX 3090): draft step 9.9 -> 2.4ms, flat at any
context; acceptance parity.
Laguna XS 2.1 Q4_K_M + q4 drafter on RTX 3090, cold-certified: decode 154.6 tok/s @30k and 152.3 @256k (8K pool, was 84.8/22.1), prefill 240K tokens in 67.3s (was 26.4 min), acceptance parity, ~1GB VRAM freed. - SWA ring caches [TAG_SWA_RING]: 30/40 layers attend a 512-token window but flash-attended the whole pooled span with 97% masked out. Window layers now keep a position-indexed ring (slot = pos % ring, ring sized from chunk+window) and bypass the pager entirely; the pager only pages full-attention layers. Verify 19.4 -> 13.7ms, prefill -27%. Snapshots are skipped when rings are active (they hold only the trailing window); kv_cap probes pick a full-attention layer. Kill: DFLASH_LAGUNA_SWA_RING=0. - Drafter ring-cache wiring behind DFLASH_DRAFT_KV (default on, =0 legacy; rebuilt on drafter-variant switch). - Pooled prefill batches multiple pager chunks per target step and the feature capture appends via set_rows with ring-row indices as input data, so prefill graphs CUDA-replay; --chunk 1024 feeds the MoE expert GEMMs better (-19% prefill wall at 240K). - Env-gated profilers: DFLASH_LAGUNA_STEP_PROF (commit/build/dwait laps), DFLASH_LAGUNA_VERIFY_PROF and DFLASH_LAGUNA_PREFILL_PROF (sub-phase laps), GGML_CUDA_GRAPH_STATS_EVERY + first-mismatch logging in ggml-cuda. Measured: the entire KVFlash host machinery costs 0.08ms/step; verify is 99.5% GPU kernel time.
Contributor
There was a problem hiding this comment.
5 issues found across 27 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/src/gemma4/gemma4_graph.cpp">
<violation number="1" location="server/src/gemma4/gemma4_graph.cpp:109">
P2: Gemma4 adaptive-experts can gate every MoE layer, including any DFlash capture/conditioning layers, because this attach call passes `il=-1` and no dense-layer default. Consider threading the Gemma4 layer index into `build_gemma4_moe_block` and passing an appropriate dense list so drafter features stay unchanged when `DFLASH_ADAPTIVE_K_TAU` is enabled.</violation>
</file>
<file name="server/deps/llama.cpp/ggml/include/ggml-cuda.h">
<violation number="1" location="server/deps/llama.cpp/ggml/include/ggml-cuda.h:48">
P3: The `[TAG_TOPK_ROWS]` prefix in the comment is inconsistent with every other function comment in this header, which uses plain descriptive text without bracket tags. Consider removing the tag or using it uniformly across the API surface to keep the documentation consistent.</violation>
</file>
<file name="server/src/laguna/laguna_internal.h">
<violation number="1" location="server/src/laguna/laguna_internal.h:308">
P1: SWA ring mode can corrupt or go out of bounds in graph builders that still leave `LagunaGraphInputs::kv_idx_swa` null. Consider requiring/filling `kv_idx_swa` whenever `cache.swa_ring_rows > 0` (or rejecting those paths) rather than letting ring-sized SWA caches fall back to pool indices.</violation>
</file>
<file name="server/src/common/dflash_draft_kv.cpp">
<violation number="1" location="server/src/common/dflash_draft_kv.cpp:290">
P2: SWA draft attention can include stale context tokens for later noise rows because `swa_lo` is fixed for the whole block instead of advancing per query position. Consider computing the SWA lower bound per `q` when filling `mask_swa`, so rows for `committed + q` enforce the same sliding-window semantics as the rest of the cache code.</violation>
</file>
<file name="server/src/server/server_main.cpp">
<violation number="1" location="server/src/server/server_main.cpp:446">
P2: `--adaptive-experts` should override the runtime tau consistently, but this `setenv(..., 0)` call leaves any pre-existing `DFLASH_ADAPTIVE_K_TAU` untouched. In deployments that already export the env var, the new CLI flag becomes a no-op.</violation>
</file>
Reply with feedback, questions, or to request a fix.
Re-trigger cubic
| // [TAG_SWA_RING] SWA ring rows (pos % cache.swa_ring_rows). When non-null, | ||
| // SWA layers write K/V through these indices and flash-attend over the | ||
| // ring span; attn_mask_swa must then be [swa_ring_rows, n_tokens]. | ||
| ggml_tensor * kv_idx_swa = nullptr; // [n_tokens] I32 |
Contributor
There was a problem hiding this comment.
P1: SWA ring mode can corrupt or go out of bounds in graph builders that still leave LagunaGraphInputs::kv_idx_swa null. Consider requiring/filling kv_idx_swa whenever cache.swa_ring_rows > 0 (or rejecting those paths) rather than letting ring-sized SWA caches fall back to pool indices.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/laguna/laguna_internal.h, line 308:
<comment>SWA ring mode can corrupt or go out of bounds in graph builders that still leave `LagunaGraphInputs::kv_idx_swa` null. Consider requiring/filling `kv_idx_swa` whenever `cache.swa_ring_rows > 0` (or rejecting those paths) rather than letting ring-sized SWA caches fall back to pool indices.</comment>
<file context>
@@ -290,6 +302,10 @@ struct LagunaGraphInputs {
+ // [TAG_SWA_RING] SWA ring rows (pos % cache.swa_ring_rows). When non-null,
+ // SWA layers write K/V through these indices and flash-attend over the
+ // ring span; attn_mask_swa must then be [swa_ring_rows, n_tokens].
+ ggml_tensor * kv_idx_swa = nullptr; // [n_tokens] I32
bool output_logits = true;
bool logits_are_output = true;
</file context>
| if (i + 1 < argc && argv[i + 1][0] != '-') { | ||
| tau = argv[++i]; | ||
| } | ||
| setenv("DFLASH_ADAPTIVE_K_TAU", tau, 0); // explicit env still wins |
Contributor
There was a problem hiding this comment.
P2: --adaptive-experts should override the runtime tau consistently, but this setenv(..., 0) call leaves any pre-existing DFLASH_ADAPTIVE_K_TAU untouched. In deployments that already export the env var, the new CLI flag becomes a no-op.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/server/server_main.cpp, line 446:
<comment>`--adaptive-experts` should override the runtime tau consistently, but this `setenv(..., 0)` call leaves any pre-existing `DFLASH_ADAPTIVE_K_TAU` untouched. In deployments that already export the env var, the new CLI flag becomes a no-op.</comment>
<file context>
@@ -435,6 +438,12 @@ int main(int argc, char ** argv) {
+ if (i + 1 < argc && argv[i + 1][0] != '-') {
+ tau = argv[++i];
+ }
+ setenv("DFLASH_ADAPTIVE_K_TAU", tau, 0); // explicit env still wins
} else if (std::strcmp(argv[i], "--verify-width") == 0 && i + 1 < argc) {
bargs.verify_width = std::atoi(argv[++i]);
</file context>
Suggested change
| setenv("DFLASH_ADAPTIVE_K_TAU", tau, 0); // explicit env still wins | |
| setenv("DFLASH_ADAPTIVE_K_TAU", tau, 1); // CLI flag overrides env default |
|
|
||
| GGML_BACKEND_API ggml_backend_reg_t ggml_backend_cuda_reg(void); | ||
|
|
||
| // [TAG_TOPK_ROWS] top-k (k <= 8) entries + softmax probabilities per row of a |
Contributor
There was a problem hiding this comment.
P3: The [TAG_TOPK_ROWS] prefix in the comment is inconsistent with every other function comment in this header, which uses plain descriptive text without bracket tags. Consider removing the tag or using it uniformly across the API surface to keep the documentation consistent.
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/include/ggml-cuda.h, line 48:
<comment>The `[TAG_TOPK_ROWS]` prefix in the comment is inconsistent with every other function comment in this header, which uses plain descriptive text without bracket tags. Consider removing the tag or using it uniformly across the API surface to keep the documentation consistent.</comment>
<file context>
@@ -45,6 +45,12 @@ GGML_BACKEND_API void ggml_backend_cuda_unregister_host_buffer(void * buffer);
GGML_BACKEND_API ggml_backend_reg_t ggml_backend_cuda_reg(void);
+// [TAG_TOPK_ROWS] top-k (k <= 8) entries + softmax probabilities per row of a
+// device-resident contiguous F32 [ncols, nrows] tensor; ~KB of readback. Not
+// a graph op: call after the producing graph_compute has returned.
</file context>
Suggested change
| // [TAG_TOPK_ROWS] top-k (k <= 8) entries + softmax probabilities per row of a | |
| // top-k (k <= 8) entries + softmax probabilities per row of a |
Wires the shared common/dflash_draft_kv module into the qwen35 backend's local-draft path (same pattern as laguna: lazy init, built_for guard on drafter switch, DFLASH_DRAFT_KV=0 kill switch, legacy path preserved). Remote-draft (IPC) requests keep the legacy path. Also adds distinct error messages to the bulk-append failure paths in dflash_draft_kv.cpp; a mismatched drafter now fails gracefully with the offending dims where the legacy path hit a GGML_ASSERT OOB write. A/B on Qwen3.6-27B Q4_K_M + dflash-draft-3.6-q4_k_m (RTX 3090): short 83.7 -> 96.3 tok/s (accept 38.0 -> 41.2%), 8K real code 26.8 -> 31.6 tok/s (19.3 -> 18.7%), outputs coherent both legs.
Contributor
There was a problem hiding this comment.
All reported issues were addressed across 3 files (changes from recent commits).
Reply with feedback, questions, or to request a fix.
Re-trigger cubic
- Reset the drafter ring at the start of every spec-decode request in both backends: the state persists across requests, so a new conversation could draft against the previous request's cached rows for positions below next_pos. The first begin_step re-appends the live window from the feature mirror (~10ms once per request). Validated: identical speed, acceptance and output across back-to-back requests. - StepGraph records built_view; the copy-mode persistent fast path refuses graphs whose topology reads features through a mirror ring view (reachable with DFLASH_DRAFT_PERSIST=0). - DraftKvState is now non-copyable: it owns raw ggml contexts, buffers and graphs, so a copy would double-free. - Document that the drafter SWA window is anchored at committed for all noise rows on purpose: it replicates the legacy one-shot graph's single windowed view, which is the trained drafter's validated attention pattern.
- Adaptive-K gates a scratch COPY of the router ids: the -1 drop sentinels never reach consumers on the non-grouped path (a sibling weight with a non-grouped quant type would cast them to uint32_t and read OOB). The shared zeroed gate weights keep every consumer bit-correct, and the already-gated marker is now weights-based since ids copies are per-op. - Grouped MUL_MAT_ID falls back to the legacy kernel above MMID_GROUPED_MAX_PAIRS instead of aborting on a hard assert. - q8 memoization re-enabled for MUL_MAT_ID (quantization is ids-independent and the memo key never included ids). - SWA ring guards: build_laguna_graph asserts ring => kv_idx_swa; laguna_layer_step refuses ring caches (unwired path fails loudly, not OOB). - Adaptive-K extras pooled per tensor address (was one leak per graph build); tau/DENSE/theta/min env knobs and --adaptive-experts get strict validated parsing; il<0 families warn once that dense exclusions cannot apply yet (layer-index threading tracked as follow-up). - Draft GGUF loader frees the metadata arena on all early validation failures; requant tool refuses src==dst (in-place corruption). - GGML_CUDA_GRAPH_STATS_EVERY clamped positive (division by zero); topk-rows sync contract + output buffer sizes documented. Validated: laguna ring A/B per-step parity (17.2/22.8 ms per step vs 17.3/23.1 pre-fix), opt-in grouped+adaptive-K leg 230 tok/s short / 126.7 @8K, coherent output, acceptance healthy.
…nv surface
The server had accumulated 153 distinct environment variables. Policy going
forward (docs/ENVIRONMENT.md): features ship as CLI flags or defaults; env
vars are reserved for burn-in kill switches (deleted after soak) and debug
instrumentation.
- DFLASH_PROF=step,verify,prefill replaces DFLASH_LAGUNA_{STEP,VERIFY,
PREFILL}_PROF (all three were new this branch; no users to migrate).
- docs/ENVIRONMENT.md documents the load-bearing variables (default,
purpose, lifecycle) and carries the full generated inventory; the
promotion/deletion audit of the legacy surface is tracked as follow-up.
Contributor
There was a problem hiding this comment.
3 issues found across 14 files (changes from recent commits).
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/mmvq.cu">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu:845">
P2: Adaptive-expert savings are lost for second/third grouped MMID ops after a prior op zeroes weights, because zero-weight slots mark the row as already gated but the scratch id copy keeps the original expert ids. Consider mirroring zero-weight markers into `idrow[j] = -1` before sorting so the grouped kernel takes its dropped-slot path.</violation>
</file>
<file name="server/docs/ENVIRONMENT.md">
<violation number="1" location="server/docs/ENVIRONMENT.md:1">
P2: The ENVIRONMENT.md entry for DFLASH_PROF claims it "replaces DFLASH_LAGUNA_{STEP,VERIFY,PREFILL}_PROF", but those three suffixed env vars are not actually read by the codebase — they only appear as TAG comments (server/src/laguna_target_graph.cpp). Only DFLASH_LAGUNA_PROFILE is getenv'd in practice (server/src/laguna_backend.cpp: ~2381, ~2411). This documentation claim is therefore misleading and could confuse maintainers looking for legacy env var behavior. Consider removing or correcting the statement to reflect the real runtime env var(s) and the true migration/supersession relationship.</violation>
</file>
<file name="server/src/common/mmid_adaptive_k.h">
<violation number="1" location="server/src/common/mmid_adaptive_k.h:85">
P2: The new process-wide pool is mutated without synchronization, so concurrent graph builds can race when they attach adaptive-k metadata for different tensors. A small mutex around the pool lookup/update would make the lifetime cache safe.</violation>
</file>
Tip: Review your code locally with the cubic CLI to iterate faster.
Re-trigger cubic
| float * wrow = gate_w + i*gate_w_stride; | ||
| bool gated = false; | ||
| for (int j = 0; j < n_slots; ++j) { | ||
| gated = gated || (idrow[j] < 0) || (wrow[j] == 0.0f); |
Contributor
There was a problem hiding this comment.
P2: Adaptive-expert savings are lost for second/third grouped MMID ops after a prior op zeroes weights, because zero-weight slots mark the row as already gated but the scratch id copy keeps the original expert ids. Consider mirroring zero-weight markers into idrow[j] = -1 before sorting so the grouped kernel takes its dropped-slot 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-cuda/mmvq.cu, line 845:
<comment>Adaptive-expert savings are lost for second/third grouped MMID ops after a prior op zeroes weights, because zero-weight slots mark the row as already gated but the scratch id copy keeps the original expert ids. Consider mirroring zero-weight markers into `idrow[j] = -1` before sorting so the grouped kernel takes its dropped-slot path.</comment>
<file context>
@@ -834,13 +834,15 @@ static __global__ void mmid_group_prep(
bool gated = false;
for (int j = 0; j < n_slots; ++j) {
- gated = gated || (idrow[j] < 0);
+ gated = gated || (idrow[j] < 0) || (wrow[j] == 0.0f);
}
if (!gated) {
</file context>
| @@ -0,0 +1,164 @@ | |||
| # Server environment variables | |||
Contributor
There was a problem hiding this comment.
P2: The ENVIRONMENT.md entry for DFLASH_PROF claims it "replaces DFLASH_LAGUNA_{STEP,VERIFY,PREFILL}_PROF", but those three suffixed env vars are not actually read by the codebase — they only appear as TAG comments (server/src/laguna_target_graph.cpp). Only DFLASH_LAGUNA_PROFILE is getenv'd in practice (server/src/laguna_backend.cpp: ~2381, ~2411). This documentation claim is therefore misleading and could confuse maintainers looking for legacy env var behavior. Consider removing or correcting the statement to reflect the real runtime env var(s) and the true migration/supersession relationship.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/docs/ENVIRONMENT.md:
<comment>The ENVIRONMENT.md entry for DFLASH_PROF claims it "replaces DFLASH_LAGUNA_{STEP,VERIFY,PREFILL}_PROF", but those three suffixed env vars are not actually read by the codebase — they only appear as TAG comments (server/src/laguna_target_graph.cpp). Only DFLASH_LAGUNA_PROFILE is getenv'd in practice (server/src/laguna_backend.cpp: ~2381, ~2411). This documentation claim is therefore misleading and could confuse maintainers looking for legacy env var behavior. Consider removing or correcting the statement to reflect the real runtime env var(s) and the true migration/supersession relationship.</comment>
<file context>
@@ -0,0 +1,164 @@
+# Server environment variables
+
+Policy (2026-07): **new features ship as CLI flags or defaults, not env vars.**
+Environment variables are reserved for two cases:
+
+1. **Burn-in kill switches** for freshly landed defaults - documented here with
+ the intent to delete them once the feature has soaked.
+2. **Debug instrumentation** (profilers, stats) - zero-cost when unset, never
+ required for correct serving.
</file context>
| // hands back the same address). Keep one allocation per distinct tensor | ||
| // address in a process-lifetime pool: bounded by the arenas' stable | ||
| // addresses instead of leaking one allocation per rebuild. | ||
| static std::unordered_map<const ggml_tensor *, mmid_gate_extra *> pool; |
Contributor
There was a problem hiding this comment.
P2: The new process-wide pool is mutated without synchronization, so concurrent graph builds can race when they attach adaptive-k metadata for different tensors. A small mutex around the pool lookup/update would make the lifetime cache safe.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/common/mmid_adaptive_k.h, line 85:
<comment>The new process-wide pool is mutated without synchronization, so concurrent graph builds can race when they attach adaptive-k metadata for different tensors. A small mutex around the pool lookup/update would make the lifetime cache safe.</comment>
<file context>
@@ -24,30 +26,66 @@ struct mmid_gate_extra {
+ // hands back the same address). Keep one allocation per distinct tensor
+ // address in a process-lifetime pool: bounded by the arenas' stable
+ // addresses instead of leaking one allocation per rebuild.
+ static std::unordered_map<const ggml_tensor *, mmid_gate_extra *> pool;
+ mmid_gate_extra *& gx = pool[ids];
+ if (gx == nullptr) gx = new mmid_gate_extra{MMID_GATE_MAGIC, tau, weights};
</file context>
Self-review pass over the review-fix batch: - malformed DFLASH_ADAPTIVE_K_DENSE entries warn once, not per graph build - the adaptive-K extras pool documents its single-threaded-builder assumption (same contract as the thread_local builder arenas) and states the boundedness argument honestly - two profiler comments still referenced the pre-consolidation env names - ENVIRONMENT.md: a raw pipe inside a table cell broke markdown rendering Validated on a full clean recompile: DFLASH_PROF=step emits the step profile (the rename's runtime behavior, previously masked by a stale object: rsync -a preserves source mtimes older than build objects, so make skipped the recompile - box-tree-only issue, the pushed branch was always verified via the fresh-clone build), opt-in grouped+adaptive leg byte-identical behavior.
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.
What
Patches the vendored ggml (
server/deps/llama.cpp).Opt-in grouped-expert path for
MUL_MAT_IDon small token batches (2-16), targeting speculative-verify workloads on quantized MoE models. Enable withDFLASH_MMID_GROUPED=1.Why
Consecutive draft tokens route to heavily overlapping expert sets, but
mul_mat_vec_q_moereads each (token, slot) pair's expert weights independently. Measured on a 256-expert top-8 MoE (Laguna XS 2.1):How
[TAG_MMID_GROUPED]inmmvq.cu). Capture-safe: ids never leave the device, all launch shapes are static.mul_mat_vec_q_moewith one change: warpwprocesses sorted pairblockIdx.y*8 + winstead of(slot = blockIdx.y, token = w). Same-expert pairs are adjacent after sorting, so the warps of a block share an expert and duplicate weight reads hit L1/L2 instead of DRAM.vec_dotsequence and warp reduction are identical tomul_mat_vec_q_moe→ bit-exact outputs, CUDA-graph capture unaffected.DFLASH_MMID_GROUPED_TYPESbitmask:1= Q4_K (default),2= Q6_K (also lifts Q6_K's MUL_MAT_ID MMVQ ceiling to 16, keeping CUDA graphs on for those batches),4= Q4_0/Q8_0/Q5_K.n_expert_used <= 16(asserted).Results
RTX 3090 (sm_86), Laguna XS 2.1 Q4_K_M + DFlash drafter, greedy HumanEval prompts. Outputs byte-identical to baseline in all cases (verified per-prompt, identical accept trajectories):
Gains grow with verify width (more overlap to dedup), which lowers the marginal cost of wide/hedged verify configurations.
Notes
ggml_cuda_mul_mat_vec_q).Update: opt-in per-token adaptive expert count (second commit)
On top of the (bit-exact, default-off) grouped kernel,
DFLASH_ADAPTIVE_K_TAU=<0..1>enables per-token expert-count gating inside the grouped prep kernel: each verify token keeps its leading experts until cumulative router mass ≥ τ, the rest are sentineled to-1(skipped, exact zero contribution) and kept weights are renormalized in place. Near-lossless, not bit-exact — off by default.server/src/common/mmid_adaptive_k.htags the router ids tensor with the combine-weights tensor viaids->extra. Wired for laguna (DFlash capture layers1,13,25,33,39dense by default), qwen35moe and gemma4 (DFLASH_ADAPTIVE_K_DENSEfor exclusions). deepseek4 (DSpark routing) not wired.DFLASH_MMID_GROUPED_TYPES=7recommended so all routed quant types participate).Measured (RTX 3090, Laguna XS 2.1 Q4_K_M + DFlash, HumanEval-10 greedy):
τ sweep: 0.90 ≈ no gain, 0.80 best, 0.70 starts costing accept length. τ is a per-model constant: laguna is validated, qwen35moe/gemma4 carry the mechanism but need their own 30-minute sweep before enabling.
Width note: with both optimizations on, the verify-width optimum shifts from 8 to 6 (HE-10: τ=0.80 w6 = 226.9 tok/s vs w8 = 224.5; true pre-update best was w5 = 217.1). The width curve is a flat plateau across w5-w8.
Update: drafter-confidence adaptive verify width (third commit)
The width plateau above exists because a fixed verify width averages two regimes: on confident steps deep slots commit often (w8 worth paying), on unsure steps they almost never do (w4 is enough). Every MoE verify row routes to its own top-k experts (~1ms/row of expert weight reads for a Q4 target on a 3090), so rows are the dominant marginal cost.
DFLASH_ADAPTIVE_WIDTH_THETA=<0..1>trims the batch per step: keep candidate slot j while the drafter's own reach-mass (product of top-1 probabilities through slot j) stays above θ, floored atDFLASH_ADAPTIVE_WIDTH_MIN(default 4) rows. Output-exactness is preserved — only the number of speculated slots changes; every committed token is still target-verified.ggml-cuda: newggml_backend_cuda_topk_rows(topk-rows.cu) — top-k (k ≤ 8) + softmax probabilities per row of the draft-head logits, ~KB readback instead of the full-vocab logits round-trip. Also acceleratesproject_hidden_to_topk(draft-tree candidates) at unit temperature.common:adaptive_verify_width.h(model-agnostic helper — any family loop with per-slot drafter top-1 probs can call it) + aDFlashTarget::project_hidden_to_tokens_topkhook that returns candidate probs alongside draft tokens; the fused domino head exposes its corrected logits for the same extraction.laguna(reference integration): candidate extraction rides the existing draft-token projection (no extra matmul), per-width verify-graph slots so each row count's CUDA graph captures once and replays, trim applied right beforeverify_batch.Measured (RTX 3090, Laguna XS 2.1 Q4_K_M + official DFlash drafter, HumanEval-10 greedy, 512 tok, same-sweep control):
--verify-width 8θ sweep: 0.15/0.25/0.30 and base widths 10 all worse than θ=0.20 w8; min 5 hurts. Mechanism confirmed by the row stats: confident steps stretch to 8 rows (commit 3.58 → 3.79) while the average row count lands below the fixed-w6 control. Cumulative PR gain vs the true pre-PR baseline (w5, 217.1): +8.7%.
Off by default; θ unset = existing fixed-width behavior. qwen35moe/gemma4 carry the helper + hook but need their loop wiring and a θ sweep before enabling.
Update: rebased onto main @ 5e302cb + two more commits
fix(draft)— Laguna XS 2.1 drafter support (4a4d519). Runtime validation on a main build (all previous numbers were measured on a lucebox working tree) exposed that main's draft loader ignores the XS 2.1 drafter's attention gates, aux hidden norms and context-KV input norms — that support only existed as uncommitted edits on the lucebox checkout. Without the feature norms the fc projection overflows in f16 and every draft hidden state is NaN, so spec decode was broken on any main build (laguna_verify_batch: embed failedfrom an all-NaN argmax). This commit ports the loader + draft-graph support; main can now run the official drafter.server:defaults consolidation (1238629). Collapses the env sprawl: the two exact optimizations are now on by default (grouped MUL_MAT_ID on CUDA, adaptive verify width θ=0.20 with base width 8 for greedy chains), envs demoted to kill-switches/debug overrides. The one output-changing knob is a CLI flag:--adaptive-experts [tau](default 0.80). A stock server invocation gets the validated exact stack with zero configuration.Measured on the rebased branch, main build, RTX 3090, HE-10 512-tok, 10/10 pass both:
--adaptive-expertsWarm per-request decode with
--adaptive-experts: 248-258 tok/s (server-side). Numbers are a hot-GPU session on the .105 box (known VRAM-throttle drift); same-protocol working-tree measurements earlier the same day: fixed-w6 control 226.4 → full stack 236.0.