perf(ds4): accelerate ROCm decode and long-context selection - #667
perf(ds4): accelerate ROCm decode and long-context selection#667davide221 wants to merge 2 commits into
Conversation
lucebox8 gfx1151 A/B (
|
| Prompt | Metric | OFF | ON | Delta |
|---|---|---|---|---|
| 32,768 | server prefill (median, n=2) | 184.56 | 191.23 tok/s | +3.62% |
| 32,768 | client decode (median, n=2) | 31.771 | 31.938 tok/s | +0.53% (noise) |
| 122,879 | server prefill | 129.46 | 133.43 tok/s | +3.07% |
| 122,879 | client decode | 26.875 | 26.540 tok/s | -1.25% (no gain) |
Correctness/integrity:
- All normal runs returned the identical expected response hash
a91b12613458d22a3fbf6d66ff2fba42124b46ea0e032f68c3564cc0c506ec43. - Captures completed with zero dropped steps, requests, or token bursts.
- 431/431 server tests and 5/5 focused profiling/GPU tests passed, including
GPU numerical equivalence forggml_moe_combine.
Matched ROCprof at 32K confirms the intended replacement: roughly 3,675 old
reduce/multiply launches disappear and 3,643 moe_combine_kernel launches are
added, saving about 1.36 s of sampled GPU work. Normal runs—not ROCprof—are the
throughput source.
Conclusion: reusable prefill win, decode-neutral. Keep this draft and
opt-in while the next gfx1151 work targets dominant quantized matmul and
grouped compact attention kernels.
|
Thank you, @davide221, for your work on Lucebox, for pushing this optimization series, and for acting on the earlier Strix Halo feedback. We qualified PR #667 head Test configuration
Matched results
The PR #667 build also fit the full 131,072-token configuration in 128 GiB. Minimum effective non-CMA headroom was 4.984 GiB. There was no OOM, GPU reset, server restart, or kernel memory-pressure event. For comparison, the stacked PR #664 base could not complete the same exact-128K fit because of a global-memory OOM. Quality findingQuality was evaluated with The previous setup stopped naturally after 1,974 generated tokens and returned For independent reproduction, our v1.5.0 release packages the exact Ansible configurations for both tested Strix Halo systems together with the benchmark definitions, pinned revisions, result data, and comparison chart. |
|
Split-KV MLA follow-up (
Retained captures:
Rejected: four-way split (repeat client deltas were only +0.17%/+0.32%) and the earlier 2-wave ROCmFP MMQ tile (+0.07%). |
|
Strix Halo follow-up at fb2f462 (PR remains draft)\n\nResult\n- Default gfx1151 path: 39.403 tok/s at exact 8K versus 37.935 tok/s with split-KV explicitly disabled: +3.87%.\n- Exact 32K observation: 36.195 tok/s.\n- All runs used DS4 model-default six routed experts, fixed q=4 verification, 128 output tokens, and the same response SHA256: 0f785a7ffa406498aafb14553966eaed0f52220fed0f7cc016b66921d104d194.\n- Split-KV is now automatic only on gfx1151. Other devices remain opt-in; GGML_CUDA_MLA_NO_SPLIT_KV and GGML_DS4_FA_NO_SPLIT_KV are opt-outs.\n\nMatched ROCprof, exact 8K, baseline to combined candidate\n- speculative wall: 4303.159 to 4019.416 ms (-6.59%)\n- verifier marker median: 104.531 to 96.781 ms (-7.41%)\n- summed GPU kernels per round: 91.441 to 83.405 ms (-8.79%)\n- ROCmFP2 q106 expert kernels: 1071.001 to 938.965 ms (-12.33%)\n- ROCmFP3 q105 expert kernels: 347.337 to 336.596 ms (-3.09%)\n- indexed attention: 191.650 to 82.393 ms (-57.01%)\n\nThe new ROCmFP2 and ROCmFP3 helpers load each activation once for two independent output rows. ROCmFP3 uses explicit FMA to preserve the previous raw float hashes; its fused-GLU specialization keeps the lower-register path because pairing regressed occupancy.\n\nValidation\n- relevant HIP targets build on ROCm 7.2.4 / gfx1151\n- test_rocmfp_mix_gateup_glu: pass\n- learned and fixed ROCmFP2/3 microbench hashes match their controls\n- full response hash matches at 8K and 32K\n- artifacts: /data/lucegraph-ds4-splitkv-runs/pairloop-q3-exact-rocprof-fixedq4-8k-r1, /data/lucegraph-ds4-splitkv-runs/pairloop-q3-exact-default-split-8k-r1, /data/lucegraph-ds4-splitkv-runs/pairloop-q3-exact-fixedq4-32k-r1 |
q5 qualification: adaptive width regresses real-code decode throughputI completed the q5 variable-acceptance evaluation requested in the PR description. Scope
The two adaptive process replicates were unusually consistent at 37.112 and 37.091 weighted tok/s. Fixed-width replicates were 39.892 and 40.084 weighted tok/s. The decode regression therefore does not look like a single-process outlier. The server explicitly logged Conclusion: the current PR667 q5 controller is not a speed win for real coding traffic on this Lucebox setup and should remain disabled in the private profile. It sometimes improves an individual low-acceptance turn, but the additive probe/backoff policy is not Pareto-better across the workload. This was a development-only matched A/B on lucebox6 with the appliance and memory-clock qualification gates bypassed because the host remains on BIOS 03.04 balanced/auto/auto. That prevents an absolute product claim, but it does not explain the matched relative regression. Artifacts: Maintained harness: |
|
Follow-up/correction on the code-workflow A/B above: all 48 measured responses across the four runs had and (the four warmups likewise hit their 128-token cap).\n\nSo the result is valid evidence for sustained decode on cumulative coding prompts with natural, non-100% acceptance, but it should not be described as completed natural-EOS workflow performance or as a quality-equivalent user-task comparison. The checked-in benchmark is being hardened to report length-limited request counts explicitly and to retain outputs/hashes for separate semantic review.\n\nThe controller conclusion for that measured regime is unchanged: PR667 q5 adaptive was slower than fixed width (37.102 vs 39.988 weighted server tok/s, -7.22%). A natural-EOS/targeted-patch follow-up is needed before generalizing beyond sustained decode. |
|
Natural-completion follow-up to the token-capped cumulative-workflow result above. Scope
Natural-EOS resultTwo 10-prompt repeats per mode, 20 requests per mode. All 40 measured requests completed with
Replicates were stable:
However, only 1/10 response hashes matched between fixed and adaptive in each paired repeat. Each mode was internally deterministic (10/10 repeat hashes matched), so the controller changes outputs on this approximate sparse-prefill/top-k4 profile. HumanEval gold tests were retained but deliberately not executed by the serving client; this run makes no quality-equivalence claim. Report SHA-256:
Artifacts remain on lucebox6 at DecisionThis is not a material product-speed win and is nowhere near 2x. Natural short code completions are effectively flat end to end (+0.38%) with a slightly worse p10, while the cumulative 512-token coding workload regressed weighted server decode by 7.22%. Keep PR667 adaptive q5 disabled in the Lucebox customer profile unless a future controller passes a broader quality-scored workload battery with a clear Pareto improvement. |
|
Draft qualification update — not ready to merge. This commit adds two narrowly scoped gfx1151 defaults for DeepSeek V4 Flash DSpark:
Matched fixed-q4, six-routed-expert results on lucebox8, ROCm 7.2.4, Strix Halo gfx1151:
The final 32K and 123K runs left the sparse-attention environment variable unset and logged the gfx1151 default activation. Both completed 128 tokens with SHA-256 LuceGraph/ROCprof attribution on the matched 32K verifier showed:
Correctness/rollback checks:
The retained disposable captures and generated host/GPU reports are under |
f456d8b to
5eb4fbe
Compare
![cubic-dev-ai[bot]](https://avatars.githubusercontent.com/in/1082092?s=60&v=4){kind=small}
There was a problem hiding this comment.
9 issues found across 24 files
You’re at about 95% of the monthly reviewed-line limit. You may want to disable incremental reviews to conserve quota. Reviews will continue until that limit is exceeded. If you need help avoiding interruptions, please contact contact@cubic.dev.
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/rocmfp3_mix.cu">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/rocmfp3_mix.cu:519">
P2: The paired path changes the per-term rounding sequence: `mix_block_accum2` uses explicit `fmaf`, while `mix_block_accum` uses the chained multiply expression. On a build that does not contract the reference expression, this changes dot products and can change greedy tokens; keep both helpers' arithmetic sequence identical.</violation>
</file>
<file name="server/src/deepseek4/deepseek4_fused_verify.inc">
<violation number="1" location="server/src/deepseek4/deepseek4_fused_verify.inc:1270">
P2: When `q - 1 - i` exceeds `w.n_swa`, this wrap calculation extends past the raw-ring segment and masks compressed or preserved rows that should remain visible. Handle `future >= w.n_swa` as a full raw-ring mask, and only apply the two wrapped ranges when `future < w.n_swa`.</violation>
</file>
<file name="server/docs/ENVIRONMENT.md">
<violation number="1" location="server/docs/ENVIRONMENT.md:51">
P3: Two of the eight newly documented DFLASH_* variables — `DFLASH_GFX1151_HC_MMVF_Q4` (table line 51, defined in ggml-cuda/mmvf.cu) and `DFLASH_ROCMFP2_ROW4` (table line 65, defined in ggml-cuda/rocmfp2_mix.cu) — are missing from the "Full inventory (generated)" section, while the other new variables in this PR were added to it. Anyone searching the index finds no cross-reference for these kill switches. Add both to the inventory with their defining file to keep the documented table and the index consistent.</violation>
</file>
<file name="server/test/bench_ds4_topk.cpp">
<violation number="1" location="server/test/bench_ds4_topk.cpp:162">
P3: This standalone bench re-implements the long-context TOP_K selected-set parity check (full sort vs block-radix) and timing that already lives in the established framework, server/tests/test_deepseek4_unit.cpp (~lines 3020-3125), which registers in ctest and is what the PR cites as passing. Keeping a separate executable with its own correctness gate means the selected-set contract must be maintained in two places and can silently drift. Move the exact-set verification into the existing framework's long-context top-k test and keep this binary for perf-only measurement over the tile sweep.</violation>
</file>
<file name="server/src/qwen35/qwen35_backend.cpp">
<violation number="1" location="server/src/qwen35/qwen35_backend.cpp:4002">
P2: When the DSpark confidence gate shortens `v_len` while shared adaptive width is off, this branch still divides by the configured `verify_cap` instead of the actual offered candidates. That under-reports acceptance and can mis-steer the PFlash residency bandit; use `n_spec_offered_sum` in all three exit paths.
(Based on your team's feedback about offered-candidate acceptance-rate denominators.)</violation>
</file>
<file name="server/src/deepseek4/deepseek4_backend.cpp">
<violation number="1" location="server/src/deepseek4/deepseek4_backend.cpp:238">
P2: When `DFLASH_DS4_SPEC=1` but the drafter is absent or fails to load, `init()` falls back to autoregressive decoding while this guard still enables SparseFlash attention. Enable the default only after a DSpark drafter loads successfully, or gate it on the actual speculative-decoder state.</violation>
</file>
<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/mmvf.cu">
<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/mmvf.cu:852">
P2: On gfx1150 and other RDNA3.5 devices, this defaults q4 hyper-connection projections to MMVF despite qualification being gfx1151-only. Gate this branch with the exact gfx1151 compute capability so unmeasured devices retain the generic crossover.</violation>
</file>
<file name="server/test/test_rocmfp_mix_gateup_glu.cpp">
<violation number="1" location="server/test/test_rocmfp_mix_gateup_glu.cpp:125">
P3: The new comment and the ntok=3 change assume ntok=3 always exercises the two-pass GLU finalizer, but that only holds on a gfx1151 strix-tuned device. In rocmfp2_mix.cu's ggml_cuda_rocmfp2_mix_mul_mat_id_glu, the one-pass kernel is selected when `!strix_tuned || n_tokens <= 2` where `strix_tuned = eu.gfx1151 && eg.gfx1151`. On any non-gfx1151 HIP device, ntok=3 still takes the one-pass path, so the two-pass finalizer is never covered there and the new one_tok block duplicates the same path (both use the `<1,2>` kernel). The test only validates the two-pass branch when run on the gfx1151 target; consider gating this test (or the ntok=3 case) on the gfx1151 device so the coverage claim matches behavior.</violation>
</file>
<file name="server/test/bench_rocmfp_mix_gateup_glu.cpp">
<violation number="1" location="server/test/bench_rocmfp_mix_gateup_glu.cpp:233">
P3: The new verification block claims it "catches mode-specific numerical drift," but it never compares any of the hashes against a reference or against each other — it only prints them. It also does not check the actual fusion invariant: fused_result (swiGLU output) is never compared to anything recomputed from the raw up/gate projections, so a wrong fusion still passes silently. For the claims in the comment to hold, compute silu(up_result)*gate_result on host and compare (with tolerance) to fused_result, or assert the raw/fused hashes match an expected value; otherwise reword the comment. This matters because the PR relies on these benchmarks to gate the fusion.</violation>
</file>
Reply with feedback, questions, or to request a fix.
Re-trigger cubic
| acc0 = fmaf(rs0 * mix_fp3_fixed(mix_fp3_code(qs0, j)), x, acc0); | ||
| acc1 = fmaf(rs1 * mix_fp3_fixed(mix_fp3_code(qs1, j)), x, acc1); |
There was a problem hiding this comment.
P2: The paired path changes the per-term rounding sequence: mix_block_accum2 uses explicit fmaf, while mix_block_accum uses the chained multiply expression. On a build that does not contract the reference expression, this changes dot products and can change greedy tokens; keep both helpers' arithmetic sequence identical.
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/rocmfp3_mix.cu, line 519:
<comment>The paired path changes the per-term rounding sequence: `mix_block_accum2` uses explicit `fmaf`, while `mix_block_accum` uses the chained multiply expression. On a build that does not contract the reference expression, this changes dot products and can change greedy tokens; keep both helpers' arithmetic sequence identical.</comment>
<file context>
@@ -435,7 +476,65 @@ __device__ __forceinline__ void mix_block_accum(
+ const float x = xc[col0 + j];
+ const float rs0 = (j < MIX_QK/2) ? s00 : s01;
+ const float rs1 = (j < MIX_QK/2) ? s10 : s11;
+ acc0 = fmaf(rs0 * mix_fp3_fixed(mix_fp3_code(qs0, j)), x, acc0);
+ acc1 = fmaf(rs1 * mix_fp3_fixed(mix_fp3_code(qs1, j)), x, acc1);
+ }
</file context>
| acc0 = fmaf(rs0 * mix_fp3_fixed(mix_fp3_code(qs0, j)), x, acc0); | |
| acc1 = fmaf(rs1 * mix_fp3_fixed(mix_fp3_code(qs1, j)), x, acc1); | |
| acc0 += rs0 * mix_fp3_fixed(mix_fp3_code(qs0, j)) * x; | |
| acc1 += rs1 * mix_fp3_fixed(mix_fp3_code(qs1, j)) * x; |
| if (future > 0) { | ||
| const int begin = (pos_i + 1) % w.n_swa; | ||
| const int first = std::min(future, w.n_swa - begin); | ||
| add_range(col + (size_t) begin, | ||
| col + (size_t) (begin + first)); | ||
| add_range(col, |
There was a problem hiding this comment.
P2: When q - 1 - i exceeds w.n_swa, this wrap calculation extends past the raw-ring segment and masks compressed or preserved rows that should remain visible. Handle future >= w.n_swa as a full raw-ring mask, and only apply the two wrapped ranges when future < w.n_swa.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/deepseek4/deepseek4_fused_verify.inc, line 1270:
<comment>When `q - 1 - i` exceeds `w.n_swa`, this wrap calculation extends past the raw-ring segment and masks compressed or preserved rows that should remain visible. Handle `future >= w.n_swa` as a full raw-ring mask, and only apply the two wrapped ranges when `future < w.n_swa`.</comment>
<file context>
@@ -1249,34 +1255,97 @@ static int ds4_try_fused_verify_step(
+ col + (size_t) w.n_swa);
+ } else {
+ const int future = e - 1 - pos_i;
+ if (future > 0) {
+ const int begin = (pos_i + 1) % w.n_swa;
+ const int first = std::min(future, w.n_swa - begin);
</file context>
| if (future > 0) { | |
| const int begin = (pos_i + 1) % w.n_swa; | |
| const int first = std::min(future, w.n_swa - begin); | |
| add_range(col + (size_t) begin, | |
| col + (size_t) (begin + first)); | |
| add_range(col, | |
| if (future > 0) { | |
| if (future >= w.n_swa) { | |
| add_range(col, col + (size_t) w.n_swa); | |
| } else { | |
| const int begin = (pos_i + 1) % w.n_swa; | |
| const int first = std::min(future, w.n_swa - begin); | |
| add_range(col + (size_t) begin, | |
| col + (size_t) (begin + first)); | |
| add_range(col, | |
| col + (size_t) (future - first)); | |
| } | |
| } |
| cache_.last_tok = out_tokens.empty() ? last_tok : out_tokens.back(); | ||
| const int total_draft_pos = std::max(1, n_spec_steps * verify_cap); | ||
| const int total_draft_pos = shared_feedback_width | ||
| ? std::max(1, n_spec_offered_sum) |
There was a problem hiding this comment.
P2: When the DSpark confidence gate shortens v_len while shared adaptive width is off, this branch still divides by the configured verify_cap instead of the actual offered candidates. That under-reports acceptance and can mis-steer the PFlash residency bandit; use n_spec_offered_sum in all three exit paths.
(Based on your team's feedback about offered-candidate acceptance-rate denominators.)
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/qwen35/qwen35_backend.cpp, line 4002:
<comment>When the DSpark confidence gate shortens `v_len` while shared adaptive width is off, this branch still divides by the configured `verify_cap` instead of the actual offered candidates. That under-reports acceptance and can mis-steer the PFlash residency bandit; use `n_spec_offered_sum` in all three exit paths.
(Based on your team's feedback about offered-candidate acceptance-rate denominators.) </comment>
<file context>
@@ -3983,7 +3998,9 @@ bool Qwen35Backend::do_spec_decode(int committed, int n_gen,
cache_.last_tok = out_tokens.empty() ? last_tok : out_tokens.back();
- const int total_draft_pos = std::max(1, n_spec_steps * verify_cap);
+ const int total_draft_pos = shared_feedback_width
+ ? std::max(1, n_spec_offered_sum)
+ : std::max(1, n_spec_steps * verify_cap);
out_accept_rate =
</file context>
| // Only the DSpark verifier configuration has completed model-level | ||
| // qualification. Preserve an explicit user choice, including the burn-in | ||
| // kill switch value 0. | ||
| if (!env_flag_enabled("DFLASH_DS4_SPEC") || |
There was a problem hiding this comment.
P2: When DFLASH_DS4_SPEC=1 but the drafter is absent or fails to load, init() falls back to autoregressive decoding while this guard still enables SparseFlash attention. Enable the default only after a DSpark drafter loads successfully, or gate it on the actual speculative-decoder state.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/src/deepseek4/deepseek4_backend.cpp, line 238:
<comment>When `DFLASH_DS4_SPEC=1` but the drafter is absent or fails to load, `init()` falls back to autoregressive decoding while this guard still enables SparseFlash attention. Enable the default only after a DSpark drafter loads successfully, or gate it on the actual speculative-decoder state.</comment>
<file context>
@@ -230,6 +230,32 @@ static bool configure_dspark_mmvq_defaults(int gpu) {
+ // Only the DSpark verifier configuration has completed model-level
+ // qualification. Preserve an explicit user choice, including the burn-in
+ // kill switch value 0.
+ if (!env_flag_enabled("DFLASH_DS4_SPEC") ||
+ std::getenv("DFLASH_DS4_SPARSE_DECODE_FLASH") != nullptr) {
+ return;
</file context>
| "DFLASH_GFX1151_HC_MMVF_Q4"); | ||
| return !value || std::atoi(value) != 0; | ||
| }(); | ||
| if (gfx1151_hc_q4_mmvf && GGML_CUDA_CC_IS_RDNA3_5(cc) && |
There was a problem hiding this comment.
P2: On gfx1150 and other RDNA3.5 devices, this defaults q4 hyper-connection projections to MMVF despite qualification being gfx1151-only. Gate this branch with the exact gfx1151 compute capability so unmeasured devices retain the generic crossover.
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/mmvf.cu, line 852:
<comment>On gfx1150 and other RDNA3.5 devices, this defaults q4 hyper-connection projections to MMVF despite qualification being gfx1151-only. Gate this branch with the exact gfx1151 compute capability so unmeasured devices retain the generic crossover.</comment>
<file context>
@@ -834,6 +836,23 @@ bool ggml_cuda_should_use_mmvf(enum ggml_type type, int cc, const int64_t * src0
+ "DFLASH_GFX1151_HC_MMVF_Q4");
+ return !value || std::atoi(value) != 0;
+ }();
+ if (gfx1151_hc_q4_mmvf && GGML_CUDA_CC_IS_RDNA3_5(cc) &&
+ src0_ne[0] == 16384 && src0_ne[1] == 24 && ne11 == 4) {
+ return true;
</file context>
| if (gfx1151_hc_q4_mmvf && GGML_CUDA_CC_IS_RDNA3_5(cc) && | |
| if (gfx1151_hc_q4_mmvf && cc == GGML_CUDA_CC_OFFSET_AMD + 0x1151 && |
| | `DFLASH_DS4_VERIFY_FORCE_GRAPH_REPLAY` | unset | OPT-IN: bypass graph property scans only after warmup; scheduler-generation checks remain mandatory. | | ||
| | `DFLASH_DS4_ROCTX` | unset | DEBUG: on HIP builds, dynamically load ROCTX and emit semantic DS4 prefill, speculative-decode, and layer-range markers for external rocprof traces. No events, timing, or device synchronization are added. | | ||
| | `DFLASH_QWEN35_ROCTX` | unset | DEBUG: on HIP builds, dynamically load ROCTX and mark Qwen concurrent steps, graph compute, and argmax readback with live, padded, and packed-prefill shape metadata. | | ||
| | `DFLASH_GFX1151_HC_MMVF_Q4` | 1 on gfx1151 for the DS4 `[16384,24]` q4 projection | BURN-IN KILL SWITCH: =0 restores the generic hipBLAS dispatch decision. | |
There was a problem hiding this comment.
P3: Two of the eight newly documented DFLASH_* variables — DFLASH_GFX1151_HC_MMVF_Q4 (table line 51, defined in ggml-cuda/mmvf.cu) and DFLASH_ROCMFP2_ROW4 (table line 65, defined in ggml-cuda/rocmfp2_mix.cu) — are missing from the "Full inventory (generated)" section, while the other new variables in this PR were added to it. Anyone searching the index finds no cross-reference for these kill switches. Add both to the inventory with their defining file to keep the documented table and the index consistent.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/docs/ENVIRONMENT.md, line 51:
<comment>Two of the eight newly documented DFLASH_* variables — `DFLASH_GFX1151_HC_MMVF_Q4` (table line 51, defined in ggml-cuda/mmvf.cu) and `DFLASH_ROCMFP2_ROW4` (table line 65, defined in ggml-cuda/rocmfp2_mix.cu) — are missing from the "Full inventory (generated)" section, while the other new variables in this PR were added to it. Anyone searching the index finds no cross-reference for these kill switches. Add both to the inventory with their defining file to keep the documented table and the index consistent.</comment>
<file context>
@@ -44,6 +48,9 @@ consolidation of this list into CLI flags is tracked as follow-up work.
| `DFLASH_DS4_VERIFY_FORCE_GRAPH_REPLAY` | unset | OPT-IN: bypass graph property scans only after warmup; scheduler-generation checks remain mandatory. |
| `DFLASH_DS4_ROCTX` | unset | DEBUG: on HIP builds, dynamically load ROCTX and emit semantic DS4 prefill, speculative-decode, and layer-range markers for external rocprof traces. No events, timing, or device synchronization are added. |
| `DFLASH_QWEN35_ROCTX` | unset | DEBUG: on HIP builds, dynamically load ROCTX and mark Qwen concurrent steps, graph compute, and argmax readback with live, padded, and packed-prefill shape metadata. |
+| `DFLASH_GFX1151_HC_MMVF_Q4` | 1 on gfx1151 for the DS4 `[16384,24]` q4 projection | BURN-IN KILL SWITCH: =0 restores the generic hipBLAS dispatch decision. |
+| `GGML_CUDA_MLA_SPLIT_KV` / `GGML_DS4_FA_SPLIT_KV` | 1 on gfx1151 indexed decode; unset elsewhere | BURN-IN: force the reusable split-KV MLA schedule. Set `GGML_CUDA_MLA_NO_SPLIT_KV=1` (or legacy `GGML_DS4_FA_NO_SPLIT_KV=1`) to disable it. |
+| `GGML_DS4_TOPK_BLOCK_RADIX` | 1 on gfx1151 | BURN-IN KILL SWITCH: =0 restores hipCUB full sort for DS4-shaped 512-row top-k selection. |
</file context>
| backend, ncols, nrows, k, false, 2, iterations, full_sort) && | ||
| compute_topk( | ||
| backend, ncols, nrows, k, true, 2, iterations, hierarchical); | ||
| const bool exact = ran && same_selected_set( |
There was a problem hiding this comment.
P3: This standalone bench re-implements the long-context TOP_K selected-set parity check (full sort vs block-radix) and timing that already lives in the established framework, server/tests/test_deepseek4_unit.cpp (~lines 3020-3125), which registers in ctest and is what the PR cites as passing. Keeping a separate executable with its own correctness gate means the selected-set contract must be maintained in two places and can silently drift. Move the exact-set verification into the existing framework's long-context top-k test and keep this binary for perf-only measurement over the tile sweep.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/test/bench_ds4_topk.cpp, line 162:
<comment>This standalone bench re-implements the long-context TOP_K selected-set parity check (full sort vs block-radix) and timing that already lives in the established framework, server/tests/test_deepseek4_unit.cpp (~lines 3020-3125), which registers in ctest and is what the PR cites as passing. Keeping a separate executable with its own correctness gate means the selected-set contract must be maintained in two places and can silently drift. Move the exact-set verification into the existing framework's long-context top-k test and keep this binary for perf-only measurement over the tile sweep.</comment>
<file context>
@@ -0,0 +1,184 @@
+ backend, ncols, nrows, k, false, 2, iterations, full_sort) &&
+ compute_topk(
+ backend, ncols, nrows, k, true, 2, iterations, hierarchical);
+ const bool exact = ran && same_selected_set(
+ full_sort.indices, hierarchical.indices, ncols, nrows, k);
+ const double speedup = hierarchical.milliseconds > 0.0
</file context>
| // in must be a multiple of 128: the wide block load reads 128 weights at a time and would | ||
| // read past the tensor on the final block (register_host enforces this). | ||
| const int in = 256, out = 64, n_experts = 6, n_used = 3, ntok = 2; | ||
| // Three tokens exercises the low-register two-pass GLU finalizer. q <= 2 |
There was a problem hiding this comment.
P3: The new comment and the ntok=3 change assume ntok=3 always exercises the two-pass GLU finalizer, but that only holds on a gfx1151 strix-tuned device. In rocmfp2_mix.cu's ggml_cuda_rocmfp2_mix_mul_mat_id_glu, the one-pass kernel is selected when !strix_tuned || n_tokens <= 2 where strix_tuned = eu.gfx1151 && eg.gfx1151. On any non-gfx1151 HIP device, ntok=3 still takes the one-pass path, so the two-pass finalizer is never covered there and the new one_tok block duplicates the same path (both use the <1,2> kernel). The test only validates the two-pass branch when run on the gfx1151 target; consider gating this test (or the ntok=3 case) on the gfx1151 device so the coverage claim matches behavior.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/test/test_rocmfp_mix_gateup_glu.cpp, line 125:
<comment>The new comment and the ntok=3 change assume ntok=3 always exercises the two-pass GLU finalizer, but that only holds on a gfx1151 strix-tuned device. In rocmfp2_mix.cu's ggml_cuda_rocmfp2_mix_mul_mat_id_glu, the one-pass kernel is selected when `!strix_tuned || n_tokens <= 2` where `strix_tuned = eu.gfx1151 && eg.gfx1151`. On any non-gfx1151 HIP device, ntok=3 still takes the one-pass path, so the two-pass finalizer is never covered there and the new one_tok block duplicates the same path (both use the `<1,2>` kernel). The test only validates the two-pass branch when run on the gfx1151 target; consider gating this test (or the ntok=3 case) on the gfx1151 device so the coverage claim matches behavior.</comment>
<file context>
@@ -122,7 +122,9 @@ TEST_CASE(RocmfpMixGateupGluFixture, fused_gateup_glu) {
// in must be a multiple of 128: the wide block load reads 128 weights at a time and would
// read past the tensor on the final block (register_host enforces this).
- const int in = 256, out = 64, n_experts = 6, n_used = 3, ntok = 2;
+ // Three tokens exercises the low-register two-pass GLU finalizer. q <= 2
+ // uses the one-pass kernel and is checked separately below.
+ const int in = 256, out = 64, n_experts = 6, n_used = 3, ntok = 3;
</file context>
|
|
||
| ggml_cuda_rocmfp2_mix_unregister(d_gate); | ||
| ggml_cuda_rocmfp2_mix_unregister(d_up); | ||
| // Hash both raw projections and the fused output after timing. This catches |
There was a problem hiding this comment.
P3: The new verification block claims it "catches mode-specific numerical drift," but it never compares any of the hashes against a reference or against each other — it only prints them. It also does not check the actual fusion invariant: fused_result (swiGLU output) is never compared to anything recomputed from the raw up/gate projections, so a wrong fusion still passes silently. For the claims in the comment to hold, compute silu(up_result)*gate_result on host and compare (with tolerance) to fused_result, or assert the raw/fused hashes match an expected value; otherwise reword the comment. This matters because the PR relies on these benchmarks to gate the fusion.
Prompt for AI agents
Check if this issue is valid — if so, understand the root cause and fix it. At server/test/bench_rocmfp_mix_gateup_glu.cpp, line 233:
<comment>The new verification block claims it "catches mode-specific numerical drift," but it never compares any of the hashes against a reference or against each other — it only prints them. It also does not check the actual fusion invariant: fused_result (swiGLU output) is never compared to anything recomputed from the raw up/gate projections, so a wrong fusion still passes silently. For the claims in the comment to hold, compute silu(up_result)*gate_result on host and compare (with tolerance) to fused_result, or assert the raw/fused hashes match an expected value; otherwise reword the comment. This matters because the PR relies on these benchmarks to gate the fusion.</comment>
<file context>
@@ -170,20 +214,51 @@ int main() {
- ggml_cuda_rocmfp2_mix_unregister(d_gate);
- ggml_cuda_rocmfp2_mix_unregister(d_up);
+ // Hash both raw projections and the fused output after timing. This catches
+ // mode-specific numerical drift even when the generated-token hash happens
+ // to remain unchanged.
</file context>
Corrected Lucebox6 result: adaptive is workload-dependent; the full performance commit does not speed up the private q5 bundleCorrection to my earlier conclusion: I had generalized from one high-acceptance workload and initially tested the adaptive controller in isolation, not the complete PR 667 performance stack. Adaptive q5 controllerThe width traces confirm that adaptive mode was active (fixed offered 4 candidates/step; adaptive averaged roughly 3.3–3.8 and starts each request at q3). On natural-EOS HumanEval, adaptive was slightly faster by the server metric:
On the capped, high-acceptance coding workload, adaptive was -7.22% by server weighted decode. Fixed q5 amortizes more work when acceptance is already high. Therefore the adaptive result is not universally negative, but it is not a robust default win either; keeping it opt-in is justified. Full PR 667 performance commit on the PR 58 private bundleI then rebuilt the exact PR 58 package on private base
All 40 measured requests ended naturally ( Correctness tests passed: server 432/432, DeepSeek4 unit suite, GPU top-k parity, and gfx1151 ROCmFP. The block-radix top-k microbenchmark is genuinely faster on gfx1151 (2.33–4.31x across tested 8K–32K widths), but that kernel win does not translate to this private q5 end-to-end path. Most of this commit is gated to gfx1151/q4 or paths bypassed by the bundled private operator. Recommendation: do not stack Reproducibility artifacts retained on lucebox6:
|
|
Thank you, @davide221. We independently qualified the current PR #667 head, Test configuration
The 2,048-token workload used one warmup followed by three measured 510-token generations. Each retrieval point used one deterministic five-key request with a 256-token output ceiling. Matched results
Both commits recovered all five values exactly at 7,680, 30,720, and 122,879 input tokens. The result confirms a decode-side improvement while input-processing throughput remains unchanged within 0.3%. Quality and compatibilityThe candidate passed:
The unmodified harness labels The complete 131,072-token configuration fit in 128 GiB. Minimum effective non-CMA memory headroom was 4,374,508 KiB, with no OOM, GPU reset, service restart, or safety-watcher event. Reproduction
The public Ansible release now builds the unmodified Thanks again for implementing and iterating on this optimization. |
2 participants
Summary
Merge-ready, stacked on #664 (
codex/ds4-pflash-long-context). This PR accelerates DeepSeek V4 Flash verification on gfx1151 without reducing the model's six routed experts. Here,q=4means seed-inclusive speculative verification width four, not four experts.Cleaned scope
The draft-only fused MoE combine and direct-visible attention experiments were removed: neither cleared the qualification threshold, and the latter's shape contract was not broad enough for merge.
Real-workflow A/B
Matched on lucebox8's Strix Halo (gfx1151), ROCm 7.2.4, the same target and DSpark GGUFs, q4 verification, six routed experts, q4 KV, sparse prefill, and 32K server context. Sparse verifier attention was explicitly enabled on the exact base (
5452b4a) to match the head (5eb4fbe) default.The two 24K agent cases averaged 22.42 -> 29.71 tok/s (+32.5%). Prefill was essentially unchanged: HumanEval 43.5 -> 44.7 tok/s, GSM8K 31.3 -> 32.5 tok/s, and agent 94.1 -> 93.9 tok/s. This PR closes a substantial decode gap; it does not claim to close the remaining raw-prefill gap to Vulkan.
On matched requests, verifier time fell from roughly 125-133 ms/step to 84-101 ms/step. Benchmark artifacts and logs are retained under
/data/pr667-merge-ready/on lucebox8.The sparse/quantized production path is not an end-to-end byte-identity mode, so task correctness is the workflow quality gate. Kernel-level gates remain strict where applicable: the new top-k returns the exact reference selected set, and ROCmFP2/3 tests validate every qualified width against the host reference.
Adaptive width result
Adaptive width remains opt-in via
DFLASH_ADAPTIVE_SPEC_WIDTH=1or the DS4-specificDFLASH_DS4_ADAPTIVE_WIDTH=1. On the same real workflows it was slower than fixed q4: HumanEval -4.7%, GSM8K -15.9%, and agent -3.5%. The implementation is retained for future backend-specific qualification, while all non-DS4 backends preserve their previous default policies.Verification
dflash_serverbuild on ROCm/gfx1151.test_server_unit: 431/431 passed after final cleanup.deepseek4_unit, including long-context top-k and indexed-attention GPU coverage: passed.test_rocmfp_mix_gateup_gluandtest_rocmfpx_mmq: all parity cases passed.git diff --check: clean.An unrelated all-target HIP build still hits the pre-existing
cudaErrorNoDevicecompatibility issue in unchangedtest_rms_norm_hip.cpp; the production server and all PR-focused targets build cleanly.Controls
DFLASH_DS4_SPARSE_DECODE_FLASH=0: restore explicit verifier attention.GGML_CUDA_MLA_NO_SPLIT_KV=1: disable split-KV MLA.GGML_DS4_TOPK_BLOCK_RADIX=0: restore hipCUB full sort.DFLASH_ROCMFP2_ROW4=0: restore the two-row ROCmFP2 kernel.DFLASH_GFX1151_HC_MMVF_Q4=0: restore generic q4 projection dispatch.DFLASH_DS4_INCREMENTAL_VERIFY_MASK=0: restore full verifier-mask transfer.