perf(ds4): accelerate ROCmFPx dual-GPU serving

[marcelormendes]

Summary

• add build-wide support for affine qtype-107 ROCmFP2 weights, including MMQ/MMVQ kernels and parity coverage on gfx1100 and gfx1151
• add grouped expert dispatch, phase-aware critical-path placement, device-side joins, and reusable scratch release for heterogeneous DeepSeek4 serving
• add bounded long-context prefill scheduling, a qualified 128 KiB RX 7900 XT + Strix Halo launch profile, and a reproducible decode benchmark harness
• report request timings through the OpenAI-compatible API and treat non-positive max_completion_tokens as the server default

Qualified results

Hardware: Radeon RX 7900 XT + Ryzen AI Max+ 395, true top-k-6, fixed DSpark q=4, 135,168 context slots, no prefix/prefill cache.

• controlled 512-token decode: 45.0, 47.0, and 47.7 tok/s; 100% speculative acceptance; byte-identical output
• true top-k-6 sparse prefill: 111.2 tok/s at 132,981 tokens
• affine MMQ parity remained finite on both gfx1100 and gfx1151

Older top-k-4 sparse-prefill records are documented separately and are not presented as equivalent launch settings. Adaptive qtype-105/106 expert tensors remain monolithic-only; the dual-GPU profile targets uniform or fixed-codebook tensors.

Verification

• built dflash_server, test_rocmfpx_mmq, test_server_unit, test_deepseek4_unit, and test_deepseek4_mmid_grouped_cuda
• ctest --test-dir server/build-rocmfpx-pr --output-on-failure -j8 — 377/377 passed
• test_rocmfpx_mmq and test_deepseek4_mmid_grouped_cuda passed on HIP_VISIBLE_DEVICES=0 and 1; grouped parity covered 70 compared cases and 2,236,416 compared bytes per device
• benchmark Python and launch shell scripts passed syntax checks
• git diff --check HEAD^ passed

[cubic-dev-ai]

2 issues found across 29 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/mmq.cu">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu:513">
P2: When a split tensor spans the qualified device and another GPU, this owner filter disables MMQ for the entire operation because callers AND all device results into one `use_mul_mat_q` flag. Apply the qualification at per-device launch/placement rather than this global capability predicate, or explicitly handle split tensors so the matching owner can still use MMQ.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh:534">
P3: The affine offset-correction math (`sumq` extraction and the `scale*sumi + (scale-offset)*sumq` formula) is copied verbatim across `vecdotq.cuh` and `mmvq.cu` (and re-derived in the MMQ loader). A future change to the affine wire format must be mirrored in each copy or the kernels silently disagree. Factor the sumq/offset formula into a shared helper to keep the MMVQ paths consistent.</violation>
</file>

<sub>Reply with feedback, questions, or to request a fix.

Re-trigger cubic</sub>

[cubic-dev-ai]

P3: The affine offset-correction math (sumq extraction and the scale*sumi + (scale-offset)*sumq formula) is copied verbatim across vecdotq.cuh and mmvq.cu (and re-derived in the MMQ loader). A future change to the affine wire format must be mirrored in each copy or the kernels silently disagree. Factor the sumq/offset formula into a shared helper to keep the MMVQ paths 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/src/ggml-cuda/vecdotq.cuh, line 534:

<comment>The affine offset-correction math (`sumq` extraction and the `scale*sumi + (scale-offset)*sumq` formula) is copied verbatim across `vecdotq.cuh` and `mmvq.cu` (and re-derived in the MMQ loader). A future change to the affine wire format must be mirrored in each copy or the kernels silently disagree. Factor the sumq/offset formula into a shared helper to keep the MMVQ paths consistent.</comment>

<file context>
@@ -525,7 +525,23 @@ static __device__ __forceinline__ float vec_dot_rocmfpx_fp2_q8_1(
+    // affine dot is scale*sum(c*q) - offset*sum(q) = scale*sumi + (scale-offset)*sumq.
+    const float scale  = rocmfpx_ue4m3_to_fp32_finite(bq2->e[0]);
+    const float offset = rocmfpx_ue4m3_to_fp32_finite(bq2->e[1]);
+    int sumq = 0;
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
</file context>

[marcelormendes]

Valid. Fixed in 96fbda3: the affine sumq extraction and correction formula now live in one device helper shared by generic and packed MMVQ. The MMQ path performs per-lane affine value decoding rather than duplicating this dot-reduction formula.

[marcelormendes]

Post-cleanup live verification at commit 2d42034:

• rebuilt and restarted the checked-in serve_ds4_dual_rocm_128k.sh profile
• one warm-up: 44.1 tok/s
• three measured 512-token runs: 47.5, 46.7, 44.8 tok/s
• every measured run completed all 512 tokens with 100% speculative acceptance
• prefix cache was disabled (cache_hit=false, cached_prefix_tokens=0)
• all outputs were byte-identical: d9b02155239dee374a8564f46bbb4f7e72a638ace42d1db415b1944a8bddaffd

This confirms the /simplify cleanup preserved 40+ tok/s decode behavior in the published commit.

[cubic-dev-ai]

All reported issues were addressed across 15 files (changes from recent commits).

<sub>Tip: Review your code locally with the cubic CLI to iterate faster.

Re-trigger cubic</sub>