perf(ds4): accelerate ROCmFPx dual-GPU serving

.github/workflows/ci.yml

@@ -357,6 +357,28 @@ jobs:
             --output-on-failure \
             -R 'rocmfp4_reference|rocmfpx_reference|rocmfp4_hip_tail|rocmfpx_mmq|deepseek4_mmid_grouped_cuda|deepseek4_unit|recurrent_snapshot|ChainRollbackPolicy|rocmfp3_mix_registry|rocmfp_mix_slice_matvec|rocmfp_mix_gateup_glu|ds4_mix_registry_teardown|test_model_smoke\.PagedAttention\.|batched_gdn$|concat_transpose$'
 
+      - name: Build + test affine ROCmFP2 GPU paths
+        # Keep the default wire-format build above and compile the affine layout
+        # separately so both configurations remain covered. The focused tests
+        # assert real MMVQ/MMQ dispatch rather than accepting a BLAS fallback.
+        run: |
+          cmake -S server -B "$RUNNER_TEMP/rocmfp-affine-build" \
+            -DDFLASH27B_GPU_BACKEND=hip \
+            -DDFLASH27B_HIP_ARCHITECTURES="$EXPECTED_HIP_ARCH" \
+            -DDFLASH27B_SERVER=OFF \
+            -DDFLASH27B_TESTS=ON \
+            -DDFLASH27B_ROCMFP2_AFFINE=ON \
+            -DGGML_HIP_GRAPHS=ON \
+            -DCMAKE_BUILD_TYPE=Release \
+            -DCMAKE_HIP_FLAGS=-DDFLASH_WAVE_SIZE=32
+          cmake --build "$RUNNER_TEMP/rocmfp-affine-build" \
+            --target test_rocmfpx test_rocmfpx_mmq \
+              test_deepseek4_mmid_grouped_cuda \
+            --parallel 8
+          ctest --test-dir "$RUNNER_TEMP/rocmfp-affine-build" \
+            --output-on-failure \
+            -R 'rocmfpx_reference|rocmfpx_mmq|deepseek4_mmid_grouped_cuda'
+
   build-windows:
     name: Build Windows (MSVC + CUDA, library + server targets)
     runs-on: windows-2022

harness/tests/test_ds4_benchmark_tools.py

@@ -14,11 +14,14 @@
 REPO_ROOT = Path(__file__).resolve().parents[2]
 BENCHMARKS_DIR = REPO_ROOT / "harness" / "benchmarks" / "deepseek4"
 QUALIFICATION_DIR = REPO_ROOT / "harness" / "qualification" / "deepseek4"
+SERVER_SCRIPTS_DIR = REPO_ROOT / "server" / "scripts"
 QUALIFIER = QUALIFICATION_DIR / "qualify_ds4_q5_amd.sh"
 sys.path.insert(0, str(BENCHMARKS_DIR))
 sys.path.insert(0, str(QUALIFICATION_DIR))
+sys.path.insert(0, str(SERVER_SCRIPTS_DIR))
 
 import analyze_rocprof_overlap  # noqa: E402
+import bench_ds4_decode  # noqa: E402
 import ds4_context_sweep  # noqa: E402
 import ds4_publication_decode_client  # noqa: E402
 
@@ -128,6 +131,39 @@ def test_zero_runs_is_rejected(self) -> None:
         self.assertEqual(error.exception.code, 2)
 
 
+class ControlledDecodeBenchmarkTests(unittest.TestCase):
+    @staticmethod
+    def valid_run(**overrides: object) -> dict:
+        run = {
+            "completion_tokens": 512,
+            "output_matches_prompt": True,
+            "cache_hit": False,
+            "cached_prefix_tokens": 0,
+            "decode_tokens_per_second": 47.0,
+            "spec_decode_ran": True,
+            "accept_rate": 1.0,
+        }
+        run.update(overrides)
+        return run
+
+    def test_speculative_run_is_accepted(self) -> None:
+        bench_ds4_decode.validate_run(self.valid_run(), max_tokens=512)
+
+    def test_autoregressive_fallback_is_rejected(self) -> None:
+        with self.assertRaisesRegex(RuntimeError, "speculative decode did not run"):
+            bench_ds4_decode.validate_run(
+                self.valid_run(spec_decode_ran=False, accept_rate=0.0),
+                max_tokens=512,
+            )
+
+    def test_missing_speculative_signal_is_rejected(self) -> None:
+        with self.assertRaisesRegex(RuntimeError, "speculative decode did not run"):
+            bench_ds4_decode.validate_run(
+                self.valid_run(spec_decode_ran=None),
+                max_tokens=512,
+            )
+
+
 class QualifierPreflightTests(unittest.TestCase):
     def run_qualifier(self, **overrides: str) -> subprocess.CompletedProcess[str]:
         fixture = str(Path(__file__).resolve())

