feat(ggml): vendor ROCmFPX quant support for DeepSeek V4 Flash

[davide221]

Summary

Ports the ROCmFP4/ROCmFPX ggml work — previously lucebox-ggml #36, which targeted the old submodule pointer — into the vendored server/deps/llama.cpp tree, now that ggml lives in-tree (#486).

• ROCmFP4/ROCmFPX quant types + CPU reference conversions (ggml/rocmfp4, ggml/rocmfpx) and type-trait registration.
• CUDA/HIP dequant, copy, getrows, MMVQ vecdot, MMVF, unary and FA paths for the new types.
• Fused DS4 hyper-connection op GGML_OP_DS4_HC (register-resident sinkhorn); inert unless the server emits it.
• DS4 SwiGLU-split op plumbing for the fused FFN matvec path.

Layered on current main (no reverts)

The branch is built on current main, so main-side ggml work is preserved: fp64 RoPE reduction, the #497 RDNA MMQ tile, LUCE_MMQ_DP_MAX_NE1, MMVQ_MAX_MOE_BATCH_SIZE, fused dual set_rows, and the raw-span guard all remain intact. (A naive tree copy of #36 would have reverted these; this is the base→tip diff only.)

Review fixes on top of #36

• ggml_ftype_to_ggml_type: handle the 11 new ROCmFPX ftypes (dominant-type mapping) so the enum switch is -Wswitch/-Werror clean.
• FP6 MMVQ vecdot: pad qs[] so the last window no longer reads one past the array. Bit-identical (the over-read bits were already masked out).

