feat(ggml): vendor ROCmFPX quant support for DeepSeek V4 Flash by davide221 · Pull Request #502 · Luce-Org/lucebox

davide221 · 2026-07-08T13:47:38Z

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.

Ports the ROCmFP4/ROCmFPX ggml work (previously lucebox-ggml #36, on the old submodule pointer) into the vendored server/deps/llama.cpp tree. - ROCmFP4/ROCmFPX quant types + CPU reference conversions (ggml/rocmfp4, ggml/rocmfpx) and the ggml type-trait registrations. - 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) with the register-resident sinkhorn kernel; inert unless emitted by the DS4 fused-decode path. - DS4 SwiGLU split op plumbing for the fused FFN matvec paths. Layered on current main, so the main-side ggml work is preserved (fp64 RoPE reduction, #497 RDNA MMQ tile, LUCE_MMQ_DP_MAX_NE1, MMVQ_MAX_MOE_BATCH_SIZE, fused dual set_rows, raw-span guard). 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[] to avoid a stack over-read of the last window (bit-identical; the over-read bits were already masked out).

The non-HIP fallback of rocmfp4_get_int_from_codebook_16 / rocmfp4_get_low_int_from_codebook_16 called get_int_from_table_16, which is defined in vecdotq.cuh. TUs that pull in this header without vecdotq.cuh (fattn-chunked.cu reaches it via the fattn dequant chain) failed to compile under nvcc: rocmfp4_hip_codebook.cuh: error: identifier "get_int_from_table_16" is undefined The HIP path never hit this (it uses __builtin_amdgcn_perm), so the ROCm CI and the Strix build stayed green while the sm_86 CUDA build broke. Fix: inline the generic table expander (the generic branch of get_int_from_table_16, verbatim) as a static helper in this header, so the fallback no longer depends on include order. Bit-identical; the HIP hot path is unchanged.

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>

_{Reply with feedback, questions, or to request a fix.

Re-trigger cubic}

cubic-dev-ai · 2026-07-09T08:27:37Z

+    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];


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 · 2026-07-09T08:27:37Z

            {
                ggml_compute_forward_swiglu_oai(params, dst);
            } break;
+        case GGML_GLU_OP_SWIGLU_DS4:


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 · 2026-07-09T08:27:37Z

+            n_tokens,
+            n_heads,
+            n_kv);
+    return true;


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 · 2026-07-09T08:27:37Z

+
+    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);


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 · 2026-07-09T08:27:37Z


+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;


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 · 2026-07-09T08:27:38Z

@@ -1,6 +1,7 @@
 #include "cpy.cuh"


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 · 2026-07-09T08:27:38Z

+
+```text
+scripts/check-rocmfpx-reference.sh        # CPU reference math
+scripts/check-rocmfpx-qwen-all.sh         # core Qwen gates


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 · 2026-07-09T08:27:38Z

+    assert(isfinite(mse3));
+    assert(isfinite(mse6));
+    assert(isfinite(mse8));
+    assert(mse8 < mse6);


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 · 2026-07-09T08:27:38Z

@@ -0,0 +1,85 @@
+#include "rocmfp4.h"


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 · 2026-07-09T08:27:38Z

+#define ROCMFP4_SCALE_E14(M) ((8 + (M)) * 0x1p3f)
+#define ROCMFP4_SCALE_E15(M) ((8 + (M)) * 0x1p4f)
+
+static const float rocmfp4_scale_ue4m3_half[127] = {


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>

This was referenced Jul 8, 2026

fix(ds4): ROCmFPX byte-safe decode + fused-decode path #503

Open

fix(ds4): keep ROCmFPX generation byte-safe #494

Closed

feat(ggml): add ROCmFPX quant support for DS4 Luce-Org/lucebox-ggml#36

Closed

davide221 added 2 commits July 8, 2026 17:23

davide221 force-pushed the codex/rocmfpx-ggml-vendor branch from 6fe5e71 to a5f1146 Compare July 8, 2026 15:46

davide221 mentioned this pull request Jul 8, 2026

feat(dflash): add direct DeepSeek4 pure-UMA runtime path #489

Open

davide221 marked this pull request as ready for review July 9, 2026 08:10

cubic-dev-ai Bot reviewed Jul 9, 2026

View reviewed changes

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

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

feat(ggml): vendor ROCmFPX quant support for DeepSeek V4 Flash#502
davide221 wants to merge 2 commits into
mainfrom
codex/rocmfpx-ggml-vendor

davide221 commented Jul 8, 2026 •

edited by cubic-dev-ai Bot

Loading

Uh oh!

cubic-dev-ai Bot left a comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

Uh oh!

Conversation

davide221 commented Jul 8, 2026 • edited by cubic-dev-ai Bot Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Layered on current main (no reverts)

Review fixes on top of #36

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

Validation

Uh oh!

cubic-dev-ai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

cubic-dev-ai Bot Jul 9, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

davide221 commented Jul 8, 2026 •

edited by cubic-dev-ai Bot

Loading