server/CMakeLists.txt

@@ -177,6 +177,16 @@ endif()
 option(DFLASH27B_FA_ALL_QUANTS "Compile ggml-cuda fattn kernels for all KV-quant pairs" ON)
 set(GGML_CUDA_FA_ALL_QUANTS ${DFLASH27B_FA_ALL_QUANTS} CACHE BOOL "" FORCE)
 
+# Some third-party qtype-107 GGUFs encode one affine (scale, offset) pair per
+# 32-weight block instead of the original two-scale layout. This is a wire
+# format choice, so keep it explicit and build-wide.
+option(DFLASH27B_ROCMFP2_AFFINE
+    "Interpret qtype-107 ROCmFP2 blocks as affine code*scale-offset weights"
+    OFF)
+if(DFLASH27B_ROCMFP2_AFFINE)
+    add_compile_definitions(ROCMFP2_AFFINE=1)
+endif()
+
 # HIP Phase 2 — opt-in rocWMMA flashprefill kernels (Strix Halo / gfx1151).
 # Default OFF (Phase 1 = ggml q8 fallback). Set ON for the 1.7-2.7× compress
 # speedup at 8K-32K context. Requires rocwmma headers installed under
@@ -357,6 +367,10 @@ if(DFLASH27B_GPU_BACKEND STREQUAL "hip" OR
     )
     target_include_directories(ggml-hip BEFORE PRIVATE
         ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_compat)
+    # HIP's clang driver otherwise finds Ubuntu's older libamdhip64 before the
+    # configured ROCm runtime when it expands --hip-link.
+    get_filename_component(_dflash_hip_runtime_dir "${hip_DIR}/../.." ABSOLUTE)
+    add_link_options("-L${_dflash_hip_runtime_dir}")
 endif()
 
 if(DFLASH27B_ENABLE_MIXED_CUDA_HIP)
@@ -989,6 +1003,11 @@ if(DFLASH27B_TESTS)
         else()
             target_link_libraries(test_deepseek4_mmid_grouped_cuda PRIVATE CUDA::cudart)
         endif()
+        # The grouped test includes rocmfpx.h unconditionally; expose it to both
+        # backends (previously only HIP resolved it, so the default CUDA build
+        # could not compile this target).
+        target_include_directories(test_deepseek4_mmid_grouped_cuda PRIVATE
+            ${CMAKE_CURRENT_SOURCE_DIR}/deps/llama.cpp/ggml/rocmfpx)
         target_link_libraries(test_deepseek4_mmid_grouped_cuda PRIVATE
             ggml ${DFLASH27B_GGML_BACKEND_TARGET})
         list(APPEND _raw_unit_test_targets test_deepseek4_mmid_grouped_cuda)

server/deps/llama.cpp/ggml/include/ggml-cuda.h