Known follow-ups (non-blocking, unchanged from #36)

• test_rocmfpx.c and rocmfp4_hip.cu are not wired into any build target yet; the imatrix/CPU-dot surface of both libs is currently unused.
• The env-gated custom FA path (DFLASH_DS4_FLASH_ATTN, default OFF) needs a shared-memory bound + post-launch error check before it is promoted off the flag.

Validation

Byte-for-byte the same ggml sources that built and ran on the Strix Halo box (DeepSeek-V4-Flash-ROCMFP2-STRIX.gguf), plus the two review fixes above.

[cubic-dev-ai]

17 issues found across 31 files

Prompt for AI agents (unresolved issues)

Check if these issues are valid — if so, understand the root cause of each and fix them. If appropriate, use sub-agents to investigate and fix each issue separately.


<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/cpy.cu">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/cpy.cu:160">
P3: The contiguous ROCmFP4 copy path is harder to keep correct because it bakes in 16 packed bytes while surrounding code derives the layout from QK_ROCMFP4. Deriving ib/j from QK_ROCMFP4/2, as the standalone ROCmFP4 HIP dequant kernel does, would avoid wrong block/qs indexing if the block layout changes.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu:286">
P1: Mode 1 HC updates mix the residual streams with the transposed Sinkhorn combine matrix. Since mode 0 writes comb as `src + dst * n_hc`, this read should use `src + h * n_hc` so each destination HC stream receives its own combine row.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/rocmfpx/README.md">

<violation number="1" location="server/deps/llama.cpp/ggml/rocmfpx/README.md:59">
P3: Validation Script Index omits three Qwen scripts referenced in the validation examples below. The index lists `check-rocmfpx-qwen-all.sh` (core Qwen gates) but the validation examples on lines 103-105 use `check-rocmfpx-qwen-coherency.sh`, `check-rocmfpx-qwen-bench.sh`, and `check-rocmfpx-qwen-strict-json.sh`. Either add those three to the index or clarify the relationship (e.g. qwen-all.sh is a wrapper).</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/common.cuh">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/common.cuh:49">
P2: Enabling `GGML_ROCMFP6_EXPANDED_DEVICE` makes Q6_0_ROCMFPX GPU reads use a different block stride than the stored tensor data. Consider keeping the device alias compact until the CUDA/HIP copy path expands each block and adjusts allocation/stride metadata consistently.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp:9878">
P1: CPU execution of ggml_swiglu_ds4_split currently aborts during task planning before reaching this new handler. The GLU scheduler also needs a GGML_GLU_OP_SWIGLU_DS4 case that assigns n_tasks = n_threads, matching the other threaded GLU ops.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c">

<violation number="1" location="server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c:14">
P2: `fill_row` writes to indices 7, 19, and 43 unconditionally but takes `n` as its size parameter. If it's ever called with `n < 44`, this silently corrupts memory. Guard the writes with `if (n > 43)` or document the minimum size requirement (`n >= 44`) in a comment so the contract is visible at the call site.</violation>

<violation number="2" location="server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c:157">
P3: MSE monotonicity is verified for FP8 < FP6 < FP3, but there's no assertion that FP3 < FP2 (i.e., 3-bit rounds more accurately than 2-bit). Add `assert(mse3 < mse2)` after the existing assertions so all four bit depths are checked.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip_scale.cuh">

<violation number="1" location="server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip_scale.cuh:29">
P2: Enabling `GGML_ROCMFP4_USE_SCALE_LUT` brings back the constant-memory scale decoder that this directory’s ROCmFP4 contract explicitly rejects after a runtime regression. Keeping only the arithmetic exponent/mantissa path would preserve the documented guardrail for this hot scale decode path.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c">

<violation number="1" location="server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c:477">
P2: FP6 quantization can produce undefined or implementation-defined results for large finite values because the rounded magnitude is converted to int before clamping. Clamping the float magnitude before lroundf keeps saturated outliers deterministic.</violation>

<violation number="2" location="server/deps/llama.cpp/ggml/rocmfpx/rocmfpx.c:610">
P2: FP8 quantization can mis-handle large finite outliers because saturation happens after the int conversion. Clamping the scaled float to [-127, 127] before lroundf avoids overflow and preserves the documented clamp contract.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip.cu">

<violation number="1" location="server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip.cu:63">
P3: The new standalone HIP dequant helpers are currently dead code because `rocmfp4_hip.cu` is not wired into any ggml build target. If these are intended for integration tests or future fused paths, consider adding the file to the ROCm/HIP target or moving it behind a tracked follow-up so it does not silently drift.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/fattn.cu">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/fattn.cu:51">
P1: The env-gated DS4 path will compute different attention scores whenever `GGML_OP_FLASH_ATTN_EXT` carries a non-default scale or logit softcap. This kernel should use `dst->op_params[0]` and preserve the existing logit-softcap handling instead of hardcoding `1/sqrt(512)`.</violation>

<violation number="2" location="server/deps/llama.cpp/ggml/src/ggml-cuda/fattn.cu:134">
P1: If this custom launch is rejected, for example because `n_kv * sizeof(float)` exceeds the device dynamic shared-memory limit, the dispatcher still reports success and skips all fallback kernels. Checking the launch error before returning would prevent silent incorrect completion on the env-gated path.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/convert.cu">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/convert.cu:688">
P1: Contiguous ROCmFP4 conversion can write past the destination for valid tensors whose element count is only QK_ROCMFP4-aligned. The row launcher rounds by QK_K and the kernel always emits QK_K values, so this path should either assert QK_K alignment or handle the tail/launch per ROCmFP4 block.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml.c">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml.c:8088">
P2: DS4_HC results can be wrong for view/strided tensors because the new helpers accept non-contiguous inputs while the CUDA implementation reads each source as a flat array. Consider asserting `ggml_is_contiguous(...)` for `mix/base/hc_state` and the corresponding `post/out` sources, or making the kernel stride-aware.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.c">

<violation number="1" location="server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.c:59">
P3: ROCmFP4 scale decode now violates the format rules documented for this directory: CPU decode is supposed to use the arithmetic UE4M3 decoder, not a LUT. Keeping the documented decoder avoids divergence from the HIP/Vulkan contract and the promoted regression assumptions.</violation>

<violation number="2" location="server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.c:458">
P2: Imatrix quantization can collapse finite blocks to zero scale when one row value is NaN, because sigma2 accumulates non-finite `x[i]*x[i]` before the guarded per-element scoring path. Consider skipping non-finite inputs in the row-energy accumulation or otherwise normalizing sigma2 before weighting.</violation>
</file>

<sub>Reply with feedback, questions, or to request a fix.

Re-trigger cubic</sub>

[cubic-dev-ai]

P1: Mode 1 HC updates mix the residual streams with the transposed Sinkhorn combine matrix. Since mode 0 writes comb as src + dst * n_hc, this read should use src + h * n_hc so each destination HC stream receives its own combine row.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-hc.cu, line 286:

<comment>Mode 1 HC updates mix the residual streams with the transposed Sinkhorn combine matrix. Since mode 0 writes comb as `src + dst * n_hc`, this read should use `src + h * n_hc` so each destination HC stream receives its own combine row.</comment>

<file context>
@@ -0,0 +1,367 @@
+    const float * comb = split + 2 * n_hc;
+    float acc = block_out[d] * post[h];
+    for (int src = 0; src < n_hc; ++src) {
+        acc += comb[h + src * n_hc] * residual[(size_t) src * n_embd + d];
+    }
+    dst[i] = acc;
</file context>

[cubic-dev-ai]

P1: CPU execution of ggml_swiglu_ds4_split currently aborts during task planning before reaching this new handler. The GLU scheduler also needs a GGML_GLU_OP_SWIGLU_DS4 case that assigns n_tasks = n_threads, matching the other threaded GLU ops.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp, line 9878:

<comment>CPU execution of ggml_swiglu_ds4_split currently aborts during task planning before reaching this new handler. The GLU scheduler also needs a GGML_GLU_OP_SWIGLU_DS4 case that assigns n_tasks = n_threads, matching the other threaded GLU ops.</comment>

<file context>
@@ -9801,6 +9875,10 @@ void ggml_compute_forward_glu(
             {
                 ggml_compute_forward_swiglu_oai(params, dst);
             } break;
+        case GGML_GLU_OP_SWIGLU_DS4:
+            {
+                ggml_compute_forward_swiglu_ds4(params, dst);
</file context>

[cubic-dev-ai]

P1: If this custom launch is rejected, for example because n_kv * sizeof(float) exceeds the device dynamic shared-memory limit, the dispatcher still reports success and skips all fallback kernels. Checking the launch error before returning would prevent silent incorrect completion on the env-gated path.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cuda/fattn.cu, line 134:

<comment>If this custom launch is rejected, for example because `n_kv * sizeof(float)` exceeds the device dynamic shared-memory limit, the dispatcher still reports success and skips all fallback kernels. Checking the launch error before returning would prevent silent incorrect completion on the env-gated path.</comment>

<file context>
@@ -7,6 +7,133 @@
+            n_tokens,
+            n_heads,
+            n_kv);
+    return true;
+}
+
</file context>

Suggested change

	return true;
	CUDA_CHECK(cudaGetLastError());
	return true;

[cubic-dev-ai]

P1: The env-gated DS4 path will compute different attention scores whenever GGML_OP_FLASH_ATTN_EXT carries a non-default scale or logit softcap. This kernel should use dst->op_params[0] and preserve the existing logit-softcap handling instead of hardcoding 1/sqrt(512).

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cuda/fattn.cu, line 51:

<comment>The env-gated DS4 path will compute different attention scores whenever `GGML_OP_FLASH_ATTN_EXT` carries a non-default scale or logit softcap. This kernel should use `dst->op_params[0]` and preserve the existing logit-softcap handling instead of hardcoding `1/sqrt(512)`.</comment>

<file context>
@@ -7,6 +7,133 @@
+
+    extern __shared__ float scores[];
+    const float * qh = q + ((size_t) h * (size_t) n_tokens + (size_t) t) * D;
+    const float scale = rsqrtf((float) D);
+
+    float local_max = sinks ? sinks[h] : -3.402823466e38f;
</file context>

[cubic-dev-ai]

P1: Contiguous ROCmFP4 conversion can write past the destination for valid tensors whose element count is only QK_ROCMFP4-aligned. The row launcher rounds by QK_K and the kernel always emits QK_K values, so this path should either assert QK_K alignment or handle the tail/launch per ROCmFP4 block.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cuda/convert.cu, line 688:

<comment>Contiguous ROCmFP4 conversion can write past the destination for valid tensors whose element count is only QK_ROCMFP4-aligned. The row launcher rounds by QK_K and the kernel always emits QK_K values, so this path should either assert QK_K alignment or handle the tail/launch per ROCmFP4 block.</comment>

<file context>
@@ -645,6 +683,18 @@ static void dequantize_row_mxfp4_cuda(const void * vx, dst_t * y, const int64_t
 
+template<typename dst_t>
+static void dequantize_row_rocmfp4_hip(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
+    const int nb = (k + QK_K - 1) / QK_K;
+    dequantize_block_rocmfp4<<<nb, 32, 0, stream>>>(vx, y);
+}
</file context>

[cubic-dev-ai]

P3: The contiguous ROCmFP4 copy path is harder to keep correct because it bakes in 16 packed bytes while surrounding code derives the layout from QK_ROCMFP4. Deriving ib/j from QK_ROCMFP4/2, as the standalone ROCmFP4 HIP dequant kernel does, would avoid wrong block/qs indexing if the block layout changes.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/src/ggml-cuda/cpy.cu, line 160:

<comment>The contiguous ROCmFP4 copy path is harder to keep correct because it bakes in 16 packed bytes while surrounding code derives the layout from QK_ROCMFP4. Deriving ib/j from QK_ROCMFP4/2, as the standalone ROCmFP4 HIP dequant kernel does, would avoid wrong block/qs indexing if the block layout changes.</comment>

<file context>
@@ -119,6 +120,72 @@ static __device__ void cpy_blck_q_f32(const char * cxi, char * cdsti) {
+        return;
+    }
+
+    const int64_t ib = packed_idx >> 4;
+    const int j = packed_idx & 0x0f;
+    const int64_t base = ib * QK_ROCMFP4;
</file context>

[cubic-dev-ai]

P3: Validation Script Index omits three Qwen scripts referenced in the validation examples below. The index lists check-rocmfpx-qwen-all.sh (core Qwen gates) but the validation examples on lines 103-105 use check-rocmfpx-qwen-coherency.sh, check-rocmfpx-qwen-bench.sh, and check-rocmfpx-qwen-strict-json.sh. Either add those three to the index or clarify the relationship (e.g. qwen-all.sh is a wrapper).

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/rocmfpx/README.md, line 59:

<comment>Validation Script Index omits three Qwen scripts referenced in the validation examples below. The index lists `check-rocmfpx-qwen-all.sh` (core Qwen gates) but the validation examples on lines 103-105 use `check-rocmfpx-qwen-coherency.sh`, `check-rocmfpx-qwen-bench.sh`, and `check-rocmfpx-qwen-strict-json.sh`. Either add those three to the index or clarify the relationship (e.g. qwen-all.sh is a wrapper).</comment>

<file context>
@@ -0,0 +1,106 @@
+
+```text
+scripts/check-rocmfpx-reference.sh        # CPU reference math
+scripts/check-rocmfpx-qwen-all.sh         # core Qwen gates
+scripts/check-rocmfpx-all.sh              # qwen-all + optional smokes
+scripts/check-rocmfpx-summary.sh          # full JSON summary runner
</file context>

[cubic-dev-ai]

P3: MSE monotonicity is verified for FP8 < FP6 < FP3, but there's no assertion that FP3 < FP2 (i.e., 3-bit rounds more accurately than 2-bit). Add assert(mse3 < mse2) after the existing assertions so all four bit depths are checked.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/rocmfpx/test_rocmfpx.c, line 157:

<comment>MSE monotonicity is verified for FP8 < FP6 < FP3, but there's no assertion that FP3 < FP2 (i.e., 3-bit rounds more accurately than 2-bit). Add `assert(mse3 < mse2)` after the existing assertions so all four bit depths are checked.</comment>

<file context>
@@ -0,0 +1,164 @@
+    assert(isfinite(mse3));
+    assert(isfinite(mse6));
+    assert(isfinite(mse8));
+    assert(mse8 < mse6);
+    assert(mse6 < mse3);
+
</file context>

[cubic-dev-ai]

P3: The new standalone HIP dequant helpers are currently dead code because rocmfp4_hip.cu is not wired into any ggml build target. If these are intended for integration tests or future fused paths, consider adding the file to the ROCm/HIP target or moving it behind a tracked follow-up so it does not silently drift.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/rocmfp4/rocmfp4_hip.cu, line 63:

<comment>The new standalone HIP dequant helpers are currently dead code because `rocmfp4_hip.cu` is not wired into any ggml build target. If these are intended for integration tests or future fused paths, consider adding the file to the ROCm/HIP target or moving it behind a tracked follow-up so it does not silently drift.</comment>

<file context>
@@ -0,0 +1,85 @@
+    }
+}
+
+extern "C" void rocmfp4_hip_dequantize_q4_0_to_f32(
+        const void  * src,
+        float       * dst,
</file context>

[cubic-dev-ai]

P3: ROCmFP4 scale decode now violates the format rules documented for this directory: CPU decode is supposed to use the arithmetic UE4M3 decoder, not a LUT. Keeping the documented decoder avoids divergence from the HIP/Vulkan contract and the promoted regression assumptions.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/deps/llama.cpp/ggml/rocmfp4/rocmfp4.c, line 59:

<comment>ROCmFP4 scale decode now violates the format rules documented for this directory: CPU decode is supposed to use the arithmetic UE4M3 decoder, not a LUT. Keeping the documented decoder avoids divergence from the HIP/Vulkan contract and the promoted regression assumptions.</comment>

<file context>
@@ -0,0 +1,797 @@
+#define ROCMFP4_SCALE_E14(M) ((8 + (M)) * 0x1p3f)
+#define ROCMFP4_SCALE_E15(M) ((8 + (M)) * 0x1p4f)
+
+static const float rocmfp4_scale_ue4m3_half[127] = {
+    ROCMFP4_SCALE_SUB(0), ROCMFP4_SCALE_SUB(1), ROCMFP4_SCALE_SUB(2), ROCMFP4_SCALE_SUB(3),
+    ROCMFP4_SCALE_SUB(4), ROCMFP4_SCALE_SUB(5), ROCMFP4_SCALE_SUB(6), ROCMFP4_SCALE_SUB(7),
</file context>