feat(ds4): add fused HIP decode for ROCmFPX

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

@@ -497,7 +497,12 @@ static constexpr __host__ __device__ int calc_rows_per_block(int ncols_dst, int
     return 1;
 }
 
-template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false>
+static bool is_gfx1151(const int cc) {
+    return cc == GGML_CUDA_CC_OFFSET_AMD + 0x1151;
+}
+
+template <ggml_type type, int ncols_dst, bool has_fusion, bool small_k = false,
+          int fixed_ncols_x = 0, bool unroll_k_loop_2 = false>
 __launch_bounds__(calc_nwarps(type, ncols_dst, get_device_table_id())*ggml_cuda_get_physical_warp_size(), 1)
 static __global__ void mul_mat_vec_q(
         const void * __restrict__ vx, const void * __restrict__ vy, const int32_t * __restrict__ ids, const ggml_cuda_mm_fusion_args_device fusion, float * __restrict__ dst,
@@ -519,9 +524,16 @@ static __global__ void mul_mat_vec_q(
 
     const     int tid = warp_size*threadIdx.y + threadIdx.x;
     const     int row0 = rows_per_cuda_block*blockIdx.x;
-    const     int blocks_per_row_x = ncols_x / qk;
+    const     int blocks_per_row_x = (fixed_ncols_x > 0 ? fixed_ncols_x : ncols_x) / qk;
     constexpr int blocks_per_iter = vdr * nwarps*warp_size / qi;
 
+    static_assert(fixed_ncols_x == 0 ||
+                  type == GGML_TYPE_Q3_0_ROCMFPX ||
+                  type == GGML_TYPE_Q2_0_ROCMFP2,
+                  "fixed-K decode specialization is limited to profiled ROCmFP2/FP3 types");
+    static_assert(!unroll_k_loop_2 || type == GGML_TYPE_Q4_0_ROCMFP4_FAST,
+                  "partial K-loop unrolling is limited to the profiled FP4FAST path");
+
     const uint32_t channel_dst = blockIdx.y;
 
     const bool has_ids = ids != nullptr;
@@ -601,23 +613,76 @@ static __global__ void mul_mat_vec_q(
 
     const block_q8_1 * y = ((const block_q8_1 *) vy) + sample_y*stride_sample_y + channel_y*stride_channel_y;
 
-    for (int kbx = tid / (qi/vdr); kbx < blocks_per_row_x; kbx += blocks_per_iter) {
-        const int kby = kbx * (qk/QK8_1); // y block index that aligns with kbx
+    if constexpr (fixed_ncols_x > 0) {
+        static_assert(fixed_ncols_x % qk == 0, "fixed K must contain whole quant blocks");
+        constexpr int fixed_blocks = fixed_ncols_x/qk;
+        constexpr int fixed_iters = (fixed_blocks + blocks_per_iter - 1) / blocks_per_iter;
+#pragma unroll
+        for (int iter = 0; iter < fixed_iters; ++iter) {
+            const int kbx = tid / (qi/vdr) + iter*blocks_per_iter;
+            if constexpr (fixed_blocks % blocks_per_iter != 0) {
+                if (kbx >= fixed_blocks) {
+                    continue;
+                }
+            }
+            const int kby = kbx * (qk/QK8_1);
+            const int kqs = vdr * (tid % (qi/vdr));
 
-        // x block quant index when casting the quants to int
-        const int kqs = vdr * (tid % (qi/vdr));
+#pragma unroll
+            for (int j = 0; j < ncols_dst; ++j) {
+                const int kbx_offset = sample_x*stride_sample_x + channel_xs[j]*stride_channel_x + row0*stride_row_x;
+#pragma unroll
+                for (int i = 0; i < rows_per_cuda_block; ++i) {
+                    tmp[j][i] += vec_dot_q_cuda(
+                        vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
+                    if constexpr (has_fusion) {
+                        if (use_gate) {
+                            tmp_gate[j][i] += vec_dot_q_cuda(
+                                vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
+                        }
+                    }
+                }
+            }
+        }
+    } else if constexpr (unroll_k_loop_2) {
+#pragma unroll 2
+        for (int kbx = tid / (qi/vdr); kbx < blocks_per_row_x; kbx += blocks_per_iter) {
+            const int kby = kbx * (qk/QK8_1);
+            const int kqs = vdr * (tid % (qi/vdr));
 
 #pragma unroll
-        for (int j = 0; j < ncols_dst; ++j) {
-            const int kbx_offset = sample_x*stride_sample_x + channel_xs[j]*stride_channel_x + row0*stride_row_x;
+            for (int j = 0; j < ncols_dst; ++j) {
+                const int kbx_offset = sample_x*stride_sample_x + channel_xs[j]*stride_channel_x + row0*stride_row_x;
 #pragma unroll
-            for (int i = 0; i < rows_per_cuda_block; ++i) {
-                tmp[j][i] += vec_dot_q_cuda(
-                    vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
-                if constexpr (has_fusion) {
-                    if (use_gate) {
-                        tmp_gate[j][i] += vec_dot_q_cuda(
-                            vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
+                for (int i = 0; i < rows_per_cuda_block; ++i) {
+                    tmp[j][i] += vec_dot_q_cuda(
+                        vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
+                    if constexpr (has_fusion) {
+                        if (use_gate) {
+                            tmp_gate[j][i] += vec_dot_q_cuda(
+                                vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
+                        }
+                    }
+                }
+            }
+        }
+    } else {
+        for (int kbx = tid / (qi/vdr); kbx < blocks_per_row_x; kbx += blocks_per_iter) {
+            const int kby = kbx * (qk/QK8_1);
+            const int kqs = vdr * (tid % (qi/vdr));
+
+#pragma unroll
+            for (int j = 0; j < ncols_dst; ++j) {
+                const int kbx_offset = sample_x*stride_sample_x + channel_xs[j]*stride_channel_x + row0*stride_row_x;
+#pragma unroll
+                for (int i = 0; i < rows_per_cuda_block; ++i) {
+                    tmp[j][i] += vec_dot_q_cuda(
+                        vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
+                    if constexpr (has_fusion) {
+                        if (use_gate) {
+                            tmp_gate[j][i] += vec_dot_q_cuda(
+                                vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs);
+                        }
                     }
                 }
             }
@@ -1138,7 +1203,8 @@ static std::pair<dim3, dim3> calc_launch_params(
     return {block_nums, block_dims};
 }
 
-template<ggml_type type, int c_ncols_dst, bool small_k = false>
+template<ggml_type type, int c_ncols_dst, bool small_k = false,
+         int fixed_ncols_x = 0, bool unroll_k_loop_2 = false>
 static void mul_mat_vec_q_switch_fusion(
         const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst,
         const uint32_t ncols_x, const uint3 nchannels_y, const uint32_t stride_row_x, const uint32_t stride_col_y,
@@ -1150,19 +1216,70 @@ static void mul_mat_vec_q_switch_fusion(
 
     const bool has_fusion = fusion.gate != nullptr || fusion.x_bias != nullptr || fusion.gate_bias != nullptr;
     if (has_fusion) {
-        mul_mat_vec_q<type, c_ncols_dst, true, small_k><<<block_nums, block_dims, nbytes_shared, stream>>>
+        mul_mat_vec_q<type, c_ncols_dst, true, small_k, fixed_ncols_x, unroll_k_loop_2><<<block_nums, block_dims, nbytes_shared, stream>>>
             (vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
              channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
              sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, ids_tokenwise_samples);
         return;
     }
 
-    mul_mat_vec_q<type, c_ncols_dst, false, small_k><<<block_nums, block_dims, nbytes_shared, stream>>>
+    mul_mat_vec_q<type, c_ncols_dst, false, small_k, fixed_ncols_x, unroll_k_loop_2><<<block_nums, block_dims, nbytes_shared, stream>>>
         (vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x, stride_col_y, stride_col_dst,
         channel_ratio, stride_channel_x, stride_channel_y, stride_channel_dst,
         sample_ratio, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, ids_tokenwise_samples);
 }
 
+template <ggml_type type, int fixed_ncols_x>
+static void mul_mat_vec_rocmfpx_fixed_k_launch(
+        const void * vx, const void * vy, const int32_t * ids,
+        const ggml_cuda_mm_fusion_args_device fusion, float * dst,
+        const uint32_t ncols_x, const uint32_t nrows_x, const uint3 nchannels_y,
+        const uint32_t nchannels_dst, const uint32_t stride_row_x,
+        const uint32_t stride_col_y, const uint32_t stride_col_dst,
+        const uint3 channel_ratio, const uint32_t stride_channel_x,
+        const uint32_t stride_channel_y, const uint32_t stride_channel_dst,
+        const uint32_t nsamples_dst, const uint3 sample_ratio,
+        const uint32_t stride_sample_x, const uint32_t stride_sample_y,
+        const uint32_t stride_sample_dst, const uint32_t ids_stride,
+        const int warp_size, cudaStream_t stream) {
+    static_assert(type == GGML_TYPE_Q3_0_ROCMFPX ||
+                  type == GGML_TYPE_Q2_0_ROCMFP2,
+                  "fixed-K helper is limited to profiled ROCmFP2/FP3 kernels");
+    static_assert(fixed_ncols_x == 2048 || fixed_ncols_x == 4096,
+                  "fixed-K helper compiled an unexpected DS4 shape");
+
+    const dim3 block_nums(nrows_x, nchannels_dst, nsamples_dst);
+    const dim3 block_dims(warp_size, 1, 1);
+    mul_mat_vec_q_switch_fusion<type, 1, false, fixed_ncols_x>(
+        vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x,
+        stride_col_y, stride_col_dst, channel_ratio, stride_channel_x,
+        stride_channel_y, stride_channel_dst, sample_ratio, stride_sample_x,
+        stride_sample_y, stride_sample_dst, block_nums, block_dims, 0,
+        ids_stride, stream);
+}
+
+static void mul_mat_vec_rocmfp4_unroll2_launch(
+        const void * vx, const void * vy, const int32_t * ids,
+        const ggml_cuda_mm_fusion_args_device fusion, float * dst,
+        const uint32_t ncols_x, const uint32_t nrows_x, const uint3 nchannels_y,
+        const uint32_t nchannels_dst, const uint32_t stride_row_x,
+        const uint32_t stride_col_y, const uint32_t stride_col_dst,
+        const uint3 channel_ratio, const uint32_t stride_channel_x,
+        const uint32_t stride_channel_y, const uint32_t stride_channel_dst,
+        const uint32_t nsamples_dst, const uint3 sample_ratio,
+        const uint32_t stride_sample_x, const uint32_t stride_sample_y,
+        const uint32_t stride_sample_dst, const uint32_t ids_stride,
+        const int warp_size, cudaStream_t stream) {
+    const dim3 block_nums(nrows_x, nchannels_dst, nsamples_dst);
+    const dim3 block_dims(warp_size, 1, 1);
+    mul_mat_vec_q_switch_fusion<GGML_TYPE_Q4_0_ROCMFP4_FAST, 1, false, 0, true>(
+        vx, vy, ids, fusion, dst, ncols_x, nchannels_y, stride_row_x,
+        stride_col_y, stride_col_dst, channel_ratio, stride_channel_x,
+        stride_channel_y, stride_channel_dst, sample_ratio, stride_sample_x,
+        stride_sample_y, stride_sample_dst, block_nums, block_dims, 0,
+        ids_stride, stream);
+}
+
 template <ggml_type type>
 static void mul_mat_vec_q_moe_launch(
         const void * vx, const void * vy, const int32_t * ids, const ggml_cuda_mm_fusion_args_device fusion, float * dst,
@@ -1333,6 +1450,60 @@ static void mul_mat_vec_q_switch_ncols_dst(
         return;
     }
 
+    // Decode-only gfx1151 specializations. Each one preserves the original
+    // per-lane K traversal and accumulation order, so exact validated shapes
+    // can use the faster kernel without a serving-time tuning flag.
+    if constexpr (type == GGML_TYPE_Q4_0_ROCMFP4_FAST) {
+        if (is_gfx1151(cc) && ncols_dst == 1) {
+            mul_mat_vec_rocmfp4_unroll2_launch(
+                vx, vy, ids, fusion, dst, ncols_x, nrows_x, nchannels_y_fd,
+                nchannels_dst, stride_row_x, stride_col_y, stride_col_dst,
+                channel_ratio_fd, stride_channel_x, stride_channel_y,
+                stride_channel_dst, nsamples_dst, sample_ratio_fd,
+                stride_sample_x, stride_sample_y, stride_sample_dst,
+                ids_stride, warp_size, stream);
+            return;
+        }
+    }
+
+    if constexpr (type == GGML_TYPE_Q3_0_ROCMFPX) {
+        if (is_gfx1151(cc) && ncols_dst == 1 && ncols_x == 2048) {
+            mul_mat_vec_rocmfpx_fixed_k_launch<type, 2048>(
+                vx, vy, ids, fusion, dst, ncols_x, nrows_x, nchannels_y_fd,
+                nchannels_dst, stride_row_x, stride_col_y, stride_col_dst,
+                channel_ratio_fd, stride_channel_x, stride_channel_y,
+                stride_channel_dst, nsamples_dst, sample_ratio_fd,
+                stride_sample_x, stride_sample_y, stride_sample_dst,
+                ids_stride, warp_size, stream);
+            return;
+        }
+    }
+
+    if constexpr (type == GGML_TYPE_Q2_0_ROCMFP2) {
+        if (is_gfx1151(cc) && ncols_dst == 1) {
+            if (ncols_x == 4096) {
+                mul_mat_vec_rocmfpx_fixed_k_launch<type, 4096>(
+                    vx, vy, ids, fusion, dst, ncols_x, nrows_x, nchannels_y_fd,
+                    nchannels_dst, stride_row_x, stride_col_y, stride_col_dst,
+                    channel_ratio_fd, stride_channel_x, stride_channel_y,
+                    stride_channel_dst, nsamples_dst, sample_ratio_fd,
+                    stride_sample_x, stride_sample_y, stride_sample_dst,
+                    ids_stride, warp_size, stream);
+                return;
+            }
+            if (ncols_x == 2048) {
+                mul_mat_vec_rocmfpx_fixed_k_launch<type, 2048>(
+                    vx, vy, ids, fusion, dst, ncols_x, nrows_x, nchannels_y_fd,
+                    nchannels_dst, stride_row_x, stride_col_y, stride_col_dst,
+                    channel_ratio_fd, stride_channel_x, stride_channel_y,
+                    stride_channel_dst, nsamples_dst, sample_ratio_fd,
+                    stride_sample_x, stride_sample_y, stride_sample_dst,
+                    ids_stride, warp_size, stream);
+                return;
+            }
+        }
+    }
+
     switch (ncols_dst) {
         case 1: {
             constexpr int c_ncols_dst = 1;

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

@@ -401,6 +401,29 @@ static __device__ __forceinline__ int rocmfpx_pack4_fp6_bits24_vec_cuda(const ui
     return *((const int *) &v);
 }
 
+static __device__ __forceinline__ int rocmfpx_pack4_fp3_bits12_vec_cuda(const uint32_t bits12) {
+#if defined(GGML_USE_HIP)
+    // Byte tables for codes {0, 1, 2, 4} and {0, -1, -2, -4}.
+    // v_perm_b32 selects a byte with each 3-bit selector, replacing four
+    // scalar code decodes without changing the packed int8 values.
+    constexpr uint32_t values_low  = 0x04020100u;
+    constexpr uint32_t values_high = 0xFCFEFF00u;
+    const uint32_t selectors =
+        ((bits12 >> 0) & 7u) |
+        (((bits12 >> 3) & 7u) << 8) |
+        (((bits12 >> 6) & 7u) << 16) |
+        (((bits12 >> 9) & 7u) << 24);
+    return (int) __builtin_amdgcn_perm(values_high, values_low, selectors);
+#else
+    const char4 v = make_char4(
+        (int8_t) rocmfpx_decode_fp3_code_vec_cuda(bits12 & 7u),
+        (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 3) & 7u),
+        (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 6) & 7u),
+        (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 9) & 7u));
+    return *((const int *) &v);
+#endif
+}
+
 static __device__ __forceinline__ int rocmfpx_pack4_fp3_vec_cuda(const uint8_t * qs, const int base) {
     const char4 v = make_char4(
         (int8_t) rocmfpx_decode_fp3_code_vec_cuda(rocmfpx_get_bits_vec_cuda(qs, (base + 0)*3, 3)),
@@ -532,12 +555,7 @@ static __device__ __forceinline__ float vec_dot_rocmfpx_fp3_q8_1(
         const uint32_t val_high = qs[reg_idx + 1];
         const uint32_t bits12 = (reg_shift == 0) ? (val_low & 0xFFFu) : (((val_low >> reg_shift) | (val_high << (32 - reg_shift))) & 0xFFFu);
 
-        const char4 v = make_char4(
-            (int8_t) rocmfpx_decode_fp3_code_vec_cuda(bits12 & 7u),
-            (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 3) & 7u),
-            (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 6) & 7u),
-            (int8_t) rocmfpx_decode_fp3_code_vec_cuda((bits12 >> 9) & 7u));
-        const int val_packed = *((const int *) &v);
+        const int val_packed = rocmfpx_pack4_fp3_bits12_vec_cuda(bits12);
 
         const int u = get_int_b4(bq8_1->qs, iqs + i);
 

server/docs/DS4.md

@@ -1,6 +1,8 @@
 # DeepSeek V4 Flash — DFlash Integration
 
-This document describes the current DeepSeek V4 Flash implementation in DFlash. Today DeepSeek4 runs through the layer-split path: a local CUDA prefix shard plus either a local follow-on shard or a remote Halo/HIP target shard.
+This document describes the current DeepSeek V4 Flash implementation in
+DFlash. DeepSeek4 supports a monolithic HIP backend for single-device Strix
+Halo systems and a layer-split backend for local or mixed-device deployments.
 
 ## Model Architecture
 
@@ -23,7 +25,7 @@ DeepSeek V4 Flash is a 43-layer MoE model with:
 
 | Area | Files |
 |------|-------|
-| Backend selection / init | `src/common/backend_factory.cpp`, `src/deepseek4/deepseek4_layer_split_adapter.{h,cpp}` |
+| Backend selection / init | `src/common/backend_factory.cpp`, `src/deepseek4/deepseek4_backend.{h,cpp}`, `src/deepseek4/deepseek4_layer_split_adapter.{h,cpp}` |
 | Per-shard forward graph | `src/deepseek4/deepseek4_graph.cpp` |
 | Model weights and metadata | `src/deepseek4/deepseek4_internal.h`, `src/deepseek4/deepseek4_loader.cpp` |
 | HC pre/post CUDA kernel | `src/deepseek4/deepseek4_hc_cuda.cu`, `.h` |
@@ -45,6 +47,32 @@ The production DeepSeek4 path does **not** use the retired per-expert worker spl
 
 ## Execution Modes
 
+### Monolithic HIP
+
+Single-device HIP launches use `DeepSeek4Backend`. Two explicit serving
+options are available:
+
+- `--ds4-fused-decode` enables the cached single-graph decode path. It keeps
+  HC, attention, MoE, and the output projection on the GPU and avoids
+  per-layer host round trips. On HIP this option requests a monolithic model
+  load because the fused graph must reference every expert tensor directly.
+  If that allocation fails, the backend logs the fallback and continues with
+  hybrid expert placement and layered decode.
+- `--ds4-expert-top-k N` keeps the highest-ranked `N` routed experts and
+  renormalizes their weights. `0` uses the model default. Reducing this value is an
+  approximate inference policy and must be quality-validated for the target
+  workload.
+
+These options currently apply only to the monolithic HIP backend. For the
+validated Strix Halo profile:
+
+```bash
+./server/build-hip/dflash_server /opt/models/DeepSeek-V4-Flash.gguf \
+  --target-device hip:0 \
+  --ds4-fused-decode \
+  --ds4-expert-top-k 4
+```
+
 ### Local single-shard
 
 If the adapter decides all 43 layers fit on one CUDA GPU, it loads a single shard locally and no IPC daemon is involved.
@@ -148,6 +176,6 @@ Explicit mixed-backend split using the generic target-shard flags:
 
 | Target | Backend | Purpose |
 |--------|---------|---------|
-| `dflash_server` | CUDA | Production server / CUDA parent |
+| `dflash_server` | CUDA or HIP | Production server |
 | `backend_ipc_daemon` | HIP | Remote Halo target shard for mixed-backend layer split |
 | `test_deepseek4_unit` | CUDA | Unit tests (no model files needed) |

server/src/common/backend_factory.cpp

@@ -232,6 +232,8 @@ std::unique_ptr<ModelBackend> create_backend(const BackendArgs & args) {
             cfg.stream_fd  = args.stream_fd;
             cfg.max_ctx    = args.device.max_ctx;
             cfg.chunk      = args.chunk;
+            cfg.expert_top_k = args.ds4_expert_top_k;
+            cfg.fused_decode = args.ds4_fused_decode;
 
             auto backend = std::make_unique<DeepSeek4Backend>(cfg);
             if (!backend->init()) {

server/src/common/backend_factory.h

@@ -45,6 +45,10 @@ struct BackendArgs {
     // Chunked prefill
     int             chunk        = 512;
 
+    // deepseek4-specific decode options
+    int             ds4_expert_top_k = 0;  // 0 = model default
+    bool            ds4_fused_decode = false;
+
     // qwen35-specific speculative decode options
     int             fa_window        = 0;  // 0 = full attention. qwen3.6 full-attn layers must see the whole context; a finite window drops the system prompt/tools -> breaks tool calls.
     int             kq_stride_pad    = 32;