perf(ds4): accelerate ROCmFPx dual-GPU serving

[marcelormendes]

Summary

• add build-wide support for affine qtype-107 ROCmFP2 weights, including MMQ/MMVQ kernels and dispatch-asserting coverage on gfx1151 and gfx12xx; gfx1100 retains the non-MMQ fallback
• 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 and whether speculative decode actually ran 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 gfx1151; gfx1100 uses the supported non-MMQ fallback and is not claimed as an MMQ qualification

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
• CI now retains the default ROCmFP2 build and adds a separate affine-enabled build on gfx1151 and gfx1201; the focused test asserts actual MMVQ/MMQ dispatch
• the controlled decode benchmark rejects missing or false spec_decode_ran; its benchmark-tool suite passes 16 tests
• post-fix CI run 32150269678 passed all seven jobs, including Windows and both affine-enabled AMD runs
• 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>

[cubic-dev-ai]

1 issue found across 10 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/test/test_rocmfpx_mmq.cpp">

<violation number="1" location="server/test/test_rocmfpx_mmq.cpp:273">
P2: On gfx1100, this test now skips instead of validating the ROCmFPx MMVQ/MMQ paths. Restore `gfx1100` to the accepted-architecture check so the advertised gfx1100 parity coverage remains enforced.</violation>
</file>

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

Re-trigger cubic</sub>

[cubic-dev-ai]

P2: On gfx1100, this test now skips instead of validating the ROCmFPx MMVQ/MMQ paths. Restore gfx1100 to the accepted-architecture check so the advertised gfx1100 parity coverage remains enforced.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/test/test_rocmfpx_mmq.cpp, line 273:

<comment>On gfx1100, this test now skips instead of validating the ROCmFPx MMVQ/MMQ paths. Restore `gfx1100` to the accepted-architecture check so the advertised gfx1100 parity coverage remains enforced.</comment>

<file context>
@@ -223,14 +270,12 @@ int main() {
     if (std::strncmp(properties.gcnArchName, "gfx1151", 7) != 0 &&
-        std::strncmp(properties.gcnArchName, "gfx1100", 7) != 0) {
-        std::printf("SKIP: ROCmFPX MMQ test expects gfx1100/gfx1151 (found %s)\n",
+        std::strncmp(properties.gcnArchName, "gfx12", 5) != 0) {
+        std::printf("SKIP: ROCmFPX dispatch test expects gfx1151/gfx12xx (found %s)\n",
                     properties.gcnArchName);
</file context>

Suggested change

	std::strncmp(properties.gcnArchName, "gfx12", 5) != 0) {
	std::strncmp(properties.gcnArchName, "gfx1100", 7) != 0 &&
	std::strncmp(properties.gcnArchName, "gfx12", 5) != 0) {

[davide221]

Intentional. Production ROCmFPX MMQ dispatch excludes RDNA 3.0 (gfx1100) and only enables RDNA 3.5/4, so accepting gfx1100 here would recreate the false-positive test this commit fixes: numerical output could pass through the non-MMQ fallback. The PR body and DS4 documentation now explicitly limit MMQ qualification to gfx1151/gfx12xx; both CI runners compile the affine layout and the test asserts the actual MMVQ/MMQ launch counters.