diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
index f6a4db82a..32187e44f 100644
--- a/.github/workflows/ci.yml
+++ b/.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
diff --git a/harness/tests/test_ds4_benchmark_tools.py b/harness/tests/test_ds4_benchmark_tools.py
index 56e5aff03..e6a9a93c4 100644
--- a/harness/tests/test_ds4_benchmark_tools.py
+++ b/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())
diff --git a/server/CMakeLists.txt b/server/CMakeLists.txt
index 46ecfda81..b0692283f 100644
--- a/server/CMakeLists.txt
+++ b/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)
diff --git a/server/deps/llama.cpp/ggml/include/ggml-cuda.h b/server/deps/llama.cpp/ggml/include/ggml-cuda.h
index bb850e2a9..0662f0f1b 100644
--- a/server/deps/llama.cpp/ggml/include/ggml-cuda.h
+++ b/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);
 
diff --git a/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c b/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c
index da798d23a..ec2690327 100644
--- a/server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c
+++ b/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
         }
     }
 }
diff --git a/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c b/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c
index a3c49ffb3..08048dca1 100644
--- a/server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c
+++ b/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();
 
diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh
index 3ed52e12e..5c092c7c3 100644
--- a/server/deps/llama.cpp/ggml/src/ggml-cuda/dequantize.cuh
+++ b/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) {
diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
index 1c55d51d1..6e6f6b458 100644
--- a/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -2582,6 +2582,14 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) {
     bool use_mul_mat_vec_q = ggml_is_quantized(src0->type) && !bad_padding_clear && src1->type == GGML_TYPE_F32 &&
                              dst->type == GGML_TYPE_F32 &&
                              ncols_dst <= (is_mul_mat_id ? MMVQ_MAX_MOE_BATCH_SIZE : MMVQ_MAX_BATCH_SIZE);
+#ifdef ROCMFP2_AFFINE
+    // The affine MMVQ dot includes the offset correction in vecdotq.cuh.
+    // Keep the conservative dequantize fallback unless explicitly enabled.
+    if (src0->type == GGML_TYPE_Q2_0_ROCMFP2 &&
+        std::getenv("DFLASH_CUDA_MMVQ_FP2_AFFINE") == nullptr) {
+        use_mul_mat_vec_q = false;
+    }
+#endif // ROCMFP2_AFFINE
 
     // fusion is not universally faster on Pascal
     const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
@@ -2773,7 +2781,9 @@ static bool ggml_cuda_try_fuse_mul_mat_glu(
         }
 
         const int64_t ncols = ids ? src1->ne[2] : src1->ne[1];
-        if (ggml_cuda_should_use_mmq(src0->type, cc, ncols, src0->ne[2])) {
+        if (ggml_cuda_should_use_mmq(
+                src0->type, cc, ncols,
+                ids ? src0->ne[2] : /*n_experts=*/0)) {
             ggml_cuda_mul_mat_q_pair(
                 ctx, up->src[0], gate->src[0], src1, ids, up, gate);
             ggml_cuda_op_swiglu_ds4(ctx, glu);
@@ -2825,6 +2835,23 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
     bool use_mul_mat_q     = ggml_is_quantized(src0->type) && !bad_padding_clear
         && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32;
 
+#ifdef ROCMFP2_AFFINE
+    // Both MMVQ and MMQ carry the affine offset correction. Keep ordinary
+    // (non-fused) MMQ behind a second A/B switch: it is numerically correct on
+    // gfx1100, but the small per-expert down projections regress model prefill
+    // throughput. The phase-scoped master switch still enables the qualified
+    // paired gate/up fusion independently.
+    if (src0->type == GGML_TYPE_Q2_0_ROCMFP2) {
+        if (std::getenv("DFLASH_CUDA_MMVQ_FP2_AFFINE") == nullptr) {
+            use_mul_mat_vec_q = false;
+        }
+        if (std::getenv("DFLASH_CUDA_MMQ_FP2_AFFINE") == nullptr ||
+            std::getenv("DFLASH_CUDA_MMQ_FP2_AFFINE_GENERAL") == nullptr) {
+            use_mul_mat_q = false;
+        }
+    }
+#endif // ROCMFP2_AFFINE
+
     bool any_gpus_with_slow_fp16 = false;
 
     if (split) {
@@ -3091,8 +3118,8 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
     const int64_t ne_get_rows = ne12 * n_expert_used;
 
     std::vector<int32_t> ids_to_sorted_host;
-    ids_to_sorted_host.reserve(2*ne_get_rows);
-    std::vector<int32_t> ids_from_sorted_host(ne_get_rows);
+    ids_to_sorted_host.reserve(ne_get_rows);
+    std::vector<int32_t> ids_from_sorted_host(ne_get_rows, -1);
 
     ggml_cuda_pool_alloc<int32_t> ids_buf_dev(ctx.pool(), 2*ne_get_rows);
 
@@ -3108,32 +3135,54 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
     for (int64_t i02 = 0; i02 < ne02; ++i02) { // expert matrices
         for (int64_t i12 = 0; i12 < ne12; ++i12) { // tokens
             for (int64_t iex = 0; iex < n_expert_used; ++iex) {
-                const int32_t expert_to_use = *(const int32_t *)(ids_host.data() + i12*ids->nb[1] + iex*ids->nb[0]);
-                assert(expert_to_use >= 0 && expert_to_use < ne02);
+                const int32_t expert_to_use = *(const int32_t *)(
+                    ids_host.data() + i12*ids->nb[1] + iex*ids->nb[0]);
+                if (expert_to_use < 0) {
+                    continue;
+                }
+                assert(expert_to_use < ne02);
                 if (expert_to_use == i02) {
-                    ids_from_sorted_host[i12*n_expert_used + iex] = ids_to_sorted_host.size();
+                    ids_from_sorted_host[i12*n_expert_used + iex] =
+                        (int32_t) ids_to_sorted_host.size();
                     ids_to_sorted_host.push_back(i12*ne11 + iex % ne11);
                     tokens_per_expert[i02]++;
-                    break;
+                    // Do not break: one token may route several slots to the
+                    // same expert.
                 }
             }
         }
     }
-    GGML_ASSERT(ids_to_sorted_host.size() == size_t(ne_get_rows));
-
-    ids_to_sorted_host.insert(ids_to_sorted_host.end(), ids_from_sorted_host.begin(), ids_from_sorted_host.end());
+    const size_t n_valid_rows = ids_to_sorted_host.size();
+    const bool has_masked_rows = n_valid_rows < (size_t) ne_get_rows;
+    GGML_ASSERT(n_valid_rows + (has_masked_rows ? 1 : 0) <=
+                (size_t) ne_get_rows);
+    if (has_masked_rows) {
+        const int32_t zero_row = (int32_t) n_valid_rows;
+        for (int32_t & row : ids_from_sorted_host) {
+            if (row < 0) row = zero_row;
+        }
+    }
 
-    CUDA_CHECK(cudaMemcpyAsync(ids_buf_dev.ptr, ids_to_sorted_host.data(), 2*ne_get_rows*sizeof(int32_t), cudaMemcpyHostToDevice, stream));
+    std::vector<int32_t> sorted_ids_host(2*ne_get_rows, 0);
+    std::copy(ids_to_sorted_host.begin(), ids_to_sorted_host.end(),
+              sorted_ids_host.begin());
+    std::copy(ids_from_sorted_host.begin(), ids_from_sorted_host.end(),
+              sorted_ids_host.begin() + ne_get_rows);
+    CUDA_CHECK(cudaMemcpyAsync(ids_buf_dev.ptr, sorted_ids_host.data(),
+                               2*ne_get_rows*sizeof(int32_t),
+                               cudaMemcpyHostToDevice, stream));
     CUDA_CHECK(cudaStreamSynchronize(stream));
 
-    const int32_t * ids_to_sorted   = ids_buf_dev.ptr + 0*ne_get_rows;
-    const int32_t * ids_from_sorted = ids_buf_dev.ptr + 1*ne_get_rows;
+    const int32_t * ids_to_sorted   = ids_buf_dev.ptr;
+    const int32_t * ids_from_sorted = ids_buf_dev.ptr + ne_get_rows;
 
-    get_rows_cuda(src1->data, src1->type, ids_to_sorted, src1_sorted.ptr, type_src1_sorted,
-        ne10, nb11, nb12, nb13,
-        ne_get_rows, 1, 1, sizeof(int32_t), ne_get_rows*sizeof(int32_t), ne_get_rows*sizeof(int32_t),
-        ne10*ts_src1_sorted, ne_get_rows*ne10*ts_src1_sorted, ne_get_rows*ne10*ts_src1_sorted, stream);
-    CUDA_CHECK(cudaGetLastError());
+    if (n_valid_rows > 0) {
+        get_rows_cuda(src1->data, src1->type, ids_to_sorted, src1_sorted.ptr, type_src1_sorted,
+            ne10, nb11, nb12, nb13,
+            n_valid_rows, 1, 1, sizeof(int32_t), n_valid_rows*sizeof(int32_t), n_valid_rows*sizeof(int32_t),
+            ne10*ts_src1_sorted, n_valid_rows*ne10*ts_src1_sorted, n_valid_rows*ne10*ts_src1_sorted, stream);
+        CUDA_CHECK(cudaGetLastError());
+    }
 
     char * src1_data_cur = (char *) src1_sorted.ptr;
     char *  dst_data_cur = (char *)  dst_sorted.ptr;
@@ -3183,6 +3232,11 @@ static void ggml_cuda_mul_mat_id(ggml_backend_cuda_context & ctx, ggml_tensor *
         src1_data_cur += src1_slice.nb[2];
         dst_data_cur  +=  dst_slice.nb[2];
     }
+    if (has_masked_rows) {
+        CUDA_CHECK(cudaMemsetAsync(
+            dst_sorted.ptr + n_valid_rows*ne0*ts_dst_sorted, 0,
+            ne0*ts_dst_sorted, stream));
+    }
 
     get_rows_cuda(dst_sorted.ptr, type_dst_sorted, ids_from_sorted, dst->data, dst->type,
         ne0, ne0*ts_dst_sorted, ne_get_rows*ne0*ts_dst_sorted, ne_get_rows*ne0*ts_dst_sorted,
diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu
index e71397bfa..64a29b31f 100644
--- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu
+++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu
@@ -5,6 +5,12 @@
 #include "rocmfp2_mix.cuh"
 #include "rocmfp3_mix.cuh"
 
+static thread_local size_t g_mmq_launch_count = 0;
+
+extern "C" size_t ggml_backend_cuda_get_mmq_launch_count(void) {
+    return g_mmq_launch_count;
+}
+
 namespace {
 
 class mix_registry_dispatch_guard {
@@ -40,6 +46,7 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con
         args.type_x == GGML_TYPE_Q2_1_ROCMFP2_MIX ||
         args.type_x == GGML_TYPE_Q3_1_ROCMFP3_MIX;
     GGML_ASSERT(!is_mix_type || (args.mix_codebooks && args.mix_modes));
+    ++g_mmq_launch_count;
     switch (args.type_x) {
         case GGML_TYPE_Q4_0:
             mul_mat_q_case<GGML_TYPE_Q4_0>(ctx, args, stream);
@@ -511,6 +518,62 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
     return false;
 #endif // GGML_CUDA_FORCE_CUBLAS
 
+#ifdef ROCMFP2_AFFINE
+    // The affine Q8_1 tile loader carries scale and -offset and uses the
+    // activation sum correction. Keep it opt-in until model-level validation.
+    if (type == GGML_TYPE_Q2_0_ROCMFP2 &&
+        std::getenv("DFLASH_CUDA_MMQ_FP2_AFFINE") == nullptr) {
+        return false;
+    }
+    // Batched expert MMQ is finite with the affine correction, but its
+    // gather/quantize path is slower than grouped MMVQ at DS4 verify width.
+    if (type == GGML_TYPE_Q2_0_ROCMFP2 && n_experts > 1) {
+        return false;
+    }
+    if (type == GGML_TYPE_Q2_0_ROCMFP2) {
+        const char * runtime_disable = std::getenv(
+            "DFLASH_CUDA_MMQ_FP2_AFFINE_RUNTIME_DISABLE");
+        if (runtime_disable && *runtime_disable &&
+            std::strcmp(runtime_disable, "0") != 0) {
+            return false;
+        }
+    }
+    if (type == GGML_TYPE_Q2_0_ROCMFP2) {
+        // Owner-isolation/qualification switch. AMD architecture codes are
+        // accepted in decimal or 0x form (for example 0x1100 or 0x1151).
+        static const int required_cc = [] {
+            const char * raw = std::getenv(
+                "DFLASH_CUDA_MMQ_FP2_AFFINE_CC");
+            if (!raw || !*raw) return 0;
+            char * end = nullptr;
+            const long parsed = std::strtol(raw, &end, 0);
+            return end && end != raw && *end == '\0' && parsed > 0 &&
+                           parsed <= INT_MAX
+                ? (int) parsed
+                : 0;
+        }();
+        if (required_cc > 0 && cc != required_cc) {
+            return false;
+        }
+    }
+    if (type == GGML_TYPE_Q2_0_ROCMFP2) {
+        static const int min_ncols = [] {
+            const char * raw = std::getenv(
+                "DFLASH_CUDA_MMQ_FP2_AFFINE_MIN_NCOLS");
+            if (!raw || !*raw) return 0;
+            char * end = nullptr;
+            const long parsed = std::strtol(raw, &end, 10);
+            return end && end != raw && *end == '\0' && parsed > 0 &&
+                           parsed <= INT_MAX
+                ? (int) parsed
+                : 0;
+        }();
+        if (ne11 < min_ncols) {
+            return false;
+        }
+    }
+#endif // ROCMFP2_AFFINE
+
     bool mmq_supported;
 
     switch (type) {
diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh
index 876a8ba45..c482ce65a 100644
--- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh
+++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh
@@ -67,11 +67,16 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) {
         case GGML_TYPE_Q8_0:
             return MMQ_Q8_1_DS_LAYOUT_D4;
         case GGML_TYPE_Q4_0_ROCMFP4_FAST:
-        case GGML_TYPE_Q2_0_ROCMFP2:
         case GGML_TYPE_Q2_1_ROCMFP2_MIX:
         case GGML_TYPE_Q3_0_ROCMFPX:
         case GGML_TYPE_Q3_1_ROCMFP3_MIX:
             return MMQ_Q8_1_DS_LAYOUT_D4;
+        case GGML_TYPE_Q2_0_ROCMFP2:
+#ifdef ROCMFP2_AFFINE
+            return MMQ_Q8_1_DS_LAYOUT_DS4;
+#else
+            return MMQ_Q8_1_DS_LAYOUT_D4;
+#endif
         case GGML_TYPE_MXFP4:
             return MMQ_Q8_1_DS_LAYOUT_D4;
         case GGML_TYPE_NVFP4:
@@ -242,7 +247,12 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml
         case GGML_TYPE_Q5_1:    return MMQ_DP4A_TXS_Q8_1;
         case GGML_TYPE_Q8_0:    return MMQ_DP4A_TXS_Q8_0;
         case GGML_TYPE_Q4_0_ROCMFP4_FAST: return MMQ_DP4A_TXS_Q8_0;
-        case GGML_TYPE_Q2_0_ROCMFP2: return MMQ_DP4A_TXS_Q8_0_16;
+        case GGML_TYPE_Q2_0_ROCMFP2:
+#ifdef ROCMFP2_AFFINE
+            return MMQ_DP4A_TXS_Q8_1;
+#else
+            return MMQ_DP4A_TXS_Q8_0_16;
+#endif
         case GGML_TYPE_Q2_1_ROCMFP2_MIX: return MMQ_DP4A_TXS_Q8_0_16;
         case GGML_TYPE_Q3_0_ROCMFPX: return MMQ_DP4A_TXS_Q8_0_16;
         case GGML_TYPE_Q3_1_ROCMFP3_MIX: return MMQ_DP4A_TXS_Q8_0_16;
@@ -291,7 +301,12 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) {
         case GGML_TYPE_Q5_1:    return MMQ_MMA_TILE_X_K_Q8_1;
         case GGML_TYPE_Q8_0:    return MMQ_MMA_TILE_X_K_Q8_0;
         case GGML_TYPE_Q4_0_ROCMFP4_FAST: return MMQ_MMA_TILE_X_K_Q8_0;
-        case GGML_TYPE_Q2_0_ROCMFP2: return MMQ_MMA_TILE_X_K_Q3_K;
+        case GGML_TYPE_Q2_0_ROCMFP2:
+#ifdef ROCMFP2_AFFINE
+            return MMQ_MMA_TILE_X_K_Q8_1;
+#else
+            return MMQ_MMA_TILE_X_K_Q3_K;
+#endif
         case GGML_TYPE_Q2_1_ROCMFP2_MIX: return MMQ_MMA_TILE_X_K_Q3_K;
         case GGML_TYPE_Q3_0_ROCMFPX: return MMQ_MMA_TILE_X_K_Q3_K;
         case GGML_TYPE_Q3_1_ROCMFP3_MIX: return MMQ_MMA_TILE_X_K_Q3_K;
@@ -1389,6 +1404,95 @@ static __device__ __forceinline__ void load_tiles_rocmfp3_mix(
     }
 }
 
+#ifdef ROCMFP2_AFFINE
+// Affine FP2 is a Q8_1-style integer tile: raw codes c in [0,3] plus
+// dm=(scale,-offset). The ordinary Q8_1 dot then evaluates
+//   scale*d*dot(c,q) - offset*(d*sum(q))
+// exactly, using the activation block's stored sum for the affine correction.
+template <int mmq_y, bool need_check>
+static __device__ __forceinline__ void load_tiles_rocmfp2_affine(
+    const char * __restrict__ x, int * __restrict__ x_tile,
+    const int kbx0, const int i_max, const int stride) {
+    constexpr int nwarps = mmq_get_nwarps_device();
+    constexpr int warp_size = ggml_cuda_get_physical_warp_size();
+    constexpr int groups_per_block = QK_ROCMFP2 / 4;
+    constexpr int blocks_per_tile = MMQ_ITER_K / QK_ROCMFP2;
+    constexpr int threads_per_row =
+        blocks_per_tile * groups_per_block / 2;
+    static_assert(threads_per_row == 32,
+                  "affine ROCmFP2 MMQ loader expects 32 lanes per row");
+
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+    int   * x_qs = (int *) x_tile;
+    half2 * x_dm = (half2 *) (x_qs + 2*MMQ_TILE_NE_K);
+#else
+    constexpr tile_x_sizes txs =
+        mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q2_0_ROCMFP2, mmq_y);
+    int   * x_qs = (int *) x_tile;
+    half2 * x_dm = (half2 *) (x_qs + txs.qs);
+#endif
+
+    constexpr int nrows = warp_size / threads_per_row;
+    const int txi = warp_size > threads_per_row
+        ? threadIdx.x % threads_per_row : threadIdx.x;
+    const int kbx = txi / (groups_per_block / 2);
+    const int group = txi % (groups_per_block / 2);
+
+    auto pack_raw4 = [](const block_rocmfp2 * block, int base) {
+        const uint32_t bits = block->qs[base >> 2];
+        return (int) (((bits >> 0) & 3u) |
+                      (((bits >> 2) & 3u) << 8) |
+                      (((bits >> 4) & 3u) << 16) |
+                      (((bits >> 6) & 3u) << 24));
+    };
+#pragma unroll
+    for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) {
+        int i = i0 + (nrows == 1
+            ? threadIdx.y
+            : threadIdx.y*nrows + threadIdx.x/threads_per_row);
+        if (need_check) i = min(i, i_max);
+
+        const block_rocmfp2 * block =
+            (const block_rocmfp2 *) x + kbx0 + i*stride + kbx;
+        const int k0 = kbx*groups_per_block + group;
+        const int q0 = pack_raw4(block, 4*group);
+        const int q1 = pack_raw4(
+            block, 4*(group + groups_per_block/2));
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+        x_qs[i*MMQ_MMA_TILE_X_K_Q8_1 + k0] = q0;
+        x_qs[i*MMQ_MMA_TILE_X_K_Q8_1 + k0 + groups_per_block/2] = q1;
+#else
+        x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0] = q0;
+        x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + groups_per_block/2] = q1;
+#endif
+    }
+
+    constexpr int scale_rows_per_warp = warp_size / blocks_per_tile;
+    const int kscale = threadIdx.x % blocks_per_tile;
+#pragma unroll
+    for (int i0 = 0; i0 < mmq_y;
+         i0 += nwarps*scale_rows_per_warp) {
+        int i = i0 + threadIdx.y*scale_rows_per_warp +
+                threadIdx.x/blocks_per_tile;
+        if (need_check) i = min(i, i_max);
+
+        const block_rocmfp2 * block =
+            (const block_rocmfp2 *) x + kbx0 + i*stride + kscale;
+        const float scale =
+            rocmfpx_ue4m3_to_fp32_finite(block->e[0]);
+        const float offset =
+            rocmfpx_ue4m3_to_fp32_finite(block->e[1]);
+        const half2 dm = make_half2(scale, -offset);
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+        x_dm[i*MMQ_MMA_TILE_X_K_Q8_1 + kscale] = dm;
+#else
+        x_dm[i*(2*MMQ_TILE_NE_K/QI8_1) +
+             i/(QI8_1/2) + kscale] = dm;
+#endif
+    }
+}
+#endif
+
 template <int mmq_y, bool need_check>
 static __device__ __forceinline__ void load_tiles_mxfp4_fp4(const char * __restrict__ x,
                                                             int * __restrict__ x_tile,
@@ -3947,9 +4051,15 @@ struct mmq_type_traits<mmq_x, mmq_y, need_check, GGML_TYPE_Q4_0_ROCMFP4_FAST> {
 template <int mmq_x, int mmq_y, bool need_check>
 struct mmq_type_traits<mmq_x, mmq_y, need_check, GGML_TYPE_Q2_0_ROCMFP2> {
     static constexpr int              vdr          = VDR_ROCMFP2_Q8_1_MMQ;
+#ifdef ROCMFP2_AFFINE
+    static constexpr load_tiles_mmq_t load_tiles   = load_tiles_rocmfp2_affine<mmq_y, need_check>;
+    static constexpr vec_dot_mmq_t    vec_dot_mma  = vec_dot_q8_1_q8_1_mma<mmq_x, mmq_y>;
+    static constexpr vec_dot_mmq_t    vec_dot_dp4a = vec_dot_q8_1_q8_1_dp4a<mmq_x, mmq_y>;
+#else
     static constexpr load_tiles_mmq_t load_tiles   = load_tiles_rocmfpx_dual<GGML_TYPE_Q2_0_ROCMFP2, mmq_y, need_check>;
     static constexpr vec_dot_mmq_t    vec_dot_mma  = vec_dot_q8_0_16_q8_1_mma<mmq_x, mmq_y>;
     static constexpr vec_dot_mmq_t    vec_dot_dp4a = vec_dot_q8_0_16_q8_1_dp4a<mmq_x, mmq_y>;
+#endif
 };
 
 template <int mmq_x, int mmq_y, bool need_check>
diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu
index b957d1ed1..f41112041 100644
--- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu
+++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu
@@ -8,6 +8,12 @@
 #include <cstdlib>
 #include <cstring>
 
+static thread_local size_t g_mmvq_launch_count = 0;
+
+extern "C" size_t ggml_backend_cuda_get_mmvq_launch_count(void) {
+    return g_mmvq_launch_count;
+}
+
 typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs);
 
 static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) {
@@ -137,7 +143,11 @@ static __device__ __forceinline__ float vec_dot_rocmfpx_fp2_q8_1_packed32(
     }
 
     const float db = __low2float(bq8_1->ds);
+#ifdef ROCMFP2_AFFINE
+    return rocmfpx_fp2_affine_dot(bq2, bq8_1, sumi, iqs, db);
+#else
     return db * rocmfpx_ue4m3_to_fp32_finite(bq2->e[iqs]) * sumi;
+#endif
 }
 
 // Dense q4 verification applies one weight row to four token activations.
@@ -409,7 +419,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type
 //                               tokens unless enabled. The ROCmFP formats are
 //                               opt-in until qualified on each AMD target.
 #define MMID_GROUPED_MAX_PAIRS 256
-#define MMID_GROUPED_MAX_TPG   8
+#define MMID_GROUPED_MAX_TPG   2
 #define MMID_META_NG 0
 #define MMID_META_GE 1
 #define MMID_META_GS (MMID_META_GE + MMID_GROUPED_MAX_PAIRS)
@@ -1485,7 +1495,7 @@ static __global__ void mmid_group_prep(
 // Weight traffic approaches (union of routed experts) instead of
 // (n_expert_used x n_tokens) expert-matrix reads.
 template <ggml_type type, int c_rows_per_block, bool has_fusion,
-          bool c_fp3_packed24 = false, bool c_fp2_packed32 = false>
+          bool fp3_packed24 = false, bool fp2_packed32 = false>
 __launch_bounds__(MMID_GROUPED_MAX_TPG*ggml_cuda_get_physical_warp_size(), 1)
 static __global__ void mul_mat_vec_q_moe_grouped(
         const void * __restrict__ vx, const void * __restrict__ vy, const int32_t * __restrict__ meta,
@@ -1501,6 +1511,10 @@ static __global__ void mul_mat_vec_q_moe_grouped(
     constexpr int warp_size = ggml_cuda_get_physical_warp_size();
 
     constexpr vec_dot_q_cuda_t vec_dot_q_cuda = get_vec_dot_q_cuda(type);
+    static_assert(!fp3_packed24 || type == GGML_TYPE_Q3_0_ROCMFPX,
+                  "packed FP3 grouped dispatch requires ROCmFP3 weights");
+    static_assert(!fp2_packed32 || type == GGML_TYPE_Q2_0_ROCMFP2,
+                  "packed FP2 grouped dispatch requires ROCmFP2 weights");
 
     const uint32_t p = blockIdx.y*MMID_GROUPED_MAX_TPG + threadIdx.y;
     if (p >= np) {
@@ -1563,13 +1577,11 @@ static __global__ void mul_mat_vec_q_moe_grouped(
 
 #pragma unroll
         for (int i = 0; i < c_rows_per_block; ++i) {
-            if constexpr (type == GGML_TYPE_Q2_0_ROCMFP2 &&
-                          c_fp2_packed32) {
+            if constexpr (fp2_packed32) {
                 tmp[i] += vec_dot_rocmfpx_fp2_q8_1_packed32(
                     vx, &y[kby],
                     kbx_offset + i*stride_row_x + kbx, kqs);
-            } else if constexpr (type == GGML_TYPE_Q3_0_ROCMFPX &&
-                                 c_fp3_packed24) {
+            } else if constexpr (fp3_packed24) {
                 tmp[i] += vec_dot_rocmfpx_fp3_q8_1_packed24(
                     vx, &y[kby],
                     kbx_offset + i*stride_row_x + kbx, kqs);
@@ -1580,13 +1592,11 @@ static __global__ void mul_mat_vec_q_moe_grouped(
             }
             if constexpr (has_fusion) {
                 if (use_gate) {
-                    if constexpr (type == GGML_TYPE_Q2_0_ROCMFP2 &&
-                                  c_fp2_packed32) {
+                    if constexpr (fp2_packed32) {
                         tmp_gate[i] += vec_dot_rocmfpx_fp2_q8_1_packed32(
                             vgate, &y[kby],
                             kbx_offset + i*stride_row_x + kbx, kqs);
-                    } else if constexpr (type == GGML_TYPE_Q3_0_ROCMFPX &&
-                                         c_fp3_packed24) {
+                    } else if constexpr (fp3_packed24) {
                         tmp_gate[i] += vec_dot_rocmfpx_fp3_q8_1_packed24(
                             vgate, &y[kby],
                             kbx_offset + i*stride_row_x + kbx, kqs);
@@ -1658,8 +1668,8 @@ static __global__ void mul_mat_vec_q_moe_grouped(
     }
 }
 
-template <ggml_type type, bool c_fp3_packed24 = false,
-          bool c_fp2_packed32 = false>
+template <ggml_type type, bool fp3_packed24 = false,
+          bool fp2_packed32 = false>
 static void mul_mat_vec_q_moe_grouped_launch(
         const void * vx, const void * vy, const int32_t * meta, const ggml_cuda_mm_fusion_args_device & fusion, float * dst,
         const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t nrows_x,
@@ -1667,26 +1677,28 @@ static void mul_mat_vec_q_moe_grouped_launch(
         const uint32_t stride_channel_x, const uint32_t stride_channel_y, const uint32_t stride_channel_dst,
         const int np, const int warp_size, cudaStream_t stream) {
 
-    constexpr int rows_per_block = 2;
+    constexpr int rows_per_block = 4;
     const int64_t nblocks_rows = (nrows_x + rows_per_block - 1)/rows_per_block;
     const dim3 block_nums(nblocks_rows, (np + MMID_GROUPED_MAX_TPG - 1)/MMID_GROUPED_MAX_TPG);
     const dim3 block_dims(warp_size, MMID_GROUPED_MAX_TPG);
 
     const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr;
     if (has_fusion) {
-        mul_mat_vec_q_moe_grouped<type, rows_per_block, true,
-                                  c_fp3_packed24,
-                                  c_fp2_packed32><<<block_nums, block_dims, 0, stream>>>(
-            vx, vy, meta, fusion, dst, (uint32_t) np, ncols_x, nchannels_y, nrows_x,
-            stride_row_x, stride_col_y, stride_col_dst,
-            stride_channel_x, stride_channel_y, stride_channel_dst);
+        mul_mat_vec_q_moe_grouped<
+            type, rows_per_block, true, fp3_packed24, fp2_packed32>
+            <<<block_nums, block_dims, 0, stream>>>(
+                vx, vy, meta, fusion, dst, (uint32_t) np, ncols_x,
+                nchannels_y, nrows_x, stride_row_x, stride_col_y,
+                stride_col_dst, stride_channel_x, stride_channel_y,
+                stride_channel_dst);
     } else {
-        mul_mat_vec_q_moe_grouped<type, rows_per_block, false,
-                                  c_fp3_packed24,
-                                  c_fp2_packed32><<<block_nums, block_dims, 0, stream>>>(
-            vx, vy, meta, fusion, dst, (uint32_t) np, ncols_x, nchannels_y, nrows_x,
-            stride_row_x, stride_col_y, stride_col_dst,
-            stride_channel_x, stride_channel_y, stride_channel_dst);
+        mul_mat_vec_q_moe_grouped<
+            type, rows_per_block, false, fp3_packed24, fp2_packed32>
+            <<<block_nums, block_dims, 0, stream>>>(
+                vx, vy, meta, fusion, dst, (uint32_t) np, ncols_x,
+                nchannels_y, nrows_x, stride_row_x, stride_col_y,
+                stride_col_dst, stride_channel_x, stride_channel_y,
+                stride_channel_dst);
     }
 }
 
@@ -1700,16 +1712,6 @@ static bool mul_mat_vec_q_grouped_dispatch(
 
     const int warp_size = ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
     const uint3 nchannels_y_fd = init_fastdiv_values((uint32_t) nchannels_y);
-    static const bool fp3_packed24 = []() {
-        const char * e =
-            std::getenv("DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24");
-        return e && e[0] == '1' && e[1] == '\0';
-    }();
-    static const bool fp2_packed32 = []() {
-        const char * e =
-            std::getenv("DFLASH_CUDA_MMVQ_MOE_FP2_PACKED32");
-        return e && e[0] == '1' && e[1] == '\0';
-    }();
 
     switch (type) {
         case GGML_TYPE_Q4_0:
@@ -1733,39 +1735,45 @@ static bool mul_mat_vec_q_grouped_dispatch(
                 stride_row_x, stride_col_y, stride_col_dst, stride_channel_x, stride_channel_y, stride_channel_dst, max_groups, warp_size, stream);
             return true;
         case GGML_TYPE_Q2_0_ROCMFP2:
-            if (fp2_packed32) {
+            if (mmvq_env_flag("DFLASH_CUDA_MMVQ_MOE_FP2_PACKED32")) {
                 mul_mat_vec_q_moe_grouped_launch<
                     GGML_TYPE_Q2_0_ROCMFP2, false, true>(
                         vx, vy, meta, fusion, dst, ncols_x, nchannels_y_fd,
                         nrows_x, stride_row_x, stride_col_y, stride_col_dst,
-                        stride_channel_x, stride_channel_y,
-                        stride_channel_dst, max_groups, warp_size, stream);
+                        stride_channel_x, stride_channel_y, stride_channel_dst,
+                        max_groups, warp_size, stream);
             } else {
                 mul_mat_vec_q_moe_grouped_launch<
                     GGML_TYPE_Q2_0_ROCMFP2>(
                         vx, vy, meta, fusion, dst, ncols_x, nchannels_y_fd,
                         nrows_x, stride_row_x, stride_col_y, stride_col_dst,
-                        stride_channel_x, stride_channel_y,
-                        stride_channel_dst, max_groups, warp_size, stream);
+                        stride_channel_x, stride_channel_y, stride_channel_dst,
+                        max_groups, warp_size, stream);
             }
             return true;
-        case GGML_TYPE_Q3_0_ROCMFPX:
-            if (fp3_packed24) {
+        case GGML_TYPE_Q3_0_ROCMFPX: {
+            const bool packed =
+                mmvq_env_flag("DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24") &&
+                std::getenv(
+                    "DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24_RUNTIME_DISABLE") ==
+                    nullptr;
+            if (packed) {
                 mul_mat_vec_q_moe_grouped_launch<
-                    GGML_TYPE_Q3_0_ROCMFPX, true, false>(
+                    GGML_TYPE_Q3_0_ROCMFPX, true>(
                         vx, vy, meta, fusion, dst, ncols_x, nchannels_y_fd,
                         nrows_x, stride_row_x, stride_col_y, stride_col_dst,
-                        stride_channel_x, stride_channel_y,
-                        stride_channel_dst, max_groups, warp_size, stream);
+                        stride_channel_x, stride_channel_y, stride_channel_dst,
+                        max_groups, warp_size, stream);
             } else {
                 mul_mat_vec_q_moe_grouped_launch<
                     GGML_TYPE_Q3_0_ROCMFPX>(
                         vx, vy, meta, fusion, dst, ncols_x, nchannels_y_fd,
                         nrows_x, stride_row_x, stride_col_y, stride_col_dst,
-                        stride_channel_x, stride_channel_y,
-                        stride_channel_dst, max_groups, warp_size, stream);
+                        stride_channel_x, stride_channel_y, stride_channel_dst,
+                        max_groups, warp_size, stream);
             }
             return true;
+        }
         default:
             return false;
     }
@@ -1996,11 +2004,13 @@ static void mul_mat_vec_q_moe_launch(
             std::getenv("DFLASH_CUDA_MMVQ_MOE_Q4_WARP_GROUPS");
         return e && e[0] == '2' && e[1] == '\0';
     }();
-    static const bool fp3_packed24 = []() {
+    static const bool fp3_packed24_configured = []() {
         const char * e =
             std::getenv("DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24");
         return e && e[0] == '1' && e[1] == '\0';
     }();
+    const bool fp3_packed24 = fp3_packed24_configured &&
+        std::getenv("DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24_RUNTIME_DISABLE") == nullptr;
     static const bool fp2_packed32 = []() {
         const char * e =
             std::getenv("DFLASH_CUDA_MMVQ_MOE_FP2_PACKED32");
@@ -2618,6 +2628,8 @@ void ggml_cuda_mul_mat_vec_q(
     GGML_ASSERT(        dst->type  == GGML_TYPE_F32);
     GGML_ASSERT(!ids || ids->type  == GGML_TYPE_I32); // Optional, used for batched GGML_MUL_MAT_ID.
 
+    ++g_mmvq_launch_count;
+
     GGML_TENSOR_BINARY_OP_LOCALS;
 
     cudaStream_t stream = ctx.stream();
diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh
index f39e8cf5f..f8dc4335d 100644
--- a/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh
+++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh
@@ -509,6 +509,28 @@ static __device__ __forceinline__ float vec_dot_rocmfp4_fast_q8_1(
     return __low2float(bq8_1->ds) * rocmfp4_ue4m3_to_fp32_half_finite(bq4->e) * sumi;
 }
 
+#ifdef ROCMFP2_AFFINE
+static __device__ __forceinline__ float rocmfpx_fp2_affine_dot(
+        const block_rocmfp2 * __restrict__ weights,
+        const block_q8_1 * __restrict__ activations,
+        const int sumi, const int iqs, const float activation_scale) {
+    int sumq = 0;
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const int q = get_int_b4(activations->qs, 4*iqs + j);
+        sumq += (int8_t)(q & 0xFF) + (int8_t)((q >> 8) & 0xFF)
+              + (int8_t)((q >> 16) & 0xFF)
+              + (int8_t)((q >> 24) & 0xFF);
+    }
+    const float scale =
+        rocmfpx_ue4m3_to_fp32_finite(weights->e[0]);
+    const float offset =
+        rocmfpx_ue4m3_to_fp32_finite(weights->e[1]);
+    return activation_scale *
+        (scale * sumi + (scale - offset) * sumq);
+}
+#endif
+
 static __device__ __forceinline__ float vec_dot_rocmfpx_fp2_q8_1(
     const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
     // int32-consistent MMVQ layout: QI_ROCMFP2=2 (8B qs = 2 int32), VDR=1.
@@ -525,7 +547,11 @@ static __device__ __forceinline__ float vec_dot_rocmfpx_fp2_q8_1(
     }
 
     const float db = __low2float(bq8_1->ds);
+#ifdef ROCMFP2_AFFINE
+    return rocmfpx_fp2_affine_dot(bq2, bq8_1, sumi, iqs, db);
+#else
     return db * rocmfpx_ue4m3_to_fp32_finite(bq2->e[iqs]) * sumi;
+#endif
 }
 
 static __device__ __forceinline__ float vec_dot_rocmfpx_fp3_q8_1(
diff --git a/server/deps/llama.cpp/ggml/src/ggml-hip/CMakeLists.txt b/server/deps/llama.cpp/ggml/src/ggml-hip/CMakeLists.txt
index 46e6a2928..67dfedbae 100644
--- a/server/deps/llama.cpp/ggml/src/ggml-hip/CMakeLists.txt
+++ b/server/deps/llama.cpp/ggml/src/ggml-hip/CMakeLists.txt
@@ -160,4 +160,11 @@ if (GGML_HIP_RCCL)
     target_link_libraries(ggml-hip PRIVATE ggml-base roc::rccl)
 endif()
 
+# clang's HIP driver searches system library directories before its implicit
+# ROCm runtime directory. Put the package's lib directory first so --hip-link
+# cannot pull an older system libamdhip64 alongside the configured runtime.
+get_filename_component(GGML_HIP_RUNTIME_DIR "${hip_DIR}/../.." ABSOLUTE)
+target_link_directories(ggml-hip BEFORE PRIVATE "${GGML_HIP_RUNTIME_DIR}")
+target_link_options(ggml-hip PRIVATE "-L${GGML_HIP_RUNTIME_DIR}")
+
 target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas)
diff --git a/server/docs/DS4.md b/server/docs/DS4.md
index 3de39b571..6bd923717 100644
--- a/server/docs/DS4.md
+++ b/server/docs/DS4.md
@@ -145,6 +145,94 @@ decode and 415.52 tok/s median sparse prefill. Those numbers require the full
 qualified manifest, including the burn-in kernel switches; they are not a
 claim for the minimal activation example above.
 
+#### Radeon RX 7900 XT + Strix Halo, true top-k-6
+
+The portable `scripts/serve_ds4_dual_rocm_128k.sh` profile captures the
+qualified fixed-codebook ROCmFPx configuration for a 20 GiB gfx1100 discrete
+GPU plus a 128 GiB gfx1151 Strix Halo. It keeps the model-default six routed
+experts, places dense work and calibrated hot experts on gfx1100, and runs the
+remaining experts and DSpark drafter on gfx1151. This profile is for uniform
+or fixed-codebook expert tensors. Adaptive qtype-105/106 expert tensors remain
+monolithic-only because their codebook registrations cannot follow a sliced
+expert allocation.
+
+Build for both targets with affine qtype-107 support enabled:
+
+```bash
+cmake -S server -B server/build-hip-dual \
+  -DDFLASH27B_GPU_BACKEND=hip \
+  -DDFLASH27B_HIP_ARCHITECTURES='gfx1100;gfx1151' \
+  -DDFLASH27B_ROCMFP2_AFFINE=ON \
+  -DGGML_HIP_GRAPHS=ON \
+  -DCMAKE_BUILD_TYPE=Release
+cmake --build server/build-hip-dual -j
+```
+
+Placement depends on the workload's routing distribution. Capture a profile
+with representative requests, stop the server so it flushes the CSV, then use
+that CSV for serving:
+
+```bash
+DFLASH_DS4_ROUTING_STATS_OUT=/tmp/ds4-routing.csv \
+  server/scripts/serve_ds4_dual_rocm_128k.sh \
+  /path/to/target.gguf /path/to/dspark.gguf
+
+DFLASH_DS4_HOTNESS_CSV=/tmp/ds4-routing.csv \
+  server/scripts/serve_ds4_dual_rocm_128k.sh \
+  /path/to/target.gguf /path/to/dspark.gguf
+```
+
+The script defaults to `hip:0` for gfx1100, peer and draft device `1` for
+gfx1151, a 10,200 MiB discrete-GPU expert budget, 135,168 context slots
+(128 KiB prompt plus 4 KiB generation headroom), sparse prefill, and disabled
+prefix/prefill caches. Every device, budget, context, build, host, and port
+setting is environment-overridable at the top of the script.
+
+The controlled decode workload is deliberately predictable: temperature zero,
+true top-k-6, fixed DSpark q=4, one warm-up, three fresh 512-token `BETA`
+sequence requests, and no prefix or prefill cache. It measures the upper-bound
+speculative path, not typical conversational or agent throughput. Run the
+checked-in harness against the qualified profile:
+
+```bash
+python3 server/scripts/bench_ds4_decode.py \
+  --url http://127.0.0.1:8016 \
+  --model dflash \
+  --warmups 1 \
+  --runs 3 \
+  --max-tokens 512
+```
+
+The harness rejects partial generations, cache hits, and autoregressive
+fallbacks, requires the measured outputs to be byte-identical, and records
+server-reported decode timing, acceptance, and output SHA-256 values as JSON.
+On a Radeon RX 7900 XT plus
+Ryzen AI Max+ 395, a clean rebase onto upstream `main` using the script's
+default 10,200 MiB expert budget and 135,168-token context produced 100%
+speculative acceptance and one byte-identical output digest:
+
+| Run | Decode time | Throughput |
+|---:|---:|---:|
+| 1 | 11.3772 s | 45.0 tok/s |
+| 2 | 10.9023 s | 47.0 tok/s |
+| 3 | 10.7302 s | 47.7 tok/s |
+
+Retain that JSON with the target and draft GGUF SHA-256 values, build revision,
+ROCm version, GPU identifiers, launch environment, and context length. An
+earlier qualified build produced three 46.8 tok/s runs; its previous
+same-hardware route/dispatch baseline was 37.7 tok/s, a 24.1% improvement.
+The checked-in kernel test verifies actual affine MMQ dispatch on gfx1151 and
+gfx12xx. gfx1100 uses the supported non-MMQ fallback and is not claimed as an
+MMQ qualification.
+
+Sparse-prefill records used several explicitly different approximation
+profiles and must not be compared as if the launch settings were identical.
+A true top-k-6 128 KiB sweep measured 111.2 tok/s at 132,981 tokens. Older
+top-k-4 profiles measured 194.51 tok/s at 12,281 tokens, 158.9-164.3 tok/s
+around 2K tokens, and 53.1 tok/s at 257,965 tokens. Sparse prefill is an
+explicit approximation; applications that require reference-exact prompt
+ingestion must use `--ds4-prefill exact` and remeasure.
+
 #### CUDA 3090 + Strix Halo in one process
 
 The mixed-vendor build links the selected target runtime normally and loads
diff --git a/server/scripts/bench_ds4_decode.py b/server/scripts/bench_ds4_decode.py
new file mode 100755
index 000000000..ec513349d
--- /dev/null
+++ b/server/scripts/bench_ds4_decode.py
@@ -0,0 +1,140 @@
+#!/usr/bin/env python3
+"""Run the controlled DeepSeek4 speculative-decode throughput benchmark."""
+
+import argparse
+import hashlib
+import json
+import sys
+import time
+import urllib.error
+import urllib.request
+
+PROMPT = (
+    "Continue this exact sequence indefinitely. Output only the word BETA "
+    "separated by single spaces and never stop before the token limit: "
+    "BETA BETA BETA BETA BETA BETA BETA BETA"
+)
+
+
+def is_beta_sequence(content: str) -> bool:
+    words = content.split()
+    return bool(words) and all(word == "BETA" for word in words)
+
+
+def run_request(url: str, model: str, max_tokens: int, timeout: float) -> dict:
+    body = json.dumps(
+        {
+            "model": model,
+            "messages": [{"role": "user", "content": PROMPT}],
+            "max_tokens": max_tokens,
+            "temperature": 0,
+            "stream": False,
+        },
+        separators=(",", ":"),
+    ).encode()
+    request = urllib.request.Request(
+        f"{url.rstrip('/')}/v1/chat/completions",
+        data=body,
+        headers={"Content-Type": "application/json"},
+    )
+
+    started = time.monotonic()
+    with urllib.request.urlopen(request, timeout=timeout) as response:
+        result = json.load(response)
+    wall_seconds = time.monotonic() - started
+
+    choice = result["choices"][0]
+    content = choice["message"]["content"]
+    usage = result["usage"]
+    timings = usage.get("timings", {})
+    return {
+        "completion_tokens": usage["completion_tokens"],
+        "decode_seconds": timings.get("decode_ms", 0.0) / 1000.0,
+        "decode_tokens_per_second": timings.get("decode_tokens_per_sec"),
+        "accept_rate": usage.get("accept_rate"),
+        "spec_decode_ran": usage.get("spec_decode_ran"),
+        "cache_hit": timings.get("cache_hit"),
+        "cached_prefix_tokens": timings.get("cached_prefix_tokens"),
+        "finish_reason": choice.get("finish_reason"),
+        "output_sha256": hashlib.sha256(content.encode()).hexdigest(),
+        "output_matches_prompt": is_beta_sequence(content),
+        "wall_seconds": wall_seconds,
+    }
+
+
+def validate_run(run: dict, max_tokens: int) -> None:
+    if run["completion_tokens"] != max_tokens:
+        raise RuntimeError(
+            f"expected {max_tokens} completion tokens, got "
+            f"{run['completion_tokens']}"
+        )
+    if not run["output_matches_prompt"]:
+        raise RuntimeError("model output did not contain only the requested BETA sequence")
+    if run["cache_hit"] or run["cached_prefix_tokens"] not in (None, 0):
+        raise RuntimeError("benchmark request reused a cached prefix")
+    if run["spec_decode_ran"] is not True:
+        raise RuntimeError("speculative decode did not run")
+    throughput = run["decode_tokens_per_second"]
+    if throughput is None or throughput <= 0:
+        raise RuntimeError("server did not report positive decode throughput")
+
+
+def main() -> int:
+    parser = argparse.ArgumentParser(
+        description=(
+            "Run one warm-up and three fresh 512-token BETA requests against "
+            "an OpenAI-compatible dflash_server."
+        )
+    )
+    parser.add_argument("--url", default="http://127.0.0.1:8016")
+    parser.add_argument("--model", default="dflash")
+    parser.add_argument("--max-tokens", type=int, default=512)
+    parser.add_argument("--warmups", type=int, default=1)
+    parser.add_argument("--runs", type=int, default=3)
+    parser.add_argument("--timeout", type=float, default=900.0)
+    args = parser.parse_args()
+
+    if args.max_tokens <= 0 or args.warmups < 0 or args.runs <= 0:
+        parser.error("max-tokens and runs must be positive; warmups cannot be negative")
+
+    try:
+        warmups = [
+            run_request(args.url, args.model, args.max_tokens, args.timeout)
+            for _ in range(args.warmups)
+        ]
+        runs = []
+        for _ in range(args.runs):
+            run = run_request(args.url, args.model, args.max_tokens, args.timeout)
+            validate_run(run, args.max_tokens)
+            runs.append(run)
+    except (KeyError, TypeError, ValueError, urllib.error.URLError, RuntimeError) as error:
+        print(f"benchmark failed: {error}", file=sys.stderr)
+        return 1
+
+    output_digests = {run["output_sha256"] for run in runs}
+    if len(output_digests) != 1:
+        print("benchmark failed: measured outputs were not byte-identical", file=sys.stderr)
+        return 1
+
+    print(
+        json.dumps(
+            {
+                "request": {
+                    "model": args.model,
+                    "prompt": PROMPT,
+                    "temperature": 0,
+                    "max_tokens": args.max_tokens,
+                    "stream": False,
+                },
+                "warmups": warmups,
+                "runs": runs,
+            },
+            indent=2,
+            sort_keys=True,
+        )
+    )
+    return 0
+
+
+if __name__ == "__main__":
+    raise SystemExit(main())
diff --git a/server/scripts/serve_ds4_dual_rocm_128k.sh b/server/scripts/serve_ds4_dual_rocm_128k.sh
new file mode 100755
index 000000000..cfd755591
--- /dev/null
+++ b/server/scripts/serve_ds4_dual_rocm_128k.sh
@@ -0,0 +1,114 @@
+#!/usr/bin/env bash
+# Qualified DeepSeek4 profile for a 20 GiB discrete AMD GPU plus 128 GiB
+# Strix Halo. The discrete GPU must be hip:0 and Strix Halo hip:1 unless the
+# device variables below are overridden.
+set -euo pipefail
+
+if [[ $# -lt 2 ]]; then
+    echo "usage: $0 TARGET_GGUF DSPARK_GGUF [dflash_server arguments...]" >&2
+    exit 2
+fi
+
+target=$1
+draft=$2
+shift 2
+
+script_dir=$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)
+server_dir=$(cd "$script_dir/.." && pwd)
+build_dir=${DFLASH_BUILD_DIR:-$server_dir/build-hip-dual}
+binary=$build_dir/dflash_server
+
+for path in "$target" "$draft"; do
+    if [[ ! -f $path ]]; then
+        echo "missing required file: $path" >&2
+        exit 1
+    fi
+done
+if [[ ! -f $binary || ! -x $binary ]]; then
+    echo "missing executable: $binary" >&2
+    exit 1
+fi
+
+main_device=${DS4_MAIN_DEVICE:-hip:0}
+peer_device=${DS4_PEER_DEVICE:-1}
+draft_device=${DS4_DRAFT_DEVICE:-1}
+context=${DS4_CONTEXT:-135168}
+expert_budget_mb=${DS4_EXPERT_BUDGET_MB:-10200}
+port=${DS4_PORT:-8016}
+
+export DFLASH_DS4_MOE_TP=1
+export DFLASH_DS4_MOE_TP_INPROC=1
+export DFLASH_DS4_MOE_TP_GPU=$peer_device
+export DFLASH_EXPERT_BUDGET_MB=$expert_budget_mb
+export DFLASH_DS4_TP_MAIN_TO_PEER_RATE=${DFLASH_DS4_TP_MAIN_TO_PEER_RATE:-100}
+export DFLASH_DS4_TP_BALANCE_MIN_HOT=${DFLASH_DS4_TP_BALANCE_MIN_HOT:-31}
+export DFLASH_DS4_TP_BALANCE_MAX_HOT=${DFLASH_DS4_TP_BALANCE_MAX_HOT:-50}
+
+# A calibrated routing profile enables the critical-path placement policy.
+# Without one, the server remains usable but placement and throughput are not
+# the qualified profile documented in docs/DS4.md.
+if [[ -n ${DFLASH_DS4_HOTNESS_CSV:-} ]]; then
+    if [[ ! -f $DFLASH_DS4_HOTNESS_CSV ]]; then
+        echo "missing routing profile: $DFLASH_DS4_HOTNESS_CSV" >&2
+        exit 1
+    fi
+    export DFLASH_DS4_TP_CRITICAL_PATH_PLACEMENT=${DFLASH_DS4_TP_CRITICAL_PATH_PLACEMENT:-1}
+else
+    echo "warning: DFLASH_DS4_HOTNESS_CSV is unset; using uncalibrated placement" >&2
+fi
+
+export DFLASH_DS4_LONG_CONTEXT_CHUNK=${DFLASH_DS4_LONG_CONTEXT_CHUNK:-2048}
+export DFLASH_DS4_DISABLE_LONG_CONTEXT_ARENA_HANDOFF=${DFLASH_DS4_DISABLE_LONG_CONTEXT_ARENA_HANDOFF:-1}
+export DFLASH_CUDA_MMQ_FP2_AFFINE_PREFILL_ONLY=${DFLASH_CUDA_MMQ_FP2_AFFINE_PREFILL_ONLY:-1}
+export DFLASH_CUDA_MMQ_FP2_AFFINE_CAPTURE=${DFLASH_CUDA_MMQ_FP2_AFFINE_CAPTURE:-1}
+export DFLASH_MOE_PREFILL_MASKED_COLD=${DFLASH_MOE_PREFILL_MASKED_COLD:-0}
+export DFLASH_DS4_HYBRID_PREFILL_GPU_HC=${DFLASH_DS4_HYBRID_PREFILL_GPU_HC:-1}
+export DFLASH_DS4_HYBRID_PREFILL_EAGER=${DFLASH_DS4_HYBRID_PREFILL_EAGER:-1}
+export DFLASH_MOE_EXPERT_MAJOR_PINNED_OUTPUT=${DFLASH_MOE_EXPERT_MAJOR_PINNED_OUTPUT:-1}
+
+export LUCE_MMVQ_MAX_NCOLS=${LUCE_MMVQ_MAX_NCOLS:-4}
+export DFLASH_HIP_NO_AUTO_UMA=${DFLASH_HIP_NO_AUTO_UMA:-1}
+export DFLASH_DS4_TP_GROUPED_MMVQ=${DFLASH_DS4_TP_GROUPED_MMVQ:-1}
+export DFLASH_MMID_GROUPED=${DFLASH_MMID_GROUPED:-1}
+export DFLASH_MMID_GROUPED_TYPES=${DFLASH_MMID_GROUPED_TYPES:-15}
+export DFLASH_CUDA_MMVQ_MOE_FP2_PACKED32=${DFLASH_CUDA_MMVQ_MOE_FP2_PACKED32:-1}
+export DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24=${DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24:-1}
+export DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24_DECODE_ONLY=${DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24_DECODE_ONLY:-1}
+export DFLASH_CUDA_MMVQ_FP4_X4=${DFLASH_CUDA_MMVQ_FP4_X4:-1}
+export DFLASH_DS4_TP_MASKED_ROUTES=${DFLASH_DS4_TP_MASKED_ROUTES:-1}
+export DFLASH_DS4_TP_DEVICE_JOIN=${DFLASH_DS4_TP_DEVICE_JOIN:-1}
+export DFLASH_DS4_TP_NATIVE_ROUTE_WIDTH=${DFLASH_DS4_TP_NATIVE_ROUTE_WIDTH:-1}
+export DFLASH_DS4_TP_SPLIT_COUNT=${DFLASH_DS4_TP_SPLIT_COUNT:-1}
+export DFLASH_DS4_TP_ROUTE_PREFORK=${DFLASH_DS4_TP_ROUTE_PREFORK:-1}
+export DFLASH_DS4_TP_DEVICE_JOIN_SPLIT=${DFLASH_DS4_TP_DEVICE_JOIN_SPLIT:-1}
+export DFLASH_DS4_TP_FUSED_HC_JOIN=${DFLASH_DS4_TP_FUSED_HC_JOIN:-1}
+export DFLASH_DS4_TP_MAIN_ROUTE_WEIGHTS=${DFLASH_DS4_TP_MAIN_ROUTE_WEIGHTS:-1}
+export DFLASH_DS4_TP_COARSE_OWNER=${DFLASH_DS4_TP_COARSE_OWNER:-1}
+export DFLASH_DS4_TP_COARSE_OWNER_SPLIT=${DFLASH_DS4_TP_COARSE_OWNER_SPLIT:-0}
+export GGML_BATCH_PEER_COPIES=${GGML_BATCH_PEER_COPIES:-1}
+export DFLASH_CUDA_MMVQ_MOE_ROWS_PER_BLOCK=${DFLASH_CUDA_MMVQ_MOE_ROWS_PER_BLOCK:-2}
+export DFLASH_DS4_TP_CAPTURE_CACHE_SLOTS=${DFLASH_DS4_TP_CAPTURE_CACHE_SLOTS:-4}
+export DFLASH_DS4_TP_FUSED_CACHE_SLOTS=${DFLASH_DS4_TP_FUSED_CACHE_SLOTS:-9}
+export DFLASH_DS4_VERIFY_FORCE_GRAPH_REPLAY=${DFLASH_DS4_VERIFY_FORCE_GRAPH_REPLAY:-1}
+export DFLASH_DS4_GPU_ARGMAX_VERIFY=${DFLASH_DS4_GPU_ARGMAX_VERIFY:-1}
+
+export DFLASH_DS4_SPEC=1
+export DFLASH_DS4_DRAFT=$draft
+export DFLASH_DS4_DRAFT_GPU=$draft_device
+export DFLASH_DS4_SPEC_Q=${DFLASH_DS4_SPEC_Q:-4}
+export DFLASH_DS4_PINNED_ROLLBACK=${DFLASH_DS4_PINNED_ROLLBACK:-1}
+export DFLASH_DS4_FUSED_VERIFY=${DFLASH_DS4_FUSED_VERIFY:-1}
+export DFLASH_DS4_TOPK=${DFLASH_DS4_TOPK:-6}
+
+exec "$binary" "$target" \
+    --target-device "$main_device" \
+    --peer-access \
+    --ds4-expert-top-k "$DFLASH_DS4_TOPK" \
+    --ds4-prefill sparse \
+    --chunk 2048 \
+    --max-ctx "$context" \
+    --prefix-cache-slots 0 \
+    --prefill-cache-slots 0 \
+    --host "${DS4_HOST:-127.0.0.1}" \
+    --port "$port" \
+    "$@"
diff --git a/server/src/common/moe_hybrid_ffn_eval.cpp b/server/src/common/moe_hybrid_ffn_eval.cpp
index b933036c0..ed2d2d4fd 100644
--- a/server/src/common/moe_hybrid_ffn_eval.cpp
+++ b/server/src/common/moe_hybrid_ffn_eval.cpp
@@ -151,6 +151,19 @@ int moe_balanced_main_slots_x4(int top_k, double main_to_peer_rate) {
     return total - peer;
 }
 
+MoeHybridOwnerMapView moe_hybrid_owner_maps(
+        const MoeHybridLayerStorage & storage,
+        bool dynamic_route_balance) {
+    return {
+        dynamic_route_balance || storage.decode_hot_local_by_global.empty()
+            ? &storage.hot_local_by_global
+            : &storage.decode_hot_local_by_global,
+        dynamic_route_balance || storage.decode_cold_local_by_global.empty()
+            ? &storage.cold_local_by_global
+            : &storage.decode_cold_local_by_global,
+    };
+}
+
 // The serial balanced-owner assignment is qualified for the q=5 DSpark
 // verifier. Wider batches retain the ordinary parallel owner remap.
 constexpr int kDynamicRouteBalanceMaxTokens = 5;
@@ -1176,19 +1189,6 @@ bool build_moe_hybrid_ffn_graph(
     // Expose the canonical tensor so the scheduler can keep it on the primary
     // backend rather than discovering it late through the secondary branch.
     out.router_weights = router_weights;
-    MoeOwnerGraphSpec primary_owner{
-        &storage.hot_local_by_global,
-        storage.gate_hot, storage.up_hot, storage.down_hot,
-        storage.gate_up_hot,
-        &out.hot_local_lut, &out.hot_valid_lut,
-        &out.hot_remap_nodes, &out.hot_nodes};
-    MoeOwnerGraphSpec secondary_owner{
-        &storage.cold_local_by_global,
-        storage.gate_cold, storage.up_cold, storage.down_cold,
-        storage.gate_up_cold,
-        &out.cold_local_lut, &out.cold_valid_lut,
-        &out.cold_remap_nodes, &out.cold_nodes};
-
     double derived_main_to_peer_rate = 0.0;
     int dynamic_main_slots_x4 =
         route_balance == MoeHybridRouteBalance::Allowed
@@ -1196,17 +1196,35 @@ bool build_moe_hybrid_ffn_graph(
                   n_used, &derived_main_to_peer_rate)
             : 0;
 
+    // Dynamic balancing routes by physical residency, not by the optional
+    // phase-specific decode placement. In particular, the decode peer map may
+    // intentionally omit experts assigned to the main owner, while the
+    // physical peer map still contains the complete fallback copy required by
+    // the dynamic complement.
+    const MoeHybridOwnerMapView physical_maps =
+        moe_hybrid_owner_maps(storage, true);
     const bool complete_secondary_map =
-        secondary_owner.local_by_global &&
-        (int) secondary_owner.local_by_global->size() == cfg.n_expert &&
+        physical_maps.peer &&
+        (int) physical_maps.peer->size() == cfg.n_expert &&
         std::none_of(
-            secondary_owner.local_by_global->begin(),
-            secondary_owner.local_by_global->end(),
+            physical_maps.peer->begin(), physical_maps.peer->end(),
             [](int32_t local) { return local < 0; });
+    const bool physical_primary_map_available =
+        physical_maps.main &&
+        (int) physical_maps.main->size() == cfg.n_expert &&
+        std::any_of(
+            physical_maps.main->begin(), physical_maps.main->end(),
+            [](int32_t local) { return local >= 0; });
+    const bool physical_primary_available = physical_primary_map_available &&
+        (storage.gate_up_hot || (storage.gate_hot && storage.up_hot)) &&
+        storage.down_hot;
+    const bool physical_secondary_available =
+        (storage.gate_up_cold || (storage.gate_cold && storage.up_cold)) &&
+        storage.down_cold;
     if (dynamic_main_slots_x4 > 0 &&
         (n_tokens > kDynamicRouteBalanceMaxTokens ||
-         !primary_owner.available() ||
-         !secondary_owner.available() || !complete_secondary_map)) {
+         !physical_primary_available ||
+         !physical_secondary_available || !complete_secondary_map)) {
         static std::once_flag fallback_log;
         std::call_once(fallback_log, [n_tokens] {
             std::fprintf(stderr,
@@ -1217,6 +1235,27 @@ bool build_moe_hybrid_ffn_graph(
         dynamic_main_slots_x4 = 0;
     }
     out.dynamic_route_balance = dynamic_main_slots_x4 > 0;
+
+    const MoeHybridOwnerMapView owner_maps =
+        moe_hybrid_owner_maps(storage, out.dynamic_route_balance);
+    const bool main_has_experts = owner_maps.main && std::any_of(
+        owner_maps.main->begin(), owner_maps.main->end(),
+        [](int32_t local) { return local >= 0; });
+    MoeOwnerGraphSpec primary_owner{
+        owner_maps.main,
+        main_has_experts ? storage.gate_hot : nullptr,
+        main_has_experts ? storage.up_hot : nullptr,
+        main_has_experts ? storage.down_hot : nullptr,
+        main_has_experts ? storage.gate_up_hot : nullptr,
+        &out.hot_local_lut, &out.hot_valid_lut,
+        &out.hot_remap_nodes, &out.hot_nodes};
+    MoeOwnerGraphSpec secondary_owner{
+        owner_maps.peer,
+        storage.gate_cold, storage.up_cold, storage.down_cold,
+        storage.gate_up_cold,
+        &out.cold_local_lut, &out.cold_valid_lut,
+        &out.cold_remap_nodes, &out.cold_nodes};
+
     if (dynamic_main_slots_x4 > 0) {
         static std::once_flag active_log;
         std::call_once(
@@ -2614,7 +2653,7 @@ static bool eval_moe_hybrid_remote_cold_batched(
 // reduced-stack MUL_MAT_ID stream-k path (which is both slow for a 24-expert
 // stack and unstable for very large batches) and lets each expert's weights be
 // reused across all of its prompt rows.
-static bool expert_major_prefill_enabled(int n_tokens) {
+bool moe_expert_major_prefill_enabled(int n_tokens) {
     static const bool enabled = []() {
         const char * raw = std::getenv("DFLASH_MOE_EXPERT_MAJOR_PREFILL");
         return !raw || !*raw || std::strcmp(raw, "0") != 0;
@@ -2638,6 +2677,17 @@ static bool expert_major_gpu_reduce_enabled() {
     return enabled;
 }
 
+
+static bool expert_major_pinned_output_enabled() {
+    static const bool enabled = []() {
+        const char * raw =
+            std::getenv("DFLASH_MOE_EXPERT_MAJOR_PINNED_OUTPUT");
+        return raw && *raw && std::strcmp(raw, "0") != 0;
+    }();
+    return enabled;
+}
+
+
 static bool full_cold_parallel_enabled() {
     static const bool enabled = []() {
         const char * raw =
@@ -2762,7 +2812,8 @@ static bool eval_moe_owner_expert_major_batched(
         }
     }
 
-    const bool gpu_reduce = expert_major_gpu_reduce_enabled() && n_pairs > 0;
+    const bool gpu_reduce =
+        expert_major_gpu_reduce_enabled() && n_pairs > 0;
 
     ggml_init_params ip{};
     ip.mem_size = 128 * 1024 * 1024;
@@ -2981,8 +3032,34 @@ static bool eval_moe_owner_expert_major_batched(
             backend, device_output_owner, combined_out, device_output);
         ggml_backend_synchronize(device_output_owner);
     } else if (gpu_reduce) {
-        ggml_backend_tensor_get(combined_out, out.data(), 0,
-                                sizeof(float) * out.size());
+        const size_t output_bytes = sizeof(float) * out.size();
+        if (expert_major_pinned_output_enabled()) {
+            ggml_backend_dev_t owner_device =
+                ggml_backend_get_device(backend);
+            ggml_backend_buffer_type_t host_buft = owner_device
+                ? ggml_backend_dev_host_buffer_type(owner_device)
+                : nullptr;
+            ggml_backend_buffer_t staging = host_buft
+                ? ggml_backend_buft_alloc_buffer(host_buft, output_bytes)
+                : nullptr;
+            if (!staging ||
+                ggml_backend_buffer_get_size(staging) < output_bytes) {
+                if (staging) ggml_backend_buffer_free(staging);
+                if (!p_alloc) ggml_gallocr_free(alloc);
+                ggml_free(ctx);
+                if (err) *err =
+                    "expert-major pinned output allocation failed";
+                return false;
+            }
+            void * staging_ptr = ggml_backend_buffer_get_base(staging);
+            ggml_backend_tensor_get(
+                combined_out, staging_ptr, 0, output_bytes);
+            std::memcpy(out.data(), staging_ptr, output_bytes);
+            ggml_backend_buffer_free(staging);
+        } else {
+            ggml_backend_tensor_get(combined_out, out.data(), 0,
+                                    output_bytes);
+        }
     } else {
         std::vector<float> packed_result(n_pairs * (size_t)n_embd);
         if (packed_out) {
@@ -3250,7 +3327,7 @@ bool eval_moe_hybrid_ffn_batched(
                             : 0;
     const bool cold_on_gpu = storage.cold_backend_kind == MoeHybridColdBackend::Gpu;
     const bool inprocess_expert_major =
-        !expert_compute && expert_major_prefill_enabled(n_tokens) &&
+        !expert_compute && moe_expert_major_prefill_enabled(n_tokens) &&
         cold_on_gpu && storage.cold_backend &&
         storage.cold_backend != gpu_backend &&
         n_hot_stack > 0 && n_cold_stack > 0 &&
@@ -3320,7 +3397,7 @@ bool eval_moe_hybrid_ffn_batched(
     // instead. The primary-local map takes priority in the core evaluator, so
     // skip_hot makes the duplicated secondary stack evaluate its routes only.
     const bool inprocess_full_cold_hot_expert_major =
-        !expert_compute && expert_major_prefill_enabled(n_tokens) &&
+        !expert_compute && moe_expert_major_prefill_enabled(n_tokens) &&
         cold_on_gpu && storage.cold_backend &&
         storage.cold_backend != gpu_backend &&
         n_hot_stack > 0 && n_cold_stack == cfg.n_expert;
diff --git a/server/src/common/moe_hybrid_ffn_eval.h b/server/src/common/moe_hybrid_ffn_eval.h
index 025f041a7..18e5613f0 100644
--- a/server/src/common/moe_hybrid_ffn_eval.h
+++ b/server/src/common/moe_hybrid_ffn_eval.h
@@ -18,6 +18,19 @@ namespace dflash::common {
 // estimated completion time. Returns zero for invalid inputs.
 int moe_balanced_main_slots_x4(int top_k, double main_to_peer_rate);
 
+// Select the phase-specific owner maps for ordinary routing, or the physical
+// residency maps required by dynamic route balancing. The peer physical map
+// must remain complete because it receives the exact complement of the
+// dynamically capped main-owner routes.
+struct MoeHybridOwnerMapView {
+    const std::vector<int32_t> * main = nullptr;
+    const std::vector<int32_t> * peer = nullptr;
+};
+
+MoeHybridOwnerMapView moe_hybrid_owner_maps(
+    const MoeHybridLayerStorage & storage,
+    bool dynamic_route_balance);
+
 // GPU-resident residual combine graph: output = residual + hot_out + cold_correction.
 struct ResidualCombineGraph {
     ggml_context * ctx = nullptr;
@@ -250,6 +263,8 @@ bool eval_moe_batched_prefill_ffn(
     int                             n_tokens,
     std::vector<float> &            out,
     std::string *                   err = nullptr);
+// Shared policy gate for paths that consume expert-major prefill outputs.
+bool moe_expert_major_prefill_enabled(int n_tokens);
 
 // Optional device-resident owner destinations for long heterogeneous prefill.
 // When present, the hot/shared and cold partials are copied directly into
diff --git a/server/src/common/moe_hybrid_placement.cpp b/server/src/common/moe_hybrid_placement.cpp
index f1525f755..967ad014c 100644
--- a/server/src/common/moe_hybrid_placement.cpp
+++ b/server/src/common/moe_hybrid_placement.cpp
@@ -462,6 +462,13 @@ bool MoeHybridPlacement::build_critical_path_balanced_from_stats(
 
     const int floor = std::clamp(
         config.min_hot_per_layer, 0, stats.n_expert);
+    const int ceiling = config.max_hot_per_layer > 0
+        ? std::clamp(config.max_hot_per_layer, 0, stats.n_expert)
+        : stats.n_expert;
+    if (ceiling < floor) {
+        if (err) *err = "max_hot_per_layer is smaller than min_hot_per_layer";
+        return false;
+    }
     uint64_t used_bytes = 0;
     for (int il = 0; il < stats.n_layer; ++il) {
         const uint64_t expert_bytes = layer_expert_bytes[(size_t) il];
@@ -544,7 +551,7 @@ bool MoeHybridPlacement::build_critical_path_balanced_from_stats(
         }
 
         double best_cost = std::numeric_limits<double>::infinity();
-        for (int hot_count = floor; hot_count <= tmp.n_expert; ++hot_count) {
+        for (int hot_count = floor; hot_count <= ceiling; ++hot_count) {
             const uint64_t extra_count = (uint64_t) (hot_count - floor);
             if (extra_count > 0 &&
                 expert_bytes > remaining / extra_count) {
@@ -642,5 +649,4 @@ bool MoeHybridPlacement::build_critical_path_balanced_from_stats(
     out = std::move(tmp);
     return true;
 }
-
 }  // namespace dflash::common
diff --git a/server/src/common/moe_hybrid_placement.h b/server/src/common/moe_hybrid_placement.h
index 9246f4032..f795c1e72 100644
--- a/server/src/common/moe_hybrid_placement.h
+++ b/server/src/common/moe_hybrid_placement.h
@@ -20,9 +20,11 @@ struct MoeHybridRoutingStats;  // forward decl
 struct MoeHybridCriticalPathConfig {
     int active_experts = 0;
     int min_hot_per_layer = 0;
+    // Zero means no additional ceiling. A finite ceiling is useful when the
+    // per-layer graph scratch grows faster than the resident-weight budget.
+    int max_hot_per_layer = 0;
     double main_to_peer_rate = 1.0;
 };
-
 inline uint64_t moe_hybrid_core_bytes_from_memory(const char * log_prefix,
                                                   size_t gpu_free,
                                                   size_t gpu_total) {
diff --git a/server/src/common/moe_hybrid_storage.cpp b/server/src/common/moe_hybrid_storage.cpp
index 57f6715cf..baa1adca5 100644
--- a/server/src/common/moe_hybrid_storage.cpp
+++ b/server/src/common/moe_hybrid_storage.cpp
@@ -175,6 +175,21 @@ static ggml_tensor * new_like_with_expert_count(ggml_context * ctx, ggml_tensor
 
 } // namespace
 
+void MoeHybridStorage::release_graph_caches() {
+    for (auto & layer : layers) {
+        layer.hot_graph.free();
+        layer.cold_graph.free();
+        for (auto & graph : layer.hot_graph_by_width) graph.free();
+        for (auto & graph : layer.cold_graph_by_width) graph.free();
+        layer.hot_graph_by_width.clear();
+        layer.cold_graph_by_width.clear();
+        layer.hot_batched_graph.free();
+        for (auto & graph : layer.hot_batched_mixed) graph.free();
+        for (auto & graph : layer.cold_batched_mixed) graph.free();
+        layer.shared_batched_graph.free();
+    }
+}
+
 MoeHybridStorage::~MoeHybridStorage() {
     // Registry entries point into the owner buffers, so remove them before the
     // buffers can be released or their addresses reused.
@@ -191,17 +206,8 @@ MoeHybridStorage::~MoeHybridStorage() {
         ggml_gallocr_free(prefill_cold_alloc);
         prefill_cold_alloc = nullptr;
     }
+    release_graph_caches();
     for (auto & layer : layers) {
-        layer.hot_graph.free();
-        layer.cold_graph.free();
-        for (auto & graph : layer.hot_graph_by_width) graph.free();
-        for (auto & graph : layer.cold_graph_by_width) graph.free();
-        layer.hot_graph_by_width.clear();
-        layer.cold_graph_by_width.clear();
-        layer.hot_batched_graph.free();
-        for (auto & g : layer.hot_batched_mixed) g.free();
-        for (auto & g : layer.cold_batched_mixed) g.free();
-        layer.shared_batched_graph.free();
         if (layer.hot_buf) {
             ggml_backend_buffer_free(layer.hot_buf);
             layer.hot_buf = nullptr;
@@ -331,6 +337,7 @@ bool build_moe_hybrid_storage(const MoeHybridConfig & cfg,
                 dst.cold_expert_ids.push_back((int32_t)expert);
             }
         }
+        dst.decode_cold_local_by_global = dst.cold_local_by_global;
 
         // Populate the model-sized VRAM bitmask from hot expert IDs.
         dst.reset_expert_vram_mask(cfg.n_expert);
@@ -547,6 +554,7 @@ bool build_moe_hybrid_storage_from_file(
                 }
             }
         }
+        dst.decode_cold_local_by_global = dst.cold_local_by_global;
 
         // Populate the model-sized VRAM bitmask from hot expert IDs.
         dst.reset_expert_vram_mask(cfg.n_expert);
diff --git a/server/src/common/moe_hybrid_storage.h b/server/src/common/moe_hybrid_storage.h
index b764ba48e..ab065175b 100644
--- a/server/src/common/moe_hybrid_storage.h
+++ b/server/src/common/moe_hybrid_storage.h
@@ -93,10 +93,11 @@ struct MoeHybridLayerStorage {
     std::vector<int32_t> hot_expert_ids;
     std::vector<int32_t> cold_expert_ids;
     std::vector<int32_t> hot_local_by_global;
-    // Optional decode ownership is a subset of the physically resident hot
-    // experts. Prefill can use every resident expert while decode retains a
-    // separately balanced fork. Empty means identical to hot_local_by_global.
+    // Optional decode ownership is independent of the physical maps used by
+    // prefill. Empty maps fall back to hot_local_by_global and
+    // cold_local_by_global respectively.
     std::vector<int32_t> decode_hot_local_by_global;
+    std::vector<int32_t> decode_cold_local_by_global;
     std::vector<int32_t> cold_local_by_global;
 
     // --- Bounded GPU expert cache (laguna) ---
@@ -226,6 +227,10 @@ struct MoeHybridStorage {
     bool matches(const MoeHybridConfig & cfg) const;
     bool empty() const;
     bool has_mmap() const { return mmap_data != nullptr && mmap_size > 0; }
+
+    // Decode/verify graph arenas are shape caches, not model state. Release
+    // them before a new bulk prefill needs substantially larger workspaces.
+    void release_graph_caches();
 };
 
 // Expert tensor file data for split loading (one entry per expert tensor).
diff --git a/server/src/deepseek4/deepseek4_backend.cpp b/server/src/deepseek4/deepseek4_backend.cpp
index 7183c4876..923fed82a 100644
--- a/server/src/deepseek4/deepseek4_backend.cpp
+++ b/server/src/deepseek4/deepseek4_backend.cpp
@@ -6,6 +6,7 @@
 #include "deepseek4_internal.h"
 #include "common/dynamic_backend.h"
 #include "common/peer_access.h"
+#include "common/platform_env.h"
 #include "common/sampler.h"
 
 #if defined(DFLASH27B_BACKEND_HIP) || defined(GGML_USE_HIP)
@@ -45,6 +46,68 @@ static bool env_flag_enabled(const char * name) {
     return value && value[0] && std::strcmp(value, "0") != 0;
 }
 
+struct AffineMmqPrefillScope {
+    bool active = false;
+
+    explicit AffineMmqPrefillScope(bool speculative_decode) {
+        // A dedicated phase-only switch matters here: setting the ordinary
+        // MMQ opt-in in the server environment exposes it while DSpark's
+        // persistent graphs are initialized, before this scope exists. Keep
+        // the master switch absent at startup and raise it only around the
+        // target prefill call.
+        static const bool configured =
+            env_flag_enabled("DFLASH_CUDA_MMQ_FP2_AFFINE_PREFILL_ONLY");
+        active = speculative_decode && configured;
+        set_enabled(true);
+    }
+
+    void set_enabled(bool enabled) {
+        if (!active) return;
+        if (enabled) {
+            set_environment_variable(
+                "DFLASH_CUDA_MMQ_FP2_AFFINE", "1", true);
+            unset_environment_variable(
+                "DFLASH_CUDA_MMQ_FP2_AFFINE_RUNTIME_DISABLE");
+        } else {
+            unset_environment_variable("DFLASH_CUDA_MMQ_FP2_AFFINE");
+            set_environment_variable(
+                "DFLASH_CUDA_MMQ_FP2_AFFINE_RUNTIME_DISABLE", "1", true);
+        }
+    }
+
+    ~AffineMmqPrefillScope() {
+        if (active) {
+            set_enabled(false);
+        }
+    }
+};
+
+struct PackedFp3DecodeScope {
+    bool active = false;
+
+    explicit PackedFp3DecodeScope(bool speculative_decode) {
+        // The packed FP3 MMVQ specialization is faster for q4 verification
+        // on gfx1151, but slower for the heterogeneous expert sub-batches
+        // used by bulk prefill. Keep the configured kernel available to the
+        // decoder while selecting the reference FP3 kernel during prefill.
+        active = speculative_decode &&
+            env_flag_enabled(
+                "DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24_DECODE_ONLY") &&
+            env_flag_enabled("DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24");
+        if (active) {
+            set_environment_variable(
+                "DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24_RUNTIME_DISABLE",
+                "1", true);
+        }
+    }
+
+    ~PackedFp3DecodeScope() {
+        if (active) {
+            unset_environment_variable(
+                "DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24_RUNTIME_DISABLE");
+        }
+    }
+};
 static bool positive_env_double(const char * name, double fallback,
                                 double & out, std::string * err) {
     out = fallback;
@@ -444,9 +507,16 @@ static uint64_t estimate_ds4_cache_bytes(const DeepSeek4Weights & w, int max_ctx
         total_bytes += window * head_dim * sizeof(float) * 2;
 
         if (ratio == 4) {
-            const size_t index_comp_width = (size_t) w.n_indexer_head * (size_t) w.n_indexer_head_dim;
-            total_bytes += comp_cap * index_comp_width * sizeof(uint16_t);
-            total_bytes += window * index_comp_width * sizeof(float) * 2;
+            // index_comp_kv is per-head: ne0 = n_indexer_head_dim (see
+            // deepseek4_graph.cpp). The full 64-head width lives only in the
+            // fixed-size state_kv/state_score scratch (index_state_rows rows),
+            // which does not scale with context. The old estimate multiplied
+            // by n_indexer_head here, overcounting 256K context by ~13x and
+            // falsely rejecting large contexts in hybrid placement.
+            total_bytes += comp_cap * (size_t) w.n_indexer_head_dim * sizeof(uint16_t);
+            total_bytes += window * (size_t) w.n_indexer_head_dim * sizeof(float) * 2;
+            total_bytes += (size_t) 2 * 2 * ratio * (size_t) w.n_indexer_head *
+                           (size_t) w.n_indexer_head_dim * sizeof(float);
         }
     }
 
@@ -584,6 +654,7 @@ static bool fill_profiled_hot_placement(const DeepSeek4Weights & w,
 static bool compute_ds4_hybrid_budget_info(const DeepSeek4Weights & w,
                                            ggml_backend_t backend,
                                            int max_ctx,
+                                           bool all_cold,
                                            Ds4HybridBudgetInfo & out,
                                            std::string * err) {
     out = {};
@@ -606,8 +677,11 @@ static bool compute_ds4_hybrid_budget_info(const DeepSeek4Weights & w,
         "deepseek4", out.gpu_free, out.gpu_total);
     out.kv_bytes = estimate_ds4_cache_bytes(w, max_ctx);
 
-    if (out.gpu_total > out.core_bytes + out.kv_bytes + out.warm_bytes + out.safety_bytes) {
-        out.expert_budget = out.gpu_total - out.core_bytes - out.kv_bytes - out.warm_bytes - out.safety_bytes;
+    // In all-cold mode the KV cache is owned by the secondary (Strix)
+    // backend, so it must not consume the primary GPU's expert budget.
+    const uint64_t main_charge = all_cold ? 0 : out.kv_bytes;
+    if (out.gpu_total > out.core_bytes + main_charge + out.warm_bytes + out.safety_bytes) {
+        out.expert_budget = out.gpu_total - out.core_bytes - main_charge - out.warm_bytes - out.safety_bytes;
     }
     if (out.expert_budget > out.mem.total_expert_bytes) {
         out.expert_budget = out.mem.total_expert_bytes;
@@ -1105,11 +1179,12 @@ bool DeepSeek4Backend::compute_uniform_hybrid_placement(const DeepSeek4Weights &
                                                        std::string * err) const {
     if (decode_out) *decode_out = {};
     Ds4HybridBudgetInfo budget;
-    if (!compute_ds4_hybrid_budget_info(w, backend_, max_ctx, budget, err)) {
+    const Ds4MoeTpConfig tp = ds4_moe_tp_config(cfg_.device.gpu);
+    if (!compute_ds4_hybrid_budget_info(w, backend_, max_ctx,
+                                        tp.all_on_secondary, budget, err)) {
         return false;
     }
 
-    const Ds4MoeTpConfig tp = ds4_moe_tp_config(cfg_.device.gpu);
     int hot_per_layer = tp.all_on_secondary ? 0 : budget.max_hot_per_layer;
     if (tp.all_on_secondary) {
         std::fprintf(stderr,
@@ -1188,6 +1263,11 @@ bool DeepSeek4Backend::compute_uniform_hybrid_placement(const DeepSeek4Weights &
                 0, w.n_expert, balance_cfg.min_hot_per_layer, err)) {
             return false;
         }
+        if (!env_int_in_range(
+                "DFLASH_DS4_TP_BALANCE_MAX_HOT", 0,
+                0, w.n_expert, balance_cfg.max_hot_per_layer, err)) {
+            return false;
+        }
 
         std::vector<uint64_t> main_fixed_bytes((size_t) w.n_layer, 0);
         for (int il = 0; il < w.n_layer; ++il) {
@@ -1239,9 +1319,10 @@ bool DeepSeek4Backend::compute_uniform_hybrid_placement(const DeepSeek4Weights &
             : (double) out.total_hot / (double) out.hot_counts.size();
         std::fprintf(stderr,
                      "[deepseek4] hybrid critical-path placement: "
-                     "profile=%s active=%d main/peer=%.3f "
+                     "profile=%s active=%d main/peer=%.3f cap=%d "
                      "hot/layer=%.1f [%d,%d]\n",
                      balance_profile_path, active_experts, main_to_peer_rate,
+                     balance_cfg.max_hot_per_layer,
                      mean_hot,
                      min_hot != out.hot_counts.end() ? *min_hot : 0,
                      max_hot != out.hot_counts.end() ? *max_hot : 0);
@@ -1485,6 +1566,9 @@ bool DeepSeek4Backend::init_hybrid_model() {
             MoeHybridLayerStorage & layer = hybrid->layers[(size_t) il];
             layer.decode_hot_local_by_global.assign(
                 (size_t) w_.n_expert, -1);
+            layer.decode_cold_local_by_global.assign(
+                (size_t) w_.n_expert, -1);
+            std::vector<uint8_t> decode_hot((size_t) w_.n_expert, 0);
             for (int32_t expert :
                     moe_decode_placement_.hot_expert_ids[(size_t) il]) {
                 if (expert < 0 || expert >= w_.n_expert ||
@@ -1495,12 +1579,54 @@ bool DeepSeek4Backend::init_hybrid_model() {
                                  (int) expert, il);
                     return fail_hybrid_init();
                 }
-                layer.decode_hot_local_by_global[(size_t) expert] =
+                decode_hot[(size_t) expert] = 1;
+            }
+            for (int expert = 0; expert < w_.n_expert; ++expert) {
+                const int32_t hot =
                     layer.hot_local_by_global[(size_t) expert];
+                const int32_t cold =
+                    layer.cold_local_by_global[(size_t) expert];
+                if (decode_hot[(size_t) expert]) {
+                    layer.decode_hot_local_by_global[(size_t) expert] = hot;
+                } else if (cold >= 0) {
+                    layer.decode_cold_local_by_global[(size_t) expert] = cold;
+                } else if (hot >= 0) {
+                    // A resident-only expert cannot move to the secondary
+                    // owner, so retain it on the primary during decode.
+                    layer.decode_hot_local_by_global[(size_t) expert] = hot;
+                } else {
+                    std::fprintf(stderr,
+                                 "[deepseek4] decode expert %d in layer %d "
+                                 "has no resident owner\n",
+                                 expert, il);
+                    return fail_hybrid_init();
+                }
             }
         }
     }
 
+    if (env_flag_enabled("DFLASH_DS4_DECODE_ALL_COLD")) {
+        for (int il = 0; il < w_.n_layer; ++il) {
+            MoeHybridLayerStorage & layer = hybrid->layers[(size_t) il];
+            if (layer.cold_local_by_global.size() !=
+                (size_t) w_.n_expert ||
+                std::any_of(layer.cold_local_by_global.begin(),
+                            layer.cold_local_by_global.end(),
+                            [](int32_t local) { return local < 0; })) {
+                std::fprintf(stderr,
+                             "[deepseek4] decode-all-cold requires a full "
+                             "secondary expert stack (layer=%d)\n", il);
+                return fail_hybrid_init();
+            }
+            layer.decode_hot_local_by_global.assign(
+                (size_t) w_.n_expert, -1);
+            layer.decode_cold_local_by_global =
+                layer.cold_local_by_global;
+        }
+        std::fprintf(stderr,
+                     "[deepseek4] speculative verifier routes all experts "
+                     "to the duplicated secondary stack\n");
+    }
     if (hybrid->has_mmap() && !hybrid->materialized_cold_experts) {
         size_t max_expert_bytes = 0;
         for (const auto & layer : hybrid->layers) {
@@ -1665,7 +1791,16 @@ int deepseek4_hybrid_prefill_chunk_tokens(
         int context_end,
         int current_cap) {
     constexpr int long_context_begin = 4096;
-    constexpr int long_context_chunk = 1024;
+    static const int long_context_chunk = [] {
+        const char * raw = std::getenv("DFLASH_DS4_LONG_CONTEXT_CHUNK");
+        if (!raw || !*raw) return 1024;
+        char * end = nullptr;
+        const long parsed = std::strtol(raw, &end, 10);
+        return end && end != raw && *end == '\0' && parsed > 0 &&
+                       parsed <= DS4_MAX_LAYER_MAJOR_PREFILL_TOKENS
+            ? (int) parsed
+            : 1024;
+    }();
     int bounded = std::max(1, requested_chunk);
     if (current_cap > 0) {
         bounded = std::min(bounded, current_cap);
@@ -1675,6 +1810,46 @@ int deepseek4_hybrid_prefill_chunk_tokens(
         : bounded;
 }
 
+int deepseek4_hybrid_prefill_step_tokens(
+        int configured_chunk,
+        int position,
+        int remaining_tokens) {
+    // Late-context pressure bound: below this position full-width chunks
+    // are used; at/above it chunks shrink to late_context_chunk. 32768 is
+    // the conservative default (the 2K attention arena needs ~1.19 GiB at
+    // ~80K and hits a fragmentation cliff on a 10.2 GB budget); operators
+    // with more headroom may raise it (DFLASH_DS4_LATE_CONTEXT_BEGIN).
+    static const int late_context_begin = [] {
+        const char * raw = std::getenv("DFLASH_DS4_LATE_CONTEXT_BEGIN");
+        if (!raw || !*raw) return 32768;
+        char * end = nullptr;
+        const long parsed = std::strtol(raw, &end, 10);
+        return end && end != raw && *end == '\0' && parsed > 0 &&
+                       parsed <= 262144
+            ? (int) parsed
+            : 32768;
+    }();
+    constexpr int late_context_chunk = 1024;
+    if (remaining_tokens <= 0) return 0;
+
+    int bounded = std::min(std::max(1, configured_chunk), remaining_tokens);
+    // Lower-residency expert placements have enough primary VRAM to retain a
+    // full 2K attention arena even at the context limit. Allow qualification
+    // runs for those placements to bypass the conservative pressure bound.
+    if (env_flag_enabled("DFLASH_DS4_DISABLE_ADAPTIVE_PREFILL")) {
+        return bounded;
+    }
+    // Stop exactly on the boundary so a non-aligned prefix or restored
+    // snapshot cannot carry one oversized attention arena into the late
+    // context region.
+    if (position < late_context_begin &&
+        position + bounded > late_context_begin) {
+        return late_context_begin - position;
+    }
+    return position >= late_context_begin
+        ? std::min(bounded, late_context_chunk)
+        : bounded;
+}
 int DeepSeek4Backend::do_prefill(const std::vector<int32_t> & tokens,
                                   const DaemonIO & io,
                                   int kv_offset,
@@ -1686,6 +1861,14 @@ int DeepSeek4Backend::do_prefill(const std::vector<int32_t> & tokens,
     const DeepSeek4RoctxRange roctx_range(
         "ds4.prefill",
         {phase, static_cast<int>(tokens.size()), 0, w_.n_layer, cfg_.device.gpu});
+    // The native affine MMQ loader is a large-prefill optimization. DSpark's
+    // fused verifier contains wider HC-expanded FP2 matmuls whose MMQ variant
+    // is not yet qualified. Keep the opt-in active throughout prefill, then
+    // disable it on every exit before speculative decode begins.
+    AffineMmqPrefillScope affine_mmq_scope(
+        spec_enabled_ && spec_drafter_ != nullptr);
+    PackedFp3DecodeScope packed_fp3_decode_scope(
+        spec_enabled_ && spec_drafter_ != nullptr);
     // The all-hot layer-range path supports causal chunked prefill. The
     // optimized graph snapshots the previous raw SWA window, attends over
     // that snapshot plus the current ubatch, and commits only the final SWA
@@ -1784,12 +1967,35 @@ int DeepSeek4Backend::do_prefill(const std::vector<int32_t> & tokens,
     const auto phase_t0 = Clock::now();
     DeepSeek4StepTelemetry tel_acc;
     int steps = 0;
+    bool capture_band_scratch_released = false;
 
     bool snapshot_saved = false;
+    bool late_context_chunk_logged = false;
     for (int i = 0; i < n_total;) {
         if (io.is_cancelled()) return pos;
 
-        int n_tok = std::min(chunk, n_total - i);
+        int n_tok = bound_hybrid_scratch
+            ? deepseek4_hybrid_prefill_step_tokens(chunk, pos, n_total - i)
+            : std::min(chunk, n_total - i);
+        if (!late_context_chunk_logged && n_tok < chunk &&
+            pos >= 32768 && n_total - i >= chunk) {
+            late_context_chunk_logged = true;
+            std::fprintf(stderr,
+                         "[deepseek4] late-context prefill pressure bound: "
+                         "chunk %d->%d at pos=%d\n",
+                         chunk, n_tok, pos);
+        }
+        // Keep the final heterogeneous band large enough for expert-major
+        // execution. A tiny (<512) remainder falls back to grouped
+        // mul_mat_id; the qualified affine MMQ path is deliberately disabled
+        // there and a full 256-expert duplicate stack makes that tail far more
+        // expensive than slightly rebalancing the preceding chunk.
+        const int tail_tokens = n_total - (i + n_tok);
+        if (moe_hybrid_ && spec_enabled_ && spec_drafter_ &&
+            tail_tokens > 0 && tail_tokens < 512 &&
+            n_tok >= 1024 - tail_tokens) {
+            n_tok -= 512 - tail_tokens;
+        }
         // A snapshot must represent an exact token boundary. Split a batched
         // prefill chunk when the requested boundary falls inside it.
         if (save_snapshot && !snapshot_saved &&
@@ -1806,6 +2012,24 @@ int DeepSeek4Backend::do_prefill(const std::vector<int32_t> & tokens,
                 spec_snap_from, spec_snap_to);
         }
 
+        // Bulk prompt graphs and the final DSpark feature-capture graph have
+        // different HC/owner arena shapes. Once all earlier chunks are
+        // complete, retire their reusable prefill arenas before entering the
+        // final capture band. Keeping both generations resident can cost more
+        // than 300 MiB and needlessly makes the highest-throughput expert
+        // placement fail only on the last chunk.
+        const bool entering_final_capture_band =
+            spec_enabled_ && spec_drafter_ && i > 0 &&
+            i + n_tok > spec_final_from;
+        if (entering_final_capture_band &&
+            !capture_band_scratch_released) {
+            deepseek4_release_prefill_scratch(cache_, moe_hybrid_.get());
+            capture_band_scratch_released = true;
+            std::fprintf(stderr,
+                         "[deepseek4] released bulk prefill scratch before "
+                         "final DSpark capture band\n");
+        }
+
         // Embed tokens
         std::vector<float> embed(w_.n_embd * n_tok);
         const auto embed_t0 = Clock::now();
@@ -1844,6 +2068,13 @@ int DeepSeek4Backend::do_prefill(const std::vector<int32_t> & tokens,
             spec_hooks.capture_token_end = capture_end;
             hp = &spec_hooks;
         }
+        // DSpark consumes the final SWA-width target features directly. Keep
+        // the high-throughput native affine MMQ on bulk chunks, but evaluate
+        // any feature-capture band through the established affine fallback so
+        // an experimental expert kernel cannot poison the speculative handoff.
+        static const bool affine_capture_enabled =
+            env_flag_enabled("DFLASH_CUDA_MMQ_FP2_AFFINE_CAPTURE");
+        affine_mmq_scope.set_enabled(hp == nullptr || affine_capture_enabled);
         if (moe_hybrid_ && (expert_runtime_.compute || expert_backend_)) {
             ok = deepseek4_step_layer_range(
                 backend_, cfg_.device.gpu, w_, cache_, hc_state,
@@ -1920,9 +2151,36 @@ int DeepSeek4Backend::do_prefill(const std::vector<int32_t> & tokens,
     }
     keep_spec_feature_tail(spec_feat_window_,
                            (size_t) std::max(0, w_.n_swa));
+    if (timing && spec_enabled_ && spec_drafter_ &&
+        !spec_feat_window_.empty()) {
+        size_t nonfinite = 0;
+        double square_sum = 0.0;
+        float min_value = std::numeric_limits<float>::infinity();
+        float max_value = -std::numeric_limits<float>::infinity();
+        for (float value : spec_feat_window_) {
+            if (!std::isfinite(value)) {
+                ++nonfinite;
+                continue;
+            }
+            square_sum += (double) value * (double) value;
+            min_value = std::min(min_value, value);
+            max_value = std::max(max_value, value);
+        }
+        const size_t finite = spec_feat_window_.size() - nonfinite;
+        std::fprintf(stderr,
+                     "[deepseek4] prefill DSpark feature tail: values=%zu "
+                     "nonfinite=%zu rms=%.4f min=%.4f max=%.4f\n",
+                     spec_feat_window_.size(), nonfinite,
+                     finite ? std::sqrt(square_sum / (double) finite) : 0.0,
+                     finite ? min_value : 0.0f,
+                     finite ? max_value : 0.0f);
+    }
     if (timing) {
         log_step_tel("prefill", n_total, steps, elapsed_s(phase_t0), tel_acc);
     }
+    if (spec_enabled_ && spec_drafter_) {
+        deepseek4_release_prefill_scratch(cache_, moe_hybrid_.get());
+    }
     return pos;
 }
 
@@ -2154,6 +2412,17 @@ GenerateResult DeepSeek4Backend::generate_from_state(
         int seed = 0;
         { float mv = last_logits_[0];
           for (int i = 1; i < w_.n_vocab; i++) if (last_logits_[i] > mv) { mv = last_logits_[i]; seed = i; } }
+        if (env_flag_enabled("DFLASH_DS4_TIMING")) {
+            size_t nonfinite_logits = 0;
+            for (float value : last_logits_) {
+                nonfinite_logits += !std::isfinite(value);
+            }
+            std::fprintf(stderr,
+                         "[deepseek4] speculative handoff: seed=%d "
+                         "seed_logit=%.6f nonfinite_logits=%zu\n",
+                         seed, last_logits_[(size_t) seed],
+                         nonfinite_logits);
+        }
         std::vector<int32_t> gen;
         gen.push_back(seed);
         out_io.emit(seed);
diff --git a/server/src/deepseek4/deepseek4_backend.h b/server/src/deepseek4/deepseek4_backend.h
index a9c58a18a..583132042 100644
--- a/server/src/deepseek4/deepseek4_backend.h
+++ b/server/src/deepseek4/deepseek4_backend.h
@@ -33,6 +33,14 @@ int deepseek4_hybrid_prefill_chunk_tokens(
     int context_end,
     int current_cap = 0);
 
+// Selects the next sparse heterogeneous prefill batch. Large batches retain
+// their throughput through the memory-light part of the prompt, then shrink
+// at the late-context boundary where one attention arena would otherwise
+// exhaust a tightly packed discrete GPU.
+int deepseek4_hybrid_prefill_step_tokens(
+    int configured_chunk,
+    int position,
+    int remaining_tokens);
 class DeepSeek4Backend : public ModelBackend {
 public:
     explicit DeepSeek4Backend(const DeepSeek4BackendConfig & cfg);
diff --git a/server/src/deepseek4/deepseek4_fused_verify.inc b/server/src/deepseek4/deepseek4_fused_verify.inc
index f831f2b68..ee270cc7d 100644
--- a/server/src/deepseek4/deepseek4_fused_verify.inc
+++ b/server/src/deepseek4/deepseek4_fused_verify.inc
@@ -167,22 +167,21 @@ static void ds4_fused_verify_refresh_hybrid_luts(
         const MoeHybridLayerStorage & layer = hybrid->layers[(size_t) il];
         const MoeHybridGraphInputs & inputs = fg.hybrid_inputs[hi];
         const bool dynamic_route_balance = inputs.dynamic_route_balance;
-        const std::vector<int32_t> & preferred_hot =
-            dynamic_route_balance ||
-                    layer.decode_hot_local_by_global.empty()
-                ? layer.hot_local_by_global
-                : layer.decode_hot_local_by_global;
+        const MoeHybridOwnerMapView owner_maps =
+            moe_hybrid_owner_maps(layer, dynamic_route_balance);
+        const std::vector<int32_t> & decode_hot = *owner_maps.main;
+        const std::vector<int32_t> & decode_cold = *owner_maps.peer;
         std::vector<int32_t> hot_lut((size_t) w.n_expert, invalid_route);
         std::vector<int32_t> cold_lut((size_t) w.n_expert, invalid_route);
         std::vector<float> hot_valid((size_t) w.n_expert, 0.0f);
         std::vector<float> cold_valid((size_t) w.n_expert, 0.0f);
         for (int ie = 0; ie < w.n_expert; ++ie) {
-            const int32_t cold = layer.cold_local_by_global[(size_t) ie];
-            int32_t hot = preferred_hot[(size_t) ie];
+            int32_t hot = decode_hot[(size_t) ie];
+            const int32_t cold = decode_cold[(size_t) ie];
             // A phase-specific placement may move a resident expert to the
             // peer only when the peer owns a copy. Otherwise retain its
             // physical main-owner route so no selected expert is dropped.
-            if (hot < 0 && cold < 0) {
+            if (!dynamic_route_balance && hot < 0 && cold < 0) {
                 hot = layer.hot_local_by_global[(size_t) ie];
             }
             if (hot >= 0) {
diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp
index f9c0c00b5..d8795846a 100644
--- a/server/src/deepseek4/deepseek4_graph.cpp
+++ b/server/src/deepseek4/deepseek4_graph.cpp
@@ -4821,6 +4821,45 @@ struct DeepSeek4LayerRangeCache {
                owns_output == candidate_owns_output;
     }
 
+    // A decoded response leaves q=1 graphs and tail-chunk prefill arenas in
+    // VRAM. Before the first large chunk of the next request, release those
+    // reproducible shape-specific allocations so the new HC/attention graphs
+    // do not have to coexist with them at peak memory.
+    void prepare_for_new_prefill() {
+        shared_prefill_attn_alloc.free();
+        prefill_hc_pre_graph.free();
+        prefill_hc_post_graph.free();
+        prefill_moe_hc_post_graph.free();
+        for (auto & graph : cached_decode_attn_hc_pre_graphs) {
+            graph.free();
+        }
+        for (auto & graph : cached_decode_ffn_hc_pre_graphs) {
+            graph.free();
+        }
+        cached_decode_hc_post_graph.free();
+        for (auto & per_layer : cached_decode_attn_graphs) {
+            for (auto & graph : per_layer) graph.free();
+        }
+        for (auto & graph : cached_decode_ffn_graphs) graph.free();
+        cached_decode_output_graph.free();
+        cached_dynamic_output_alloc.free();
+        fused_decode_graph_cache.evict_graphs();
+        fused_verify_graph_cache.destroy();
+        fused_capture_graph_cache.destroy();
+        decode_shared_inputs.free();
+    }
+
+    // Successful heterogeneous prefill no longer needs any batch-width graph
+    // arenas. Retire them before speculative verification constructs its
+    // q<=8 graphs; weights, HC mirrors, KV, and captured features are owned by
+    // other objects and intentionally remain resident.
+    void release_prefill_scratch() {
+        shared_prefill_attn_alloc.free();
+        prefill_hc_pre_graph.free();
+        prefill_hc_post_graph.free();
+        prefill_moe_hc_post_graph.free();
+    }
+
     void reset() {
         for (auto & alloc : cached_attn_allocs) {
             alloc.free();
@@ -5600,7 +5639,7 @@ static bool eval_ds4_layer_range_hybrid_ffn(
         std::strcmp(device_input_env, "0") != 0;
     const bool device_ffn_input =
         device_input_enabled &&
-        !expert_compute && n_tokens >= 512 &&
+        !expert_compute && moe_expert_major_prefill_enabled(n_tokens) &&
         layer_storage.cold_backend_kind == MoeHybridColdBackend::Gpu &&
         layer_storage.cold_backend && layer_storage.cold_backend != backend &&
         hot_stack_ref && hot_stack_ref->ne[2] > 0 &&
@@ -7016,6 +7055,44 @@ bool deepseek4_step_layer_range(
     auto & fused_decode_graph_cache = layer_range_cache.fused_decode_graph_cache;
     auto & decode_shared_inputs = layer_range_cache.decode_shared_inputs;
 
+    // Tiny request prefixes comfortably coexist with the verifier's cached
+    // q<=8 graphs. Evicting those graphs for every multi-token request turns
+    // steady-state speculative decode back into a cold graph-build path. The
+    // attention workspace starts its bulk growth regime at 512 tokens, so
+    // reserve the destructive cleanup for those genuinely large prefills.
+    constexpr int k_bulk_prefill_cleanup_tokens = 512;
+    if (heterogeneous_sparse_prefill && kv_start == 0 &&
+        n_tokens >= k_bulk_prefill_cleanup_tokens) {
+        ggml_backend_synchronize(backend);
+        if (moe_hybrid && moe_hybrid->cold_backend &&
+            moe_hybrid->cold_backend != backend) {
+            ggml_backend_synchronize(moe_hybrid->cold_backend);
+        }
+        layer_range_cache.prepare_for_new_prefill();
+        if (moe_hybrid) {
+            // Speculative verification also owns per-layer q<=8 MoE graph
+            // caches outside DeepSeek4LayerRangeCache. They can retain enough
+            // primary VRAM to make the next 2K-token attention arena fail,
+            // especially when the full cold verifier stack is duplicated.
+            moe_hybrid->release_graph_caches();
+            if (moe_hybrid->prefill_route_alloc) {
+                ggml_gallocr_free(moe_hybrid->prefill_route_alloc);
+                moe_hybrid->prefill_route_alloc = nullptr;
+            }
+            if (moe_hybrid->prefill_hot_alloc) {
+                ggml_gallocr_free(moe_hybrid->prefill_hot_alloc);
+                moe_hybrid->prefill_hot_alloc = nullptr;
+            }
+            if (moe_hybrid->prefill_cold_alloc) {
+                ggml_gallocr_free(moe_hybrid->prefill_cold_alloc);
+                moe_hybrid->prefill_cold_alloc = nullptr;
+            }
+        }
+        std::fprintf(stderr,
+                     "[deepseek4] released prior decode/tail arenas before "
+                     "new heterogeneous prefill\n");
+    }
+
     // Per-layer execution with CPU-side HC
     DeepSeek4LayerRangeScratch & scratch = layer_range_cache.scratch;
     const int n_expert_used = ds4_effective_expert_count(w);
@@ -7416,11 +7493,42 @@ bool deepseek4_step_layer_range(
                     if (!build_cached_decode_attn_graph(candidate, backend, w, L, lc, il, kv_start,
                                                         n_raw, n_comp_attn, n_index_comp,
                                                         shared_inputs)) {
-                        std::fprintf(stderr, "[deepseek4] cached attn graph alloc failed layer %d\n", il);
-                        return false;
+                        // Out of memory (tight primary GPU in split mode):
+                        // the decode attention graphs accumulate one ~16 MiB
+                        // entry per (layer, shape) as n_comp grows. Evict all
+                        // reproducible decode/prefill caches and retry once
+                        // before giving up.
+                        std::fprintf(stderr,
+                                     "[deepseek4] cached attn graph alloc failed layer %d; "
+                                     "evicting decode caches and retrying\n", il);
+                        ggml_backend_synchronize(backend);
+                        for (auto & other : cached_decode_attn_graphs) {
+                            for (auto & g : other) g.free();
+                            other.clear();
+                        }
+                        layer_range_cache.shared_prefill_attn_alloc.free();
+                        for (auto & alloc : layer_range_cache.cached_attn_allocs) {
+                            alloc.free();
+                        }
+                        layer_range_cache.fused_decode_graph_cache.evict_graphs();
+                        layer_range_cache.fused_verify_graph_cache.destroy();
+                        layer_range_cache.fused_capture_graph_cache.destroy();
+                        per_layer.emplace_back();
+                        auto & candidate2 = per_layer.back();
+                        if (!build_cached_decode_attn_graph(
+                                candidate2, backend, w, L, lc, il, kv_start,
+                                n_raw, n_comp_attn, n_index_comp,
+                                shared_inputs)) {
+                            std::fprintf(stderr,
+                                         "[deepseek4] cached attn graph alloc failed layer %d "
+                                         "after eviction\n", il);
+                            return false;
+                        }
+                        it = std::prev(per_layer.end());
+                    } else {
+                        it = std::prev(per_layer.end());
                     }
                     if (telemetry) telemetry->attn_build_us += ds4_elapsed_us(attn_build_t0, Ds4TimingClock::now());
-                    it = std::prev(per_layer.end());
                 }
                 cached_attn = &*it;
                 gf = cached_attn->sg.gf;
@@ -7536,8 +7644,43 @@ bool deepseek4_step_layer_range(
                         ggml_backend_cuda_get_device_memory(
                             device, &free_bytes, &total_bytes);
                         (void) total_bytes;
+                        // The route and hot-owner arenas belong to the
+                        // preceding layer at this point. Aggregate HIP free
+                        // memory can look sufficient while no individual
+                        // chunk is large enough for the replacement attention
+                        // arena, so retire these completed workspaces on every
+                        // growth instead of relying solely on the byte count.
+                        // They are recreated lazily by the FFN later in this
+                        // layer and contain no persistent model state.
+                        if (moe_hybrid &&
+                            (moe_hybrid->prefill_route_alloc ||
+                             moe_hybrid->prefill_hot_alloc)) {
+                            ggml_backend_synchronize(backend);
+                            if (moe_hybrid->cold_backend &&
+                                moe_hybrid->cold_backend != backend) {
+                                ggml_backend_synchronize(
+                                    moe_hybrid->cold_backend);
+                            }
+                            if (moe_hybrid->prefill_route_alloc) {
+                                ggml_gallocr_free(
+                                    moe_hybrid->prefill_route_alloc);
+                                moe_hybrid->prefill_route_alloc = nullptr;
+                            }
+                            if (moe_hybrid->prefill_hot_alloc) {
+                                ggml_gallocr_free(
+                                    moe_hybrid->prefill_hot_alloc);
+                                moe_hybrid->prefill_hot_alloc = nullptr;
+                            }
+                        }
+                        // HIP may split a ~650 MiB gallocr reservation into
+                        // several backend buffers. At tightly packed DS4
+                        // placements the nominal free-byte check can pass
+                        // while the largest follow-up chunk still fails due
+                        // to fragmentation and layer-local owner workspaces.
+                        // Evict those reproducible workspaces earlier so the
+                        // final DSpark capture band has a contiguous cushion.
                         constexpr size_t growth_margin =
-                            128u * 1024u * 1024u;
+                            384u * 1024u * 1024u;
                         const size_t replaceable_bytes =
                             free_bytes + attn_bytes_before;
                         if (replaceable_bytes < required_bytes + growth_margin) {
@@ -7558,6 +7701,22 @@ bool deepseek4_step_layer_range(
                             // their F16 HC mirrors are required by the active
                             // prefill and must remain resident.
                             layer_range_cache.fused_decode_graph_cache.evict_graphs();
+                            // The previous layer's routing and hot-owner
+                            // workspaces are also complete. On large chunks
+                            // they can otherwise prevent the shared attention
+                            // arena from growing by only a few dozen MiB.
+                            if (moe_hybrid) {
+                                if (moe_hybrid->prefill_route_alloc) {
+                                    ggml_gallocr_free(
+                                        moe_hybrid->prefill_route_alloc);
+                                    moe_hybrid->prefill_route_alloc = nullptr;
+                                }
+                                if (moe_hybrid->prefill_hot_alloc) {
+                                    ggml_gallocr_free(
+                                        moe_hybrid->prefill_hot_alloc);
+                                    moe_hybrid->prefill_hot_alloc = nullptr;
+                                }
+                            }
                             size_t free_after_evict = 0;
                             ggml_backend_cuda_get_device_memory(
                                 device, &free_after_evict, &total_bytes);
@@ -7572,6 +7731,30 @@ bool deepseek4_step_layer_range(
                                 free_bytes / (1024.0 * 1024.0),
                                 free_after_evict / (1024.0 * 1024.0));
                         }
+                        // ggml_gallocr grows its backend buffer by allocating
+                        // the replacement before releasing the current arena.
+                        // On tightly packed heterogeneous placements even a
+                        // tiny shape increase can therefore require roughly
+                        // twice the whole prefill scratch allocation. No node
+                        // from the previous layer/chunk remains live here, so
+                        // release the old arena first and recreate the
+                        // allocator before assigning this graph's tensors.
+                        if (attn_bytes_before > 0) {
+                            ggml_backend_synchronize(backend);
+                            attn_alloc.free();
+                            attn_alloc.alloc =
+                                ggml_gallocr_new_with_max_chunk_size(
+                                    ggml_backend_get_default_buffer_type(
+                                        backend),
+                                    shared_prefill_max_chunk);
+                            attn_alloc.owner_ctx = w.ctx;
+                            attn_alloc.backend = backend;
+                            std::fprintf(stderr,
+                                "[deepseek4] released %.1f MiB shared "
+                                "prefill scratch before %.1f MiB growth\n",
+                                attn_bytes_before / (1024.0 * 1024.0),
+                                required_bytes / (1024.0 * 1024.0));
+                        }
                     }
                 }
                 const bool attn_allocated = attn_alloc.alloc &&
@@ -7693,6 +7876,38 @@ bool deepseek4_step_layer_range(
             }
         }
 
+        // At long contexts the shared attention arena grows past the point
+        // where it can coexist with the primary-owner expert workspace on a
+        // tightly packed discrete GPU. The attention result has already been
+        // copied into, and consumed by, the persistent HC-post graph above;
+        // no attention tensor remains live for the FFN. Retire only these
+        // large arenas here. The next layer recreates one after releasing the
+        // preceding hot-owner workspace, keeping the normal <=12K fast path
+        // allocation-free across layers.
+        constexpr int k_long_context_attn_release_pos = 14336;
+        constexpr size_t k_long_context_attn_release_bytes =
+            384u * 1024u * 1024u;
+        if (!ds4_env_flag(
+                "DFLASH_DS4_DISABLE_LONG_CONTEXT_ARENA_HANDOFF") &&
+            heterogeneous_sparse_prefill &&
+            kv_start >= k_long_context_attn_release_pos &&
+            shared_prefill_attn_alloc.alloc &&
+            ggml_gallocr_get_buffer_size(
+                shared_prefill_attn_alloc.alloc, 0) >=
+                k_long_context_attn_release_bytes) {
+            const size_t released_bytes = ggml_gallocr_get_buffer_size(
+                shared_prefill_attn_alloc.alloc, 0);
+            ggml_backend_synchronize(backend);
+            shared_prefill_attn_alloc.free();
+            static std::atomic<bool> logged_long_context_attn_release{false};
+            if (!logged_long_context_attn_release.exchange(true)) {
+                std::fprintf(stderr,
+                    "[deepseek4] long-context attention/FFN arena handoff "
+                    "active (released %.1f MiB)\n",
+                    released_bytes / (1024.0 * 1024.0));
+            }
+        }
+
         // ── HC pre (FFN) ────────────────────────────────────────────
         const auto hc_pre_ffn_t0 = Ds4TimingClock::now();
         if (trace_prefill) {
@@ -7819,11 +8034,39 @@ bool deepseek4_step_layer_range(
                     : layer_storage.gate_cold;
                 const bool local_expert_runtime =
                     !expert_runtime || !expert_runtime->compute_ptr();
+                // The device-resident owner join is only populated when the
+                // FFN actually writes hot/cold outputs to device tensors
+                // (eval_ds4_layer_range_hybrid_ffn forwards device_outputs
+                // only when its device_ffn_input conditions hold). Mirror
+                // those conditions here: otherwise the moe HC-post graph
+                // reads never-written block_out/block_out_cold and the HC
+                // state becomes garbage.
+                ggml_tensor * hot_stack = layer_storage.gate_up_hot
+                    ? layer_storage.gate_up_hot
+                    : layer_storage.gate_hot;
+                const char * device_input_env =
+                    std::getenv("DFLASH_MOE_PREFILL_DEVICE_INPUT");
+                const bool device_input_enabled =
+                    !device_input_env || !*device_input_env ||
+                    std::strcmp(device_input_env, "0") != 0;
+                const bool ffn_device_join_possible =
+                    device_input_enabled &&
+                    moe_expert_major_prefill_enabled(n_tokens) &&
+                    layer_storage.cold_backend_kind ==
+                        MoeHybridColdBackend::Gpu &&
+                    layer_storage.cold_backend &&
+                    layer_storage.cold_backend != backend &&
+                    hot_stack && hot_stack->ne[2] > 0 &&
+                    // Only the full-secondary owner path publishes hot/cold
+                    // results into the device join tensors. A genuine split
+                    // uses the expert-major host-combine path; treating it as
+                    // device-resident makes HC-post read stale tensors.
+                    cold_stack && cold_stack->ne[2] == w.n_expert;
                 ffn_device_join =
                     use_backend_prefill_hc && ffn_in_backend &&
                     local_expert_runtime &&
                     prefill_moe_hc_post_graph.valid() &&
-                    cold_stack && cold_stack->ne[2] == w.n_expert;
+                    ffn_device_join_possible;
                 if (ffn_device_join) {
                     static bool logged_device_join = false;
                     if (!logged_device_join) {
@@ -8228,6 +8471,60 @@ void reset_deepseek4_cache(DeepSeek4Cache & c) {
     }
 }
 
+void deepseek4_release_prefill_scratch(
+        DeepSeek4Cache & c,
+        MoeHybridStorage * moe_hybrid) {
+    DeepSeek4LayerRangeCache * runtime = c.layer_range_cache;
+    ggml_backend_t primary_backend = runtime ? runtime->backend : nullptr;
+    ggml_backend_t cold_backend = moe_hybrid ? moe_hybrid->cold_backend : nullptr;
+
+    // The last owner/attention graph can still be executing asynchronously.
+    // All arenas below are backend allocations, so synchronize both owners
+    // before destroying either side.
+    if (primary_backend) {
+        ggml_backend_synchronize(primary_backend);
+    }
+    if (cold_backend && cold_backend != primary_backend) {
+        ggml_backend_synchronize(cold_backend);
+    }
+
+    size_t free_before = 0;
+    size_t total_bytes = 0;
+    if (runtime && runtime->device >= 0) {
+        ggml_backend_cuda_get_device_memory(
+            runtime->device, &free_before, &total_bytes);
+    }
+
+    if (runtime) {
+        runtime->release_prefill_scratch();
+    }
+    if (moe_hybrid) {
+        if (moe_hybrid->prefill_route_alloc) {
+            ggml_gallocr_free(moe_hybrid->prefill_route_alloc);
+            moe_hybrid->prefill_route_alloc = nullptr;
+        }
+        if (moe_hybrid->prefill_hot_alloc) {
+            ggml_gallocr_free(moe_hybrid->prefill_hot_alloc);
+            moe_hybrid->prefill_hot_alloc = nullptr;
+        }
+        if (moe_hybrid->prefill_cold_alloc) {
+            ggml_gallocr_free(moe_hybrid->prefill_cold_alloc);
+            moe_hybrid->prefill_cold_alloc = nullptr;
+        }
+    }
+
+    if (runtime && runtime->device >= 0) {
+        size_t free_after = 0;
+        ggml_backend_cuda_get_device_memory(
+            runtime->device, &free_after, &total_bytes);
+        std::fprintf(stderr,
+                     "[deepseek4] released prefill scratch before decode: "
+                     "primary free %.1f->%.1f MiB\n",
+                     free_before / (1024.0 * 1024.0),
+                     free_after / (1024.0 * 1024.0));
+    }
+}
+
 namespace {
 
 ggml_tensor * clone_snapshot_tensor(ggml_context * ctx,
diff --git a/server/src/deepseek4/deepseek4_internal.h b/server/src/deepseek4/deepseek4_internal.h
index eea667423..268ee4101 100644
--- a/server/src/deepseek4/deepseek4_internal.h
+++ b/server/src/deepseek4/deepseek4_internal.h
@@ -339,6 +339,10 @@ bool create_deepseek4_cache(ggml_backend_t backend,
 
 void free_deepseek4_cache(DeepSeek4Cache & c);
 void reset_deepseek4_cache(DeepSeek4Cache & c);
+// Release only reproducible large-batch graph arenas after prefill. KV/model
+// state and the DSpark feature tail remain live for the following decode.
+void deepseek4_release_prefill_scratch(DeepSeek4Cache & c,
+                                       MoeHybridStorage * moe_hybrid);
 int deepseek4_previous_raw_ring_spans(
     int kv_start,
     int n_swa,
diff --git a/server/src/server/http_server.cpp b/server/src/server/http_server.cpp
index 9d6f9f12d..d6c629063 100644
--- a/server/src/server/http_server.cpp
+++ b/server/src/server/http_server.cpp
@@ -511,14 +511,22 @@ static size_t json_array_size(const json & value) {
 }
 
 int resolve_max_output_tokens(const json & body, int default_max_tokens) {
+    // OpenAI-compatible clients (e.g. PocketPal's "Unlimited") send
+    // max_completion_tokens: -1 to mean "no explicit limit"; 0 is also
+    // invalid as a budget. Treat non-positive values as unset so they
+    // fall back to the server default instead of yielding zero tokens.
+    auto field_or_default = [&](const char * key) {
+        const int value = body.at(key).get<int>();
+        return value > 0 ? value : default_max_tokens;
+    };
     if (body.contains("max_tokens")) {
-        return body.at("max_tokens").get<int>();
+        return field_or_default("max_tokens");
     }
     if (body.contains("max_output_tokens")) {
-        return body.at("max_output_tokens").get<int>();
+        return field_or_default("max_output_tokens");
     }
     if (body.contains("max_completion_tokens")) {
-        return body.at("max_completion_tokens").get<int>();
+        return field_or_default("max_completion_tokens");
     }
     return default_max_tokens;
 }
@@ -2245,8 +2253,9 @@ json build_openai_completion_response(
     // usage.completion_tokens_details.reasoning_tokens — OpenAI o1/o3
     // standard location; kept in sync with finish_details.thinking_tokens.
     // usage.timings — per-request prefill/decode wall clock, additive to
-    // the OpenAI shape (ignored by clients that don't recognize it). See
-    // docs/specs/thinking-budget.md §6.3.
+    // the OpenAI shape (ignored by clients that don't recognize it).
+    // spec_decode_ran disambiguates a real zero-acceptance speculative run
+    // from an autoregressive fallback. See docs/specs/thinking-budget.md §6.3.
     const int prompt_tokens = (int) req.prompt_tokens.size();
     const json usage = {
         {"prompt_tokens", prompt_tokens},
@@ -2257,6 +2266,7 @@ json build_openai_completion_response(
         }},
         {"timings", build_timings_json(timings, counts.total)},
         {"accept_rate", result.accept_rate},
+        {"spec_decode_ran", result.spec_decode_ran},
     };
     return {
         {"id", req.response_id},
@@ -2324,6 +2334,7 @@ json build_anthropic_response(
         {"output_tokens", counts.total},
         {"timings", build_timings_json(timings, counts.total)},
         {"accept_rate", result.accept_rate},
+        {"spec_decode_ran", result.spec_decode_ran},
     };
     return {
         {"id", req.response_id},
@@ -2373,6 +2384,7 @@ json build_responses_api_response(
         {"total_tokens", prompt_tokens + counts.total},
         {"timings", build_timings_json(timings, counts.total)},
         {"accept_rate", result.accept_rate},
+        {"spec_decode_ran", result.spec_decode_ran},
     };
     return {
         {"id", req.response_id},
diff --git a/server/test/test_deepseek4_mmid_grouped_cuda.cpp b/server/test/test_deepseek4_mmid_grouped_cuda.cpp
index 975b772fd..0a555d0ba 100644
--- a/server/test/test_deepseek4_mmid_grouped_cuda.cpp
+++ b/server/test/test_deepseek4_mmid_grouped_cuda.cpp
@@ -2,6 +2,7 @@
 #include "ggml-backend.h"
 #include "ggml-cuda.h"
 #include "ggml.h"
+#include "rocmfpx.h"
 
 #include <cuda_runtime.h>
 
@@ -81,8 +82,18 @@ static bool run_case(
     }
 
     std::vector<uint8_t> weights_q(ggml_nbytes(weights));
-    const size_t quantized = ggml_quantize_chunk(
-        type, weights_f.data(), weights_q.data(), 0, n_rows * n_experts, k_dim, nullptr);
+    const size_t quantized =
+        type == GGML_TYPE_Q2_0_ROCMFP2
+            ? rocmfpx_quantize_fp2(
+                  weights_f.data(), weights_q.data(), n_rows * n_experts,
+                  k_dim, nullptr)
+        : type == GGML_TYPE_Q3_0_ROCMFPX
+            ? rocmfpx_quantize_fp3(
+                  weights_f.data(), weights_q.data(), n_rows * n_experts,
+                  k_dim, nullptr)
+            : ggml_quantize_chunk(
+                  type, weights_f.data(), weights_q.data(), 0,
+                  n_rows * n_experts, k_dim, nullptr);
     if (quantized != weights_q.size()) {
         std::fprintf(stderr, "quantize size mismatch type=%s got=%zu expected=%zu\n",
                      ggml_type_name(type), quantized, weights_q.size());
@@ -152,7 +163,8 @@ static int run_child(const char * mode, const char * output_path) {
 
     std::ofstream output(output_path, std::ios::binary | std::ios::trunc);
     const ggml_type types[] = {
-        GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q5_K,
+        GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0,
+        GGML_TYPE_Q5_K, GGML_TYPE_Q2_0_ROCMFP2, GGML_TYPE_Q3_0_ROCMFPX,
     };
     const int widths[] = {2, 4, 8, 9, 16};
     bool ok = output.good();
@@ -294,13 +306,17 @@ static std::string child_command(
         const std::string & output_path,
         const std::string & log_path) {
 #if defined(_WIN32)
-    return "set \"DFLASH_MMID_TELEMETRY=1\" && set \"DFLASH_MMID_GROUPED_TYPES=7\" && "
+    return "set \"DFLASH_MMID_TELEMETRY=1\" && set \"DFLASH_MMID_GROUPED_TYPES=15\" && "
+        "set \"DFLASH_CUDA_MMVQ_MOE_FP2_PACKED32=1\" && "
+        "set \"DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24=1\" && "
         "set \"DFLASH_MMID_GROUPED=" +
         std::string(std::strcmp(mode, "grouped") == 0 ? "1" : "0") + "\" && " +
         shell_quote(executable) + " --child " + mode + " " + shell_quote(output_path) +
         " 2>" + shell_quote(log_path);
 #else
-    return "DFLASH_MMID_TELEMETRY=1 DFLASH_MMID_GROUPED_TYPES=7 DFLASH_MMID_GROUPED=" +
+    return "DFLASH_MMID_TELEMETRY=1 DFLASH_MMID_GROUPED_TYPES=15 "
+        "DFLASH_CUDA_MMVQ_MOE_FP2_PACKED32=1 DFLASH_CUDA_MMVQ_MOE_FP3_PACKED24=1 "
+        "DFLASH_MMID_GROUPED=" +
         std::string(std::strcmp(mode, "grouped") == 0 ? "1" : "0") + " " +
         shell_quote(executable) + " --child " + mode + " " + shell_quote(output_path) +
         " 2>" + shell_quote(log_path);
@@ -347,7 +363,8 @@ int main(int argc, char ** argv) {
     const size_t grouped_grouped = count_records(grouped_log, "variant=grouped");
 
     const ggml_type types[] = {
-        GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q5_K,
+        GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0,
+        GGML_TYPE_Q5_K, GGML_TYPE_Q2_0_ROCMFP2, GGML_TYPE_Q3_0_ROCMFPX,
     };
     const int widths[] = {2, 4, 8, 9, 16};
     size_t offset = 0;
@@ -383,9 +400,9 @@ int main(int argc, char ** argv) {
             }
         }
     }
-    output_parity = output_parity && offset == legacy.size() && compared_cases == 50;
+    output_parity = output_parity && offset == legacy.size() && compared_cases == 70;
     const bool pass = legacy_status == 0 && grouped_status == 0 &&
-        output_parity && grouped_dispatch && legacy_grouped == 0 && grouped_grouped == 75;
+        output_parity && grouped_dispatch && legacy_grouped == 0 && grouped_grouped == 105;
     if (pass) {
         std::remove(legacy_path.c_str());
         std::remove(grouped_path.c_str());
diff --git a/server/test/test_moe_hybrid_storage.cpp b/server/test/test_moe_hybrid_storage.cpp
index 4e7a1d136..a0d98c58b 100644
--- a/server/test/test_moe_hybrid_storage.cpp
+++ b/server/test/test_moe_hybrid_storage.cpp
@@ -2,6 +2,7 @@
 #include "../src/common/moe_hybrid_ffn_eval.h"
 #include "../src/common/moe_hybrid_storage.h"
 
+#include <algorithm>
 #include <cstdint>
 #include <vector>
 
@@ -44,6 +45,27 @@ TEST_CASE(MoeHybridStorageFixture, heterogeneous_route_balance_scales_with_model
     REQUIRE(moe_balanced_main_slots_x4(6, 0.0) == 0);
 }
 
+TEST_CASE(MoeHybridStorageFixture, dynamic_route_balance_uses_physical_owner_maps) {
+    MoeHybridLayerStorage storage;
+    storage.hot_local_by_global = {0, -1, 1, -1};
+    storage.cold_local_by_global = {0, 1, 2, 3};
+    storage.decode_hot_local_by_global = {-1, -1, 1, -1};
+    storage.decode_cold_local_by_global = {0, 1, -1, 3};
+
+    const MoeHybridOwnerMapView static_maps =
+        moe_hybrid_owner_maps(storage, false);
+    REQUIRE(static_maps.main == &storage.decode_hot_local_by_global);
+    REQUIRE(static_maps.peer == &storage.decode_cold_local_by_global);
+
+    const MoeHybridOwnerMapView dynamic_maps =
+        moe_hybrid_owner_maps(storage, true);
+    REQUIRE(dynamic_maps.main == &storage.hot_local_by_global);
+    REQUIRE(dynamic_maps.peer == &storage.cold_local_by_global);
+    REQUIRE(std::none_of(
+        dynamic_maps.peer->begin(), dynamic_maps.peer->end(),
+        [](int32_t local) { return local < 0; }));
+}
+
 TEST_CASE(MoeHybridStorageFixture, fractional_route_quota_rounds_over_the_batch) {
     ggml_init_params params{
         /*mem_size=*/1024 * 1024,
diff --git a/server/test/test_qwen35moe_expert_placement.cpp b/server/test/test_qwen35moe_expert_placement.cpp
index 9ec6b45e3..4f50db97f 100644
--- a/server/test/test_qwen35moe_expert_placement.cpp
+++ b/server/test/test_qwen35moe_expert_placement.cpp
@@ -71,7 +71,6 @@ TEST_CASE(Qwen35MoeExpertPlacementFixture, moe_expert_placement_suite) {
     REQUIRE(!MoeHybridPlacement::load_json(tmp.string(), loaded, &err));
     REQUIRE(err == "placement contains a duplicate expert at layer 0");
     std::filesystem::remove(tmp);
-
     // Aggregate hit-rate placement can overfeed a highly skewed layer while a
     // flat layer remains peer-bound. The critical-path model stops at each
     // layer's branch crossover and may deliberately leave spare memory unused.
@@ -127,6 +126,14 @@ TEST_CASE(Qwen35MoeExpertPlacementFixture, moe_expert_placement_suite) {
     REQUIRE(unequal_size.hot_counts == std::vector<int>({1, 0}));
     REQUIRE(unequal_size.is_hot(0, 0));
 
+    balance_cfg.max_hot_per_layer = 2;
+    MoeHybridPlacement capped;
+    REQUIRE(MoeHybridPlacement::build_critical_path_balanced_from_stats(
+        balance_stats, {100, 100}, {100, 100}, 600,
+        balance_cfg, capped, &err));
+    REQUIRE(capped.hot_counts == std::vector<int>({2, 1}));
+    balance_cfg.max_hot_per_layer = 0;
+
     // A phase-specific decode placement remains intact while a second profile
     // spends otherwise-unused bytes on prefill residency.
     MoeHybridRoutingStats residency_stats;
@@ -181,11 +188,9 @@ TEST_CASE(Qwen35MoeExpertPlacementFixture, moe_expert_placement_suite) {
         overflow_csv.string(), loaded_overflow, &err));
     REQUIRE(err == "routing profile count overflow at layer 0");
     std::filesystem::remove(overflow_csv);
-
     balance_cfg.main_to_peer_rate = 0.0;
     REQUIRE(!MoeHybridPlacement::build_critical_path_balanced_from_stats(
         balance_stats, {100, 100}, {100, 100}, 600,
         balance_cfg, balanced, &err));
-
     std::printf("OK\n");
 }
diff --git a/server/test/test_rocmfpx_mmq.cpp b/server/test/test_rocmfpx_mmq.cpp
index baa77d8db..fd2e92f8a 100644
--- a/server/test/test_rocmfpx_mmq.cpp
+++ b/server/test/test_rocmfpx_mmq.cpp
@@ -31,6 +31,15 @@ struct Shape {
     const char * label;
 };
 
+enum class DispatchPath {
+    MMVQ,
+    MMQ,
+};
+
+static const char * dispatch_path_name(DispatchPath path) {
+    return path == DispatchPath::MMVQ ? "MMVQ" : "MMQ";
+}
+
 static std::vector<float> make_values(size_t count, int stride, float scale) {
     std::vector<float> values(count);
     for (size_t i = 0; i < count; ++i) {
@@ -58,7 +67,8 @@ static bool run_backend(
         const Shape & shape,
         const std::vector<uint8_t> & weights_data,
         const std::vector<float> & input_data,
-        std::vector<float> & output_data) {
+        std::vector<float> & output_data,
+        DispatchPath expected_path) {
     ggml_init_params params{};
     params.mem_size = 16 * 1024 * 1024;
     params.no_alloc = true;
@@ -98,9 +108,38 @@ static bool run_backend(
     if (ok) {
         ggml_backend_tensor_set(weights, weights_data.data(), 0, weights_data.size());
         ggml_backend_tensor_set(input, input_data.data(), 0, input_data.size() * sizeof(float));
+
+        const size_t mmvq_before = ggml_backend_cuda_get_mmvq_launch_count();
+        const size_t mmq_before = ggml_backend_cuda_get_mmq_launch_count();
+        const int previous_mmvq_max =
+            ggml_backend_cuda_set_mmvq_max_ncols_override(
+                expected_path == DispatchPath::MMVQ ? 8 : 1);
+        const bool previous_graphs_disabled =
+            ggml_backend_cuda_set_graphs_disabled_override(true);
         ok = ggml_backend_graph_compute(backend, graph) == GGML_STATUS_SUCCESS;
+        ggml_backend_cuda_set_graphs_disabled_override(previous_graphs_disabled);
+        ggml_backend_cuda_set_mmvq_max_ncols_override(previous_mmvq_max);
+        const size_t mmvq_delta =
+            ggml_backend_cuda_get_mmvq_launch_count() - mmvq_before;
+        const size_t mmq_delta =
+            ggml_backend_cuda_get_mmq_launch_count() - mmq_before;
         if (!ok) {
             std::fprintf(stderr, "backend graph compute failed\n");
+        } else {
+            const bool dispatch_matches =
+                expected_path == DispatchPath::MMVQ
+                    ? mmvq_delta == 1 && mmq_delta == 0
+                    : mmvq_delta == 0 && mmq_delta == 1;
+            if (!dispatch_matches) {
+                std::fprintf(
+                    stderr,
+                    "%s: expected %s dispatch, observed MMVQ=%zu MMQ=%zu\n",
+                    shape.label,
+                    dispatch_path_name(expected_path),
+                    mmvq_delta,
+                    mmq_delta);
+                ok = false;
+            }
         }
     }
 
@@ -178,7 +217,8 @@ static bool compare_outputs(
 static bool test_case(
         ggml_backend_t hip_backend,
         const QuantCase & quant,
-        const Shape & shape) {
+        const Shape & shape,
+        DispatchPath expected_path) {
     const std::vector<float> weights_f32 =
         make_values((size_t) shape.k * shape.m, 37, 0.015625f);
     const std::vector<float> input_f32 =
@@ -207,26 +247,35 @@ static bool test_case(
         reference_mul_mat(quant, shape, weights_quantized, input_f32);
     std::vector<float> actual;
     if (!run_backend(
-            hip_backend, quant.type, shape, weights_quantized, input_f32, actual)) {
-        std::fprintf(stderr, "%s/%s: HIP MMQ run failed\n", quant.label, shape.label);
+            hip_backend, quant.type, shape, weights_quantized, input_f32, actual,
+            expected_path)) {
+        std::fprintf(
+            stderr,
+            "%s/%s: HIP %s run failed\n",
+            quant.label,
+            shape.label,
+            dispatch_path_name(expected_path));
         return false;
     }
     return compare_outputs(quant, shape, expected, actual);
 }
 
 int main() {
+    setenv("DFLASH_CUDA_MMVQ_FP2_AFFINE", "1", 1);
+    setenv("DFLASH_CUDA_MMQ_FP2_AFFINE", "1", 1);
+    setenv("DFLASH_CUDA_MMQ_FP2_AFFINE_GENERAL", "1", 1);
     hipDeviceProp_t properties{};
     if (hipGetDeviceProperties(&properties, 0) != hipSuccess) {
         std::fprintf(stderr, "failed to query HIP device 0\n");
         return 1;
     }
-    if (std::strncmp(properties.gcnArchName, "gfx1151", 7) != 0) {
-        std::printf("SKIP: ROCmFPX MMQ is gfx1151-only (found %s)\n", properties.gcnArchName);
+    if (std::strncmp(properties.gcnArchName, "gfx1151", 7) != 0 &&
+        std::strncmp(properties.gcnArchName, "gfx12", 5) != 0) {
+        std::printf("SKIP: ROCmFPX dispatch test expects gfx1151/gfx12xx (found %s)\n",
+                    properties.gcnArchName);
         return 0;
     }
 
-    setenv("LUCE_MMVQ_MAX_NCOLS", "1", 1);
-
     ggml_backend_t hip_backend = ggml_backend_cuda_init(0);
     if (!hip_backend) {
         std::fprintf(stderr, "failed to initialize HIP backend\n");
@@ -242,14 +291,41 @@ int main() {
     const Shape shapes[] = {
         {256, 64, 64, "full"},
         {288, 70, 67, "tail"},
+        {4096, 64, 7, "expert_gate_up_n7"},
+        {4096, 64, 16, "expert_gate_up_n16"},
+        {4096, 64, 31, "expert_gate_up_n31"},
+        {4096, 64, 33, "expert_gate_up_n33"},
+        {2048, 64, 7, "expert_down_n7"},
+        {2048, 64, 16, "expert_down_n16"},
+        {2048, 64, 31, "expert_down_n31"},
+        {2048, 64, 33, "expert_down_n33"},
+        {4096, 4096, 31, "actual_gate_up_n31"},
+        {2048, 4096, 31, "actual_down_n31"},
     };
+    const char * shape_filter = std::getenv("DFLASH_TEST_SHAPE");
+    bool matched_shape = false;
 
     bool ok = true;
     for (const QuantCase & quant : quant_cases) {
         for (const Shape & shape : shapes) {
-            ok = test_case(hip_backend, quant, shape) && ok;
+            if (shape_filter && std::strcmp(shape.label, shape_filter) != 0) {
+                continue;
+            }
+            matched_shape = true;
+            if (std::strncmp(shape.label, "actual_", 7) == 0 &&
+                quant.type != GGML_TYPE_Q2_0_ROCMFP2) {
+                continue;
+            }
+            const DispatchPath expected_path =
+                shape.n <= 8 ? DispatchPath::MMVQ : DispatchPath::MMQ;
+            ok = test_case(hip_backend, quant, shape, expected_path) && ok;
         }
     }
+    if (shape_filter && !matched_shape) {
+        std::fprintf(stderr, "DFLASH_TEST_SHAPE matched no shape: %s\n",
+                     shape_filter);
+        ok = false;
+    }
 
     ggml_backend_free(hip_backend);
     std::printf("%s\n", ok ? "ALL TESTS PASSED" : "SOME TESTS FAILED");
diff --git a/server/test/test_server_unit.cpp b/server/test/test_server_unit.cpp
index 8373859dc..07e30b98b 100644
--- a/server/test/test_server_unit.cpp
+++ b/server/test/test_server_unit.cpp
@@ -2544,6 +2544,18 @@ TEST_CASE(ServerUnitFixture, test_max_output_alias_precedence_ignores_shadowed_i
     TEST_ASSERT(
         resolve_max_output_tokens({{"max_output_tokens", 200}}, 400) == 200);
     TEST_ASSERT(resolve_max_output_tokens(json::object(), 400) == 400);
+    // "Unlimited" sentinels from clients such as PocketPal must fall back
+    // to the default rather than yielding a zero-token budget.
+    TEST_ASSERT(
+        resolve_max_output_tokens({{"max_completion_tokens", -1}}, 400) == 400);
+    TEST_ASSERT(
+        resolve_max_output_tokens({{"max_completion_tokens", 0}}, 400) == 400);
+    TEST_ASSERT(
+        resolve_max_output_tokens({{"max_tokens", -1}}, 400) == 400);
+    TEST_ASSERT(
+        resolve_max_output_tokens({{"max_output_tokens", -1}}, 400) == 400);
+    TEST_ASSERT(
+        resolve_max_output_tokens({{"max_completion_tokens", 8}}, 400) == 8);
 }
 
 TEST_CASE(ServerUnitFixture, test_pflash_placement_same_backend_local) {
@@ -5112,7 +5124,7 @@ TEST_CASE(ServerUnitFixture, test_model_backend_retries_empty_visible_spec_resto
     TEST_ASSERT(backend.restore_saw_force_ar);
 }
 
-// GenerateResult.accept_rate plumbing tests (Day 1 of bandit MVP)
+// GenerateResult speculative telemetry plumbing tests (Day 1 of bandit MVP)
 // ═══════════════════════════════════════════════════════════════════════
 
 TEST_CASE(ServerUnitFixture, test_generate_result_accept_rate_defaults_to_zero) {
@@ -5141,6 +5153,7 @@ TEST_CASE(ServerUnitFixture, test_generate_result_accept_rate_in_usage_openai) {
     result.succeed();
     result.tokens = {1, 2, 3};
     result.accept_rate = 0.75f;
+    result.spec_decode_ran = true;
 
     std::vector<int32_t> prompt_tokens = {10, 20};
 
@@ -5150,12 +5163,14 @@ TEST_CASE(ServerUnitFixture, test_generate_result_accept_rate_in_usage_openai) {
             {"prompt_tokens", (int)prompt_tokens.size()},
             {"completion_tokens", (int)result.tokens.size()},
             {"total_tokens", (int)(prompt_tokens.size() + result.tokens.size())},
-            {"accept_rate", result.accept_rate}
+            {"accept_rate", result.accept_rate},
+            {"spec_decode_ran", result.spec_decode_ran}
         }}
     };
 
     TEST_ASSERT(resp["usage"].contains("accept_rate"));
     TEST_ASSERT(std::abs(resp["usage"]["accept_rate"].get<float>() - 0.75f) < 1e-6f);
+    TEST_ASSERT(resp["usage"]["spec_decode_ran"].get<bool>());
 }
 
 TEST_CASE(ServerUnitFixture, test_generate_result_accept_rate_in_usage_anthropic) {
@@ -5163,6 +5178,7 @@ TEST_CASE(ServerUnitFixture, test_generate_result_accept_rate_in_usage_anthropic
     result.succeed();
     result.tokens = {1, 2};
     result.accept_rate = 0.60f;
+    result.spec_decode_ran = true;
 
     std::vector<int32_t> prompt_tokens = {5};
 
@@ -5170,20 +5186,23 @@ TEST_CASE(ServerUnitFixture, test_generate_result_accept_rate_in_usage_anthropic
         {"usage", {
             {"input_tokens", (int)prompt_tokens.size()},
             {"output_tokens", (int)result.tokens.size()},
-            {"accept_rate", result.accept_rate}
+            {"accept_rate", result.accept_rate},
+            {"spec_decode_ran", result.spec_decode_ran}
         }}
     };
 
     TEST_ASSERT(resp["usage"].contains("accept_rate"));
     TEST_ASSERT(std::abs(resp["usage"]["accept_rate"].get<float>() - 0.60f) < 1e-6f);
+    TEST_ASSERT(resp["usage"]["spec_decode_ran"].get<bool>());
 }
 
 TEST_CASE(ServerUnitFixture, test_generate_result_accept_rate_zero_when_no_spec_decode) {
     // When spec decode doesn't run (no draft model), accept_rate stays 0.
     GenerateResult r;
     r.succeed();
-    // accept_rate not set → must be 0.0f
+    // Telemetry not set → accept_rate is zero and speculative decode is false.
     TEST_ASSERT(r.accept_rate == 0.0f);
+    TEST_ASSERT(!r.spec_decode_ran);
 }
 
 TEST_CASE(ServerUnitFixture, test_generate_result_error_state_is_consistent) {
diff --git a/server/tests/test_deepseek4_unit.cpp b/server/tests/test_deepseek4_unit.cpp
index 4a2b5f448..9a5f68322 100644
--- a/server/tests/test_deepseek4_unit.cpp
+++ b/server/tests/test_deepseek4_unit.cpp
@@ -1692,9 +1692,19 @@ static void test_hybrid_prefill_chunk_tokens() {
     TEST_ASSERT(deepseek4_hybrid_prefill_chunk_tokens(0, 8192) == 1);
     TEST_ASSERT(deepseek4_hybrid_prefill_chunk_tokens(
                     2048, 2048, 1024) == 1024);
+    TEST_ASSERT(deepseek4_hybrid_prefill_step_tokens(2048, 0, 70000) == 2048);
+    TEST_ASSERT(deepseek4_hybrid_prefill_step_tokens(
+                    2048, 30720, 70000) == 2048);
+    TEST_ASSERT(deepseek4_hybrid_prefill_step_tokens(
+                    2048, 31744, 70000) == 1024);
+    TEST_ASSERT(deepseek4_hybrid_prefill_step_tokens(
+                    2048, 32768, 70000) == 1024);
+    TEST_ASSERT(deepseek4_hybrid_prefill_step_tokens(
+                    512, 65536, 57) == 57);
+    TEST_ASSERT(deepseek4_hybrid_prefill_step_tokens(
+                    2048, 32768, 0) == 0);
     std::fprintf(stderr, g_failures ? " done\n" : " ok\n");
 }
-
 static void test_dspark_park_all_releases_drafter() {
     std::fprintf(stderr, "  test_dspark_park_all_releases_drafter ...");