@@ -58,6 +58,12 @@ GGML_BACKEND_API bool ggml_backend_cuda_set_graphs_disabled_override(bool disabl
 // Intended for focused correctness tests of the dispatch guard.
 GGML_BACKEND_API size_t ggml_backend_cuda_get_concat_transpose_f32_count(void);
 
+// Calling-thread launch counters for quantized matrix-vector (MMVQ) and
+// matrix-matrix (MMQ) kernels. Intended for focused tests that must prove
+// which dispatch path executed rather than only checking numerical output.
+GGML_BACKEND_API size_t ggml_backend_cuda_get_mmvq_launch_count(void);
+GGML_BACKEND_API size_t ggml_backend_cuda_get_mmq_launch_count(void);
+
 // device buffer
 GGML_BACKEND_API ggml_backend_buffer_type_t ggml_backend_cuda_buffer_type(int device);
 

server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c

@@ -716,8 +716,137 @@ static uint8_t rocmfpx_choose_scale_fp8_weighted_mse(const float * x, int n, con
     return best_e;
 }
 
+#ifdef ROCMFP2_AFFINE
+static uint8_t rocmfpx_quantize_fp2_affine_code(
+        float x, float scale, float offset) {
+    if (!(scale > 0.0f)) {
+        return 0;
+    }
+    if (!isfinite(x)) {
+        x = 0.0f;
+    }
+    const float q = floorf((x + offset) / scale + 0.5f);
+    return (uint8_t) fminf(3.0f, fmaxf(0.0f, q));
+}
+
+static float rocmfpx_fp2_affine_block_mse(
+        const float * x, int n, const float * mse_weights,
+        uint8_t scale_e, uint8_t offset_e, float best_err) {
+    const float scale = rocmfpx_ue4m3_to_fp32(scale_e);
+    const float offset = rocmfpx_ue4m3_to_fp32(offset_e);
+    float err = 0.0f;
+    for (int i = 0; i < n; ++i) {
+        if (!isfinite(x[i])) continue;
+        const float weight = mse_weights ? mse_weights[i] : 1.0f;
+        if (!(weight > 0.0f)) continue;
+        const uint8_t code =
+            rocmfpx_quantize_fp2_affine_code(x[i], scale, offset);
+        const float delta = x[i] - ((float) code * scale - offset);
+        err += weight * delta * delta;
+        if (err > best_err) return err;
+    }
+    return err;
+}
+
+static void rocmfpx_choose_fp2_affine_params(
+        const float * x, int n, const float * mse_weights,
+        uint8_t * scale_e, uint8_t * offset_e) {
+    float min_value = INFINITY;
+    float max_value = -INFINITY;
+    for (int pass = 0; pass < 2 && min_value == INFINITY; ++pass) {
+        for (int i = 0; i < n; ++i) {
+            if (!isfinite(x[i])) continue;
+            if (pass == 0 && mse_weights && !(mse_weights[i] > 0.0f)) {
+                continue;
+            }
+            min_value = fminf(min_value, x[i]);
+            max_value = fmaxf(max_value, x[i]);
+        }
+        if (!mse_weights) break;
+    }
+    if (min_value == INFINITY ||
+        (min_value == 0.0f && max_value == 0.0f)) {
+        *scale_e = 0;
+        *offset_e = 0;
+        return;
+    }
+
+    const float initial_offset = fmaxf(0.0f, -min_value);
+    const float initial_scale =
+        fmaxf((max_value + initial_offset) / 3.0f,
+              rocmfpx_ue4m3_to_fp32(1));
+    uint8_t best_scale = rocmfpx_nearest_scale_ue4m3(initial_scale);
+    uint8_t best_offset = rocmfpx_nearest_scale_ue4m3(initial_offset);
+    float best_err = rocmfpx_fp2_affine_block_mse(
+        x, n, mse_weights, best_scale, best_offset, INFINITY);
+
+    for (int e = 1; e <= 0x7e; ++e) {
+        const float err = rocmfpx_fp2_affine_block_mse(
+            x, n, mse_weights, (uint8_t) e, best_offset, best_err);
+        if (err < best_err || (err == best_err && e < best_scale)) {
+            best_err = err;
+            best_scale = (uint8_t) e;
+        }
+    }
+    for (int e = 0; e <= 0x7e; ++e) {
+        const float err = rocmfpx_fp2_affine_block_mse(
+            x, n, mse_weights, best_scale, (uint8_t) e, best_err);
+        if (err < best_err || (err == best_err && e < best_offset)) {
+            best_err = err;
+            best_offset = (uint8_t) e;
+        }
+    }
+    for (int e = 1; e <= 0x7e; ++e) {
+        const float err = rocmfpx_fp2_affine_block_mse(
+            x, n, mse_weights, (uint8_t) e, best_offset, best_err);
+        if (err < best_err || (err == best_err && e < best_scale)) {
+            best_err = err;
+            best_scale = (uint8_t) e;
+        }
+    }
+    *scale_e = best_scale;
+    *offset_e = best_offset;
+}
+
+static void rocmfpx_quantize_row_fp2_affine(
+        const float * GGML_RESTRICT x,
+        block_rocmfp2 * GGML_RESTRICT y, int64_t k,
+        const float * GGML_RESTRICT quant_weights, float sigma2) {
+    const int64_t nb = k / QK_ROCMFP2;
+    for (int64_t ib = 0; ib < nb; ++ib) {
+        const float * xb = x + ib*QK_ROCMFP2;
+        block_rocmfp2 * yb = y + ib;
+        float mse_weights[QK_ROCMFP2];
+        const float * block_weights = NULL;
+        if (quant_weights) {
+            float max_abs;
+            float max_abs_weight;
+            rocmfpx_prepare_mse_weights(
+                mse_weights, xb, QK_ROCMFP2,
+                quant_weights + ib*QK_ROCMFP2, sigma2,
+                &max_abs, &max_abs_weight);
+            block_weights = mse_weights;
+        }
+        rocmfpx_choose_fp2_affine_params(
+            xb, QK_ROCMFP2, block_weights, &yb->e[0], &yb->e[1]);
+        const float scale = rocmfpx_ue4m3_to_fp32(yb->e[0]);
+        const float offset = rocmfpx_ue4m3_to_fp32(yb->e[1]);
+        memset(yb->qs, 0, sizeof(yb->qs));
+        for (int i = 0; i < QK_ROCMFP2; ++i) {
+            const uint8_t code =
+                rocmfpx_quantize_fp2_affine_code(xb[i], scale, offset);
+            yb->qs[i >> 2] |= (uint8_t) (code << (2*(i & 3)));
+        }
+    }
+}
+#endif
+
 void rocmfpx_quantize_row_fp2_ref(const float * GGML_RESTRICT x, block_rocmfp2 * GGML_RESTRICT y, int64_t k) {
     assert(k % QK_ROCMFP2 == 0);
+#ifdef ROCMFP2_AFFINE
+    rocmfpx_quantize_row_fp2_affine(x, y, k, NULL, 0.0f);
+    return;
+#endif
 
     const int64_t nb = k / QK_ROCMFP2;
     for (int64_t ib = 0; ib < nb; ++ib) {
@@ -745,6 +874,11 @@ void rocmfpx_quantize_row_fp2_ref(const float * GGML_RESTRICT x, block_rocmfp2 *
 static void rocmfpx_quantize_row_fp2_weighted(
         const float * GGML_RESTRICT x, block_rocmfp2 * GGML_RESTRICT y, int64_t k, const float * GGML_RESTRICT quant_weights) {
     assert(k % QK_ROCMFP2 == 0);
+#ifdef ROCMFP2_AFFINE
+    rocmfpx_quantize_row_fp2_affine(
+        x, y, k, quant_weights, rocmfpx_row_sigma2(x, k));
+    return;
+#endif
 
     const float sigma2 = rocmfpx_row_sigma2(x, k);
 
@@ -785,9 +919,17 @@ void rocmfpx_dequantize_row_fp2(const block_rocmfp2 * GGML_RESTRICT x, float * G
         float * yb = y + ib*QK_ROCMFP2;
 
         for (int i = 0; i < QK_ROCMFP2; ++i) {
+#ifdef ROCMFP2_AFFINE
+            // affine type-107: value = code*scale - offset, e[0]=scale, e[1]=offset
+            const float scale  = rocmfpx_ue4m3_to_fp32(xb->e[0]);
+            const float offset = rocmfpx_ue4m3_to_fp32(xb->e[1]);
+            const uint8_t code = (uint8_t) ((xb->qs[i >> 2] >> (2*(i & 3))) & 3u);
+            yb[i] = (float) code * scale - offset;
+#else
             const float scale = rocmfpx_ue4m3_to_fp32(xb->e[i >= QK_ROCMFP2/2]);
             const uint8_t code = (uint8_t) ((xb->qs[i >> 2] >> (2*(i & 3))) & 3u);
             yb[i] = kvalues_rocmfp2[code] * scale;
+#endif
         }
     }
 }

server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c

@@ -103,6 +103,47 @@ static void check_weighted_imatrix_fp2(void) {
     assert(weighted_err < plain_err);
 }
 
+#ifdef ROCMFP2_AFFINE
+static void check_fp2_affine_encoding(void) {
+    float src[QK_ROCMFP2];
+    block_rocmfp2 quantized;
+
+    for (int i = 0; i < QK_ROCMFP2; ++i) {
+        src[i] = -2.0f + 4.0f*(float) i/(float) (QK_ROCMFP2 - 1);
+    }
+    rocmfpx_quantize_row_fp2_ref(src, &quantized, QK_ROCMFP2);
+
+    const float scale = rocmfpx_ue4m3_to_fp32(quantized.e[0]);
+    const float offset = rocmfpx_ue4m3_to_fp32(quantized.e[1]);
+    assert(scale > 0.0f);
+    assert(offset > 0.0f);
+    for (int i = 0; i < QK_ROCMFP2; ++i) {
+        const float q = floorf((src[i] + offset)/scale + 0.5f);
+        const uint8_t expected =
+            (uint8_t) fminf(3.0f, fmaxf(0.0f, q));
+        const uint8_t actual =
+            (quantized.qs[i >> 2] >> (2*(i & 3))) & 0x3;
+        assert(actual == expected);
+    }
+
+    src[0] = NAN;
+    src[1] = INFINITY;
+    src[2] = -INFINITY;
+    rocmfpx_quantize_row_fp2_ref(src, &quantized, QK_ROCMFP2);
+
+    const float nonfinite_scale = rocmfpx_ue4m3_to_fp32(quantized.e[0]);
+    const float nonfinite_offset = rocmfpx_ue4m3_to_fp32(quantized.e[1]);
+    const float zero_q = floorf(nonfinite_offset/nonfinite_scale + 0.5f);
+    const uint8_t zero_code =
+        (uint8_t) fminf(3.0f, fmaxf(0.0f, zero_q));
+    for (int i = 0; i < 3; ++i) {
+        const uint8_t actual =
+            (quantized.qs[i >> 2] >> (2*(i & 3))) & 0x3;
+        assert(actual == zero_code);
+    }
+}
+#endif
+
 static void check_large_finite_values(void) {
     float src[QK_ROCMFPX] = { 0 };
     float imatrix[QK_ROCMFPX];
@@ -189,6 +230,9 @@ int main(void) {
     assert(mse3 < mse2);
 
     check_weighted_imatrix_fp2();
+#ifdef ROCMFP2_AFFINE
+    check_fp2_affine_encoding();
+#endif
     check_weighted_imatrix_fp3();
     check_large_finite_values();
 

server/deps/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh

@@ -168,11 +168,23 @@ static __device__ __forceinline__ void dequantize_rocmfpx_fp2(const void * vx, c
 
     const int i0 = iqs + 0;
     const int i1 = iqs + 1;
+
+#ifdef ROCMFP2_AFFINE
+    // Affine qtype-107 stores one (scale, offset) pair per 32-weight block:
+    // value = code*scale - offset, with code in [0, 3]. This wire format is
+    // incompatible with the default two-scale ROCmFP2 layout.
+    const float scale  = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[0]);
+    const float offset = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[1]);
+
+    v.x = (float) (rocmfpx_get_fp2_code_cuda(x[ib].qs, i0) & 3u) * scale - offset;
+    v.y = (float) (rocmfpx_get_fp2_code_cuda(x[ib].qs, i1) & 3u) * scale - offset;
+#else
     const float d0 = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[i0 >= QK_ROCMFP2/2]);
     const float d1 = rocmfpx_ue4m3_to_fp32_finite(x[ib].e[i1 >= QK_ROCMFP2/2]);
 
     v.x = d0 * rocmfpx_decode_fp2_code_cuda(rocmfpx_get_fp2_code_cuda(x[ib].qs, i0));
     v.y = d1 * rocmfpx_decode_fp2_code_cuda(rocmfpx_get_fp2_code_cuda(x[ib].qs, i1));
+#endif
 }
 
 static __device__ __forceinline__ void dequantize_rocmfpx_fp6(const void * vx, const int64_t ib, const int iqs, float2 & v) {