perf(dflash): enable grouped MMID on RDNA

.github/workflows/ci.yml

@@ -172,13 +172,18 @@ jobs:
             -DDFLASH27B_ENABLE_BSA=OFF \
             -DDFLASH27B_FA_ALL_QUANTS=OFF \
             -DCMAKE_BUILD_TYPE=Release
-          cmake --build build --target test_flash_attn_sparse -j"$(nproc)"
+          cmake --build build \
+            --target test_flash_attn_sparse test_deepseek4_mmid_grouped_cuda \
+            -j"$(nproc)"
 
       - name: Run flash-attn sparse kernel test on the 3090
         # Self-contained kernel correctness test (no model weights), so it is
         # safe and fast in CI. This is the execution the hosted jobs cannot do.
         run: ./server/build/test_flash_attn_sparse
 
+      - name: Run grouped MMID dispatch and parity test on the 3090
+        run: ./server/build/test_deepseek4_mmid_grouped_cuda
+
       # Optional model-backed end-to-end smoke (real spec-decode on the 3090),
       # disabled by default because it builds dflash_server and lazy-loads the
       # ~16 GB Qwen3.6-27B target + draft (~1-2 min). The weights are already
@@ -285,7 +290,8 @@ jobs:
             -DCMAKE_HIP_FLAGS=-DDFLASH_WAVE_SIZE=32
           cmake --build "$RUNNER_TEMP/rocmfp-build" \
             --target test_rocmfp4 test_rocmfpx test_rocmfp4_hip_tail test_rocmfpx_mmq \
+              test_deepseek4_mmid_grouped_cuda \
             --parallel 8
           ctest --test-dir "$RUNNER_TEMP/rocmfp-build" \
             --output-on-failure \
-            -R 'rocmfp4_reference|rocmfpx_reference|rocmfp4_hip_tail|rocmfpx_mmq'
+            -R 'rocmfp4_reference|rocmfpx_reference|rocmfp4_hip_tail|rocmfpx_mmq|deepseek4_mmid_grouped_cuda'

server/CMakeLists.txt

@@ -749,6 +749,22 @@ if(DFLASH27B_TESTS)
         target_link_libraries(test_deepseek4_hc_cuda PRIVATE dflash_common ${DFLASH27B_GGML_BACKEND_TARGET})
         add_test(NAME deepseek4_hc_cuda COMMAND test_deepseek4_hc_cuda)
     endif()
+    if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/test/test_deepseek4_mmid_grouped_cuda.cpp")
+        add_executable(test_deepseek4_mmid_grouped_cuda test/test_deepseek4_mmid_grouped_cuda.cpp)
+        if(DFLASH27B_GPU_BACKEND STREQUAL "hip")
+            set_source_files_properties(test/test_deepseek4_mmid_grouped_cuda.cpp PROPERTIES LANGUAGE HIP)
+            set_target_properties(test_deepseek4_mmid_grouped_cuda PROPERTIES HIP_ARCHITECTURES "${_dflash_archs}")
+            target_include_directories(test_deepseek4_mmid_grouped_cuda PRIVATE
+                ${CMAKE_CURRENT_SOURCE_DIR}/hip_compat)
+            target_link_libraries(test_deepseek4_mmid_grouped_cuda PRIVATE hip::host)
+        else()
+            target_link_libraries(test_deepseek4_mmid_grouped_cuda PRIVATE CUDA::cudart)
+        endif()
+        target_link_libraries(test_deepseek4_mmid_grouped_cuda PRIVATE
+            ggml ${DFLASH27B_GGML_BACKEND_TARGET})
+        add_test(NAME deepseek4_mmid_grouped_cuda COMMAND test_deepseek4_mmid_grouped_cuda)
+        set_tests_properties(deepseek4_mmid_grouped_cuda PROPERTIES SKIP_RETURN_CODE 77)
+    endif()
     # GPU draft top-K kernel vs CPU reference (extract_draft_topk). Built on both
     # backends: geometric_draft_topk_cuda.cu is compiled into dflash_common on the
     # cuda backend directly and on hip via LANGUAGE HIP + the hip_compat shim.

server/deps/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu

@@ -2462,6 +2462,10 @@ static bool ggml_cuda_should_fuse_mul_mat_vec_q(const ggml_tensor * tensor) {
 
     // fusion is not universally faster on Pascal
     const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
+    if (is_mul_mat_id && ggml_cuda_mmvq_mmid_grouped_enabled(
+            src0->type, cc, dst->ne[2], dst->ne[1]*dst->ne[2])) {
+        return false;
+    }
     if (cc <= GGML_CUDA_CC_PASCAL) {
         return false;
     }
@@ -2498,6 +2502,12 @@ static bool ggml_cuda_try_fuse_mul_mat_glu(
     const ggml_tensor * src1 = up->src[1];
     const ggml_tensor * ids  = up->src[2];
 
+    const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
+    if (ids && ggml_cuda_mmvq_mmid_grouped_enabled(
+            src0->type, cc, up->ne[2], up->ne[1]*up->ne[2])) {
+        return false;
+    }
+
     if (ggml_cuda_should_fuse_mul_mat_vec_f(up)) {
         ggml_cuda_mm_fusion_args_host fusion_data{};
         fusion_data.gate = gate->src[0];
@@ -2531,7 +2541,6 @@ static bool ggml_cuda_try_fuse_mul_mat_glu(
             }
         }
 
-        const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc;
         const int64_t ncols = ids ? src1->ne[2] : src1->ne[1];
         if (ggml_cuda_should_use_mmq(src0->type, cc, ncols, src0->ne[2])) {
             ggml_cuda_mul_mat_q_pair(

server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu

@@ -256,7 +256,8 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type
 // weight traffic toward the union of routed experts. Bit-exact per
 // (row, token) vs mul_mat_vec_q_moe (same vec_dot sequence and reduction).
 // Model-agnostic: applies to any MoE with n_expert_used*n_tokens <= 256.
-//   DFLASH_MMID_GROUPED=1       opt-in
+//   DFLASH_MMID_GROUPED         1 = enable, 0 = disable
+//                               (CUDA default on; HIP default off)
 //   DFLASH_MMID_GROUPED_TYPES   bitmask, 1 = Q4_K (default), 2 = Q6_K,
 //                               4 = Q4_0/Q8_0/Q5_K. Q6_K stays on its tuned
 //                               MMQ route above 5 tokens unless enabled.
@@ -285,7 +286,9 @@ struct mmid_gate_extra {
 static bool mmid_grouped_env() {
     // Bit-exact and measured equal-or-faster on small MoE verify batches, so
     // enabled by default on CUDA; DFLASH_MMID_GROUPED=0 is the kill switch.
-    // HIP is unvalidated and stays opt-in (DFLASH_MMID_GROUPED=1).
+    // HIP (RDNA3/RDNA4) correctness is validated on gfx1151, but performance is
+    // still unmeasured, so it stays opt-in and default-off
+    // (DFLASH_MMID_GROUPED=1 to enable).
     static const bool on = []() {
         const char * e = std::getenv("DFLASH_MMID_GROUPED");
         if (e != nullptr) {
@@ -324,11 +327,26 @@ static bool mmid_grouped_type_ok(ggml_type type) {
     }
 }
 
+static bool mmid_grouped_arch_ok(int cc) {
+    return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_TURING) ||
+        GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc);
+}
+
+bool ggml_cuda_mmvq_mmid_grouped_enabled(
+        ggml_type type, int cc, int64_t ncols_dst, int64_t routed_pairs) {
+    return ncols_dst >= 2 && ncols_dst <= MMVQ_MAX_MOE_BATCH_SIZE &&
+        routed_pairs <= MMID_GROUPED_MAX_PAIRS &&
+        mmid_grouped_env() && mmid_grouped_type_ok(type) && mmid_grouped_arch_ok(cc);
+}
+
 // Host function: returns the max batch size for the current arch+type at runtime.
 int get_mmvq_mmid_max_batch(ggml_type type, int cc) {
     // [TAG_MMID_GROUPED] the grouped kernel handles any supported type up to the
     // MoE batch ceiling; this also keeps CUDA graphs on for these batches.
-    if (mmid_grouped_env() && mmid_grouped_type_ok(type) && GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_TURING) {
+    // RDNA3/RDNA4 (wave32) share the non-grouped kernel's wave-width warp_reduce.
+    // IS_RDNA3/IS_RDNA4 are pure cc-range checks, safe above the IS_AMD guard below.
+    // The HIP path remains opt-in until on-hardware parity and performance validation.
+    if (mmid_grouped_env() && mmid_grouped_type_ok(type) && mmid_grouped_arch_ok(cc)) {
         return MMVQ_MAX_MOE_BATCH_SIZE;
     }
     // Dedicated multi-token MoE kernel: extend the MUL_MAT_ID ceiling to 16
@@ -1122,6 +1140,8 @@ static __global__ void mul_mat_vec_q_moe_grouped(
     const float * x_bias = nullptr;
     const float * gate_bias = nullptr;
     ggml_glu_op active_glu;
+    float glu_param0 = 0.0f;
+    float glu_param1 = 0.0f;
 
     if constexpr (has_fusion) {
         use_gate      = fusion.gate      != nullptr;
@@ -1131,6 +1151,8 @@ static __global__ void mul_mat_vec_q_moe_grouped(
         x_bias        = (const float *) fusion.x_bias;
         gate_bias     = (const float *) fusion.gate_bias;
         active_glu    = fusion.glu_op;
+        glu_param0    = fusion.glu_param0;
+        glu_param1    = fusion.glu_param1;
     }
 
     float tmp[c_rows_per_block] = {0.0f};
@@ -1180,7 +1202,11 @@ static __global__ void mul_mat_vec_q_moe_grouped(
                         result *= ggml_cuda_op_gelu_single(gate_value);
                         break;
                     case GGML_GLU_OP_SWIGLU_OAI:
-                        result = ggml_cuda_op_swiglu_oai_single(gate_value, result);
+                        result = ggml_cuda_op_swiglu_oai_single(
+                            gate_value, result, glu_param0, glu_param1);
+                        break;
+                    case GGML_GLU_OP_SWIGLU_DS4:
+                        result = ggml_cuda_op_swiglu_ds4_single(gate_value, result, glu_param0);
                         break;
                     default:
                         result = result * gate_value;
@@ -1192,7 +1218,9 @@ static __global__ void mul_mat_vec_q_moe_grouped(
     }
 
     if constexpr (!has_fusion) {
-        GGML_UNUSED_VARS(use_gate, use_bias, use_gate_bias, vgate, x_bias, gate_bias, active_glu, tmp_gate);
+        GGML_UNUSED_VARS(
+            use_gate, use_bias, use_gate_bias, vgate, x_bias, gate_bias,
+            active_glu, glu_param0, glu_param1, tmp_gate);
     }
 }
 
@@ -2000,15 +2028,15 @@ void ggml_cuda_mul_mat_vec_q(
     const int64_t stride_channel_y   = ids ? s11  : s12;
 
     const int64_t ids_stride = ids ? ids->nb[1] / ggml_type_size(ids->type) : 0;
+    const int cc = ggml_cuda_info().devices[ctx.device].cc;
     static const bool mmid_telemetry = []() {
         const char * value = std::getenv("DFLASH_MMID_TELEMETRY");
         return value != nullptr && std::strcmp(value, "0") != 0;
     }();
 
     // [TAG_MMID_GROUPED] grouped-expert path for small MUL_MAT_ID batches.
-    if (ids && ncols_dst >= 2 && ncols_dst <= MMVQ_MAX_MOE_BATCH_SIZE &&
-        (int) (nchannels_dst*ncols_dst) <= MMID_GROUPED_MAX_PAIRS &&
-        mmid_grouped_env() && mmid_grouped_type_ok(src0->type)) {
+    if (ids && ggml_cuda_mmvq_mmid_grouped_enabled(
+            src0->type, cc, ncols_dst, nchannels_dst*ncols_dst)) {
         // Batches above MMID_GROUPED_MAX_PAIRS fall through to the legacy
         // per-expert kernel instead of aborting the request.
         const int np = (int) (nchannels_dst*ncols_dst);
@@ -2080,7 +2108,8 @@ void ggml_cuda_mul_mat_vec_q(
                 ncols_dst > MMVQ_MAX_MOE_BATCH_SIZE ? "width_gt_16" :
                 (int) (nchannels_dst*ncols_dst) > MMID_GROUPED_MAX_PAIRS ? "pairs_gt_256" :
                 !mmid_grouped_env() ? "flag_off" :
-                !mmid_grouped_type_ok(src0->type) ? "unsupported_type" : "dispatch_rejected";
+                !mmid_grouped_type_ok(src0->type) ? "unsupported_type" :
+                !mmid_grouped_arch_ok(cc) ? "unsupported_arch" : "dispatch_rejected";
             std::fprintf(stderr,
                 "[dflash-mmid] event=mmvq type=%s width=%lld pairs=%lld variant=%s reason=%s\n",
                 ggml_type_name(src0->type), (long long) ncols_dst,

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

@@ -11,6 +11,9 @@
 // based on the quantization type and GPU architecture (compute capability).
 int get_mmvq_mmid_max_batch(ggml_type type, int cc);
 
+bool ggml_cuda_mmvq_mmid_grouped_enabled(
+    ggml_type type, int cc, int64_t ncols_dst, int64_t routed_pairs);
+
 void ggml_cuda_mul_mat_vec_q(ggml_backend_cuda_context & ctx,
     const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst, const ggml_cuda_mm_fusion_args_host * fusion = nullptr);