perf(deepseek4): enable causal Flash prefill

[weicj]

Summary

Enable causal Flash prefill for DeepSeek4 ROCmFP2. The implementation defaults to an internal prefill batch of 128 tokens, the best passing prefill result from the sweep and workload comparison below; deployments can reduce it when tuning.

Stack: main -> #502 -> #503 -> #496 -> causal Flash prefill at chunk 128.

Changes

• Enable causal Flash prefill for the all-hot DeepSeek4 prefill path.
• Batch the causal Flash prefill path internally, with 128 tokens as the prefill-selected default. DFLASH_DS4_FLASH_PREFILL_CHUNK can reduce it.
• Route DeepSeek4 causal prefill through Flash attention while preserving per-token causal visibility.

Validation

Validation used DeepSeek-V4-Flash-ROCMFP2-STRIX.gguf on with ROCm 7.2.4, gfx1151, max_ctx=1048576, chunk=512, and fused decode.

Flash prefill internal batch	Prefill p50	Result
16	24.20 tok/s	pass
32	34.45 tok/s	pass
64	43.49 tok/s	pass
128	54.96 tok/s	default: fastest passing prefill
256	25.67 tok/s	rejected: semantic mismatch

The same model, GPU target, server arguments, and prompts are used for before (baseline), 64-token and 128-token comparison. The different workload shows how different prompt impact the prefill speed.

Workload	PP / TG	Baseline (prefill / decode)	Causal Flash, 64-token batch (prefill / decode)	Causal Flash, 128-token batch (prefill / decode; default)
Repeated hello text	4k / 128	20.84 / 21.2 tok/s	40.41 / 19.2 tok/s	60.45 / 19.5 tok/s
Repeated hello text	8k / 256	20.32 / 21.4 tok/s	40.37 / 18.1 tok/s	52.10 / 18.4 tok/s
Creative writing	1065 / 128	21.28 / 21.2 tok/s	22.87 / 19.8 tok/s	36.41 / 20.4 tok/s
Mathematics worksheet	932 / 128	21.32 / 21.2 tok/s	31.43 / 20.1 tok/s	38.68 / 20.5 tok/s
Bug-fix request with source	755 / 128	21.35 / 21.3 tok/s	33.24 / 20.1 tok/s	39.70 / 20.5 tok/s

The 4k throughput prompt repeats a small token pattern and concentrates MoE routing. Prose, mathematics, and code distribute tokens across more experts, which adds ROCmFP2 dequantization and routed-FFN work not accelerated by Flash attention; this accounts for their lower prefill rates than 60.45 tok/s.

All prose, mathematics, and code responses remained in the requested content type and contained no malformed text. Their generation cap was 128 tokens.

[cubic-dev-ai]

12 issues found across 48 files

Not reviewed (too large): server/src/deepseek4/deepseek4_graph.cpp (~2,569 lines), server/deps/llama.cpp/ggml/src/ggml-cpu/ops.cpp (~78 lines) - if these are generated or fixture files, add them to ignored paths to exclude them from future reviews.

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/CMakeLists.txt">

<violation number="1" location="server/deps/llama.cpp/ggml/src/CMakeLists.txt:211">
P2: The `rocmfpx` sources are being added to `ggml-base`, but the companion test file `ggml/rocmfpx/test_rocmfpx.c` is not wired into any CMake target or CTest suite. That leaves the test uncompiled and unrunnable in CI, so regressions in the ROCmFP2/FP3/FP6/FP8 quantization paths could go undetected. Consider adding an `add_executable` and `add_test` target for `test_rocmfpx.c` so it builds and runs alongside the rest of the test suite.</violation>
</file>

<file name="server/src/deepseek4/deepseek4_dspark.h">

<violation number="1" location="server/src/deepseek4/deepseek4_dspark.h:80">
P2: The documentation for `deepseek4_dspark_draft_forward` documents a `seed_tok` parameter that does not exist in the function signature, which is confusing for API consumers. The `noise_embed` parameter already encompasses the seed embedding, so the `seed_tok` doc line should be removed or updated to describe how the seed is passed via `noise_embed`.</violation>
</file>

<file name="server/docs/DS4.md">

<violation number="1" location="server/docs/DS4.md:138">
P1: The documented conversion command references the wrong converter script. `convert_dflash_to_gguf.py` hard-codes the Qwen draft architecture (`qwen35-dflash-draft`) and Qwen tensor mappings, while the DeepSeek4 DSpark loader (`load_deepseek4_dspark_drafter()`) requires `deepseek4-dflash-draft` and DeepSeek4 MLA/HC/MoE tensors. Running the documented command will produce a GGUF that the DS4 runtime rejects with an unexpected-architecture error. Replace this command with the correct DS4-specific converter, or add a DS4 mode to the converter script.</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: The `dequantize_row_rocmfp4_hip` and `dequantize_row_rocmfp4_fast_hip` wrappers round the element count `k` up to the next 256-element boundary and launch kernels that unconditionally read and write a full 256-element block. Since the kernels have no bounds argument and no conditional early return, a final partial block reads past the end of the source buffer and overwrites memory past the destination buffer. Established patterns in the same file avoid this either by passing an exact block count to the kernel (e.g. `dequantize_row_q4_0_cuda` → `dequantize_block_q4_0` with `nb32` and an `if (ib >= nb32) return;` guard), or by using the generic `dequantize_block_cont_cuda` path whose `dequantize_block` kernel checks `if (i00 >= ne00) return;`. The ROCmFP4 dedicated paths should add a tail-safe bounds check in the kernel, accept an element-limit argument, or at minimum assert `k % QK_K == 0` before rounding up.</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:99">
P1: `rocmfp4_ue4m3_to_fp32_half_finite` defaults to the non-LUT fallback branch (since `GGML_ROCMFP4_USE_SCALE_LUT` is 0 by default), and that branch does not guard against the reserved UE4M3 NaN byte `0x7f`. The LUT branch and companion `rocmfpx_ue4m3_to_fp32_finite` both map `0x7f` to `0.0f`, but the fallback branch silently decodes it as a large finite scale (`480.0f`). This produces divergent, corrupted output depending on build configuration if an invalid scale ever reaches the device. Add an explicit `if (x > 0x7e) return 0.0f;` guard to the `#else` fallback path to match the LUT branch and the CPU-side `rocmfp4_ue4m3_to_fp32_half`.</violation>
</file>

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

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml.c:724">
P1: New ROCmFP types are registered but lack quantize-dispatcher cases in `ggml_quantize_chunk`, so any standard quantization path targeting these types will hit `assert(false)` and abort at runtime. Consider adding explicit `case` labels for all six new formats and routing them to the appropriate `rocmfp4_quantize_*` or `rocmfpx_quantize_*` functions, or else guard callers so they do not reach the unsupported default path.</violation>
</file>

<file name="server/src/deepseek4/deepseek4_backend.cpp">

<violation number="1" location="server/src/deepseek4/deepseek4_backend.cpp:336">
P1: init() now tries a full model load on HIP before falling back to hybrid, but unpark() still unconditionally restores HIP models as hybrid. After a park/unpark cycle, a HIP server that was running in all-hot mode silently switches to hybrid mode, which disables the causal Flash prefill path and changes behavior. Consider making unpark() mirror init()'s placement logic—trying full load first and respecting DFLASH_DEEPSEEK4_FORCE_HYBRID—so that load mode remains stable across the object's lifetime.</violation>

<violation number="2" location="server/src/deepseek4/deepseek4_backend.cpp:638">
P1: Speculative feature capture blindly indexes `spec_cap` using `n_target_layers` without confirming the buffer actually contains that many floats per token. The underlying `deepseek4_step_layer_range` sizes `capture_out` as `capture_layer_ids.size() * n_embd * n_tokens`, but the backend assumes each row is `n_target_layers * n_embd`. If these two counts ever diverge, the insert-loop constructs iterators beyond `spec_cap.end()`, causing undefined behavior and potential memory corruption. Add a size check (e.g. `spec_cap.size() == (size_t) n_tok * feat_row`) before indexing, or clamp to `spec_cap.end()` to stay safe.</violation>

<violation number="3" location="server/src/deepseek4/deepseek4_backend.cpp:786">
P1: The speculative decode fast-path in generate_impl emits tokens without checking io.cancelled or enforcing req.budget_hook, unlike the normal do_decode loop which checks both on every token. A cancelled request or a request with a remaining-token budget can produce and emit speculative tokens that the normal path would have suppressed. Consider adding io.cancelled checks before emitting tokens and respecting the budget hook (or falling back to the regular decode loop) so speculative decoding does not bypass output controls.</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:18">
P1: Removing the `gqa_opt_applies` guard for D=512 exposes the MMA path to inputs it cannot handle. On GPUs with Turing tensor cores, a D=512 prefill that does not satisfy GQA optimizations now selects `BEST_FATTN_KERNEL_MMA_F16` instead of being reported as unsupported. The MMA template switch for 512 dimensions ultimately calls `GGML_ABORT("fatal error")` for the non-GQA path, turning what used to be a backend fallback into a process crash. Consider restoring the `if (!gqa_opt_applies) { return BEST_FATTN_KERNEL_NONE; }` guard inside the `case 512:` block, or providing a non-GQA MMA fallback so the dispatch is safe.</violation>

<violation number="2" location="server/deps/llama.cpp/ggml/src/ggml-cuda/fattn.cu:103">
P1: The custom fast-path matcher checks `ggml_is_contiguous(mask)` but only validates `nb[0]` for Q, K, V, and dst. The kernel directly treats these tensors as fully packed arrays (`q + ((h * n_tokens + t) * D)`, `k + r * D`, etc.), so any non-unit higher stride would silently access incorrect memory. Add `ggml_is_contiguous()` checks for Q, K, V, and dst (and sinks if applicable), or compute offsets using the actual `nb[]` strides.</violation>

<violation number="3" location="server/deps/llama.cpp/ggml/src/ggml-cuda/fattn.cu:123">
P1: The custom DS4 Flash Attention kernel requests unbounded dynamic shared memory (`n_kv * sizeof(float)`) without checking device limits or launch status. For contexts where `n_kv` exceeds a few thousand, the allocation will surpass typical per-block shared memory limits and cause an invalid kernel launch. The matcher returns `true` regardless, so no fallback kernel is attempted. Cap `n_kv` to a device-specific shared-memory budget, or check `cudaGetLastError()` before returning success.</violation>
</file>

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

Re-trigger cubic</sub>

[cubic-dev-ai]

P1: The documented conversion command references the wrong converter script. convert_dflash_to_gguf.py hard-codes the Qwen draft architecture (qwen35-dflash-draft) and Qwen tensor mappings, while the DeepSeek4 DSpark loader (load_deepseek4_dspark_drafter()) requires deepseek4-dflash-draft and DeepSeek4 MLA/HC/MoE tensors. Running the documented command will produce a GGUF that the DS4 runtime rejects with an unexpected-architecture error. Replace this command with the correct DS4-specific converter, or add a DS4 mode to the converter script.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/docs/DS4.md, line 138:

<comment>The documented conversion command references the wrong converter script. `convert_dflash_to_gguf.py` hard-codes the Qwen draft architecture (`qwen35-dflash-draft`) and Qwen tensor mappings, while the DeepSeek4 DSpark loader (`load_deepseek4_dspark_drafter()`) requires `deepseek4-dflash-draft` and DeepSeek4 MLA/HC/MoE tensors. Running the documented command will produce a GGUF that the DS4 runtime rejects with an unexpected-architecture error. Replace this command with the correct DS4-specific converter, or add a DS4 mode to the converter script.</comment>

<file context>
@@ -110,6 +110,62 @@ The runtime logs the chosen split with a `[deepseek4-split] auto-split:` banner.
+Example conversion with the DS4 MTP shard that contains the DSpark heads:
+
+```bash
+python server/scripts/convert_dflash_to_gguf.py \
+  /path/to/dflash-draft/model.safetensors \
+  /path/to/dflash-draft.gguf \
</file context>

[cubic-dev-ai]

P1: The dequantize_row_rocmfp4_hip and dequantize_row_rocmfp4_fast_hip wrappers round the element count k up to the next 256-element boundary and launch kernels that unconditionally read and write a full 256-element block. Since the kernels have no bounds argument and no conditional early return, a final partial block reads past the end of the source buffer and overwrites memory past the destination buffer. Established patterns in the same file avoid this either by passing an exact block count to the kernel (e.g. dequantize_row_q4_0_cuda → dequantize_block_q4_0 with nb32 and an if (ib >= nb32) return; guard), or by using the generic dequantize_block_cont_cuda path whose dequantize_block kernel checks if (i00 >= ne00) return;. The ROCmFP4 dedicated paths should add a tail-safe bounds check in the kernel, accept an element-limit argument, or at minimum assert k % QK_K == 0 before rounding up.

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>The `dequantize_row_rocmfp4_hip` and `dequantize_row_rocmfp4_fast_hip` wrappers round the element count `k` up to the next 256-element boundary and launch kernels that unconditionally read and write a full 256-element block. Since the kernels have no bounds argument and no conditional early return, a final partial block reads past the end of the source buffer and overwrites memory past the destination buffer. Established patterns in the same file avoid this either by passing an exact block count to the kernel (e.g. `dequantize_row_q4_0_cuda` → `dequantize_block_q4_0` with `nb32` and an `if (ib >= nb32) return;` guard), or by using the generic `dequantize_block_cont_cuda` path whose `dequantize_block` kernel checks `if (i00 >= ne00) return;`. The ROCmFP4 dedicated paths should add a tail-safe bounds check in the kernel, accept an element-limit argument, or at minimum assert `k % QK_K == 0` before rounding up.</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]

P1: rocmfp4_ue4m3_to_fp32_half_finite defaults to the non-LUT fallback branch (since GGML_ROCMFP4_USE_SCALE_LUT is 0 by default), and that branch does not guard against the reserved UE4M3 NaN byte 0x7f. The LUT branch and companion rocmfpx_ue4m3_to_fp32_finite both map 0x7f to 0.0f, but the fallback branch silently decodes it as a large finite scale (480.0f). This produces divergent, corrupted output depending on build configuration if an invalid scale ever reaches the device. Add an explicit if (x > 0x7e) return 0.0f; guard to the #else fallback path to match the LUT branch and the CPU-side rocmfp4_ue4m3_to_fp32_half.

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_scale.cuh, line 99:

<comment>`rocmfp4_ue4m3_to_fp32_half_finite` defaults to the non-LUT fallback branch (since `GGML_ROCMFP4_USE_SCALE_LUT` is 0 by default), and that branch does not guard against the reserved UE4M3 NaN byte `0x7f`. The LUT branch and companion `rocmfpx_ue4m3_to_fp32_finite` both map `0x7f` to `0.0f`, but the fallback branch silently decodes it as a large finite scale (`480.0f`). This produces divergent, corrupted output depending on build configuration if an invalid scale ever reaches the device. Add an explicit `if (x > 0x7e) return 0.0f;` guard to the `#else` fallback path to match the LUT branch and the CPU-side `rocmfp4_ue4m3_to_fp32_half`.</comment>

<file context>
@@ -0,0 +1,157 @@
+#if defined(GGML_USE_HIP) && GGML_ROCMFP4_USE_SCALE_LUT
+    return x <= 0x7e ? rocmfp4_scale_ue4m3_half_lut[x] : 0.0f;
+#else
+    const int exp = (x >> 3) & 0xF;
+    const int man = x & 0x7;
+
</file context>

Suggested change

	const int exp = (x >> 3) & 0xF;
	if (x > 0x7e) {
	return 0.0f;
	}
	const int exp = (x >> 3) & 0xF;

[cubic-dev-ai]

P1: New ROCmFP types are registered but lack quantize-dispatcher cases in ggml_quantize_chunk, so any standard quantization path targeting these types will hit assert(false) and abort at runtime. Consider adding explicit case labels for all six new formats and routing them to the appropriate rocmfp4_quantize_* or rocmfpx_quantize_* functions, or else guard callers so they do not reach the unsupported default 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.c, line 724:

<comment>New ROCmFP types are registered but lack quantize-dispatcher cases in `ggml_quantize_chunk`, so any standard quantization path targeting these types will hit `assert(false)` and abort at runtime. Consider adding explicit `case` labels for all six new formats and routing them to the appropriate `rocmfp4_quantize_*` or `rocmfpx_quantize_*` functions, or else guard callers so they do not reach the unsupported default path.</comment>

<file context>
@@ -719,6 +721,54 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = {
         .to_float                 = (ggml_to_float_t) dequantize_row_q8_0,
         .from_float_ref           = (ggml_from_float_t) quantize_row_q8_0_ref,
     },
+    [GGML_TYPE_Q4_0_ROCMFP4] = {
+        .type_name                = "q4_0_rocmfp4",
+        .blck_size                = QK_ROCMFP4,
</file context>

[cubic-dev-ai]

P1: Speculative feature capture blindly indexes spec_cap using n_target_layers without confirming the buffer actually contains that many floats per token. The underlying deepseek4_step_layer_range sizes capture_out as capture_layer_ids.size() * n_embd * n_tokens, but the backend assumes each row is n_target_layers * n_embd. If these two counts ever diverge, the insert-loop constructs iterators beyond spec_cap.end(), causing undefined behavior and potential memory corruption. Add a size check (e.g. spec_cap.size() == (size_t) n_tok * feat_row) before indexing, or clamp to spec_cap.end() to stay safe.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/src/deepseek4/deepseek4_backend.cpp, line 638:

<comment>Speculative feature capture blindly indexes `spec_cap` using `n_target_layers` without confirming the buffer actually contains that many floats per token. The underlying `deepseek4_step_layer_range` sizes `capture_out` as `capture_layer_ids.size() * n_embd * n_tokens`, but the backend assumes each row is `n_target_layers * n_embd`. If these two counts ever diverge, the insert-loop constructs iterators beyond `spec_cap.end()`, causing undefined behavior and potential memory corruption. Add a size check (e.g. `spec_cap.size() == (size_t) n_tok * feat_row`) before indexing, or clamp to `spec_cap.end()` to stay safe.</comment>

<file context>
@@ -556,13 +608,42 @@ int DeepSeek4Backend::do_prefill(const std::vector<int32_t> & tokens,
+                                            tokens.data() + i,
+                                            timing ? &step_tel : nullptr,
+                                            /*allow_decode_graph_reuse=*/true, hp);
+            if (hp && !spec_cap.empty()) {
+                const int feat_row = spec_drafter_->n_target_layers * w_.n_embd;
+                for (int t = 0; t < n_tok; t++)
</file context>

[cubic-dev-ai]

P1: The custom fast-path matcher checks ggml_is_contiguous(mask) but only validates nb[0] for Q, K, V, and dst. The kernel directly treats these tensors as fully packed arrays (q + ((h * n_tokens + t) * D), k + r * D, etc.), so any non-unit higher stride would silently access incorrect memory. Add ggml_is_contiguous() checks for Q, K, V, and dst (and sinks if applicable), or compute offsets using the actual nb[] strides.

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 103:

<comment>The custom fast-path matcher checks `ggml_is_contiguous(mask)` but only validates `nb[0]` for Q, K, V, and dst. The kernel directly treats these tensors as fully packed arrays (`q + ((h * n_tokens + t) * D)`, `k + r * D`, etc.), so any non-unit higher stride would silently access incorrect memory. Add `ggml_is_contiguous()` checks for Q, K, V, and dst (and sinks if applicable), or compute offsets using the actual `nb[]` strides.</comment>

<file context>
@@ -7,6 +7,133 @@
+        K->ne[2] != 1 || V->ne[2] != 1 ||
+        Q->ne[3] != 1 || K->ne[3] != 1 || V->ne[3] != 1 ||
+        dst->ne[0] != 512 || dst->ne[1] != Q->ne[2] || dst->ne[2] != Q->ne[1] ||
+        Q->nb[0] != (int64_t) sizeof(float) ||
+        K->nb[0] != (int64_t) sizeof(float) ||
+        V->nb[0] != (int64_t) sizeof(float) ||
</file context>

Suggested change

	Q->nb[0] != (int64_t) sizeof(float) ||
	!ggml_is_contiguous(Q) ||
	!ggml_is_contiguous(K) ||
	!ggml_is_contiguous(V) ||
	!ggml_is_contiguous(dst) ||

[cubic-dev-ai]

P1: The custom DS4 Flash Attention kernel requests unbounded dynamic shared memory (n_kv * sizeof(float)) without checking device limits or launch status. For contexts where n_kv exceeds a few thousand, the allocation will surpass typical per-block shared memory limits and cause an invalid kernel launch. The matcher returns true regardless, so no fallback kernel is attempted. Cap n_kv to a device-specific shared-memory budget, or check cudaGetLastError() before returning success.

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 123:

<comment>The custom DS4 Flash Attention kernel requests unbounded dynamic shared memory (`n_kv * sizeof(float)`) without checking device limits or launch status. For contexts where `n_kv` exceeds a few thousand, the allocation will surpass typical per-block shared memory limits and cause an invalid kernel launch. The matcher returns `true` regardless, so no fallback kernel is attempted. Cap `n_kv` to a device-specific shared-memory budget, or check `cudaGetLastError()` before returning success.</comment>

<file context>
@@ -7,6 +7,133 @@
+
+    cudaStream_t stream = ctx.stream();
+    dim3 grid((unsigned) n_tokens, (unsigned) n_heads, 1);
+    const size_t shmem = (size_t) n_kv * sizeof(float);
+    ds4_flash_attn_d512_f32_shared_kv_kernel<<<grid, 256, shmem, stream>>>(
+            (float *) dst->data,
</file context>

[cubic-dev-ai]

P1: Removing the gqa_opt_applies guard for D=512 exposes the MMA path to inputs it cannot handle. On GPUs with Turing tensor cores, a D=512 prefill that does not satisfy GQA optimizations now selects BEST_FATTN_KERNEL_MMA_F16 instead of being reported as unsupported. The MMA template switch for 512 dimensions ultimately calls GGML_ABORT("fatal error") for the non-GQA path, turning what used to be a backend fallback into a process crash. Consider restoring the if (!gqa_opt_applies) { return BEST_FATTN_KERNEL_NONE; } guard inside the case 512: block, or providing a non-GQA MMA fallback so the dispatch is safe.

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 18:

<comment>Removing the `gqa_opt_applies` guard for D=512 exposes the MMA path to inputs it cannot handle. On GPUs with Turing tensor cores, a D=512 prefill that does not satisfy GQA optimizations now selects `BEST_FATTN_KERNEL_MMA_F16` instead of being reported as unsupported. The MMA template switch for 512 dimensions ultimately calls `GGML_ABORT("fatal error")` for the non-GQA path, turning what used to be a backend fallback into a process crash. Consider restoring the `if (!gqa_opt_applies) { return BEST_FATTN_KERNEL_NONE; }` guard inside the `case 512:` block, or providing a non-GQA MMA fallback so the dispatch is safe.</comment>

<file context>
@@ -7,6 +7,133 @@
+    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
+        if (tid < stride) smem[tid] += smem[tid + stride];
+        __syncthreads();
+    }
+    return smem[0];
+}
</file context>

[cubic-dev-ai]

P2: The rocmfpx sources are being added to ggml-base, but the companion test file ggml/rocmfpx/test_rocmfpx.c is not wired into any CMake target or CTest suite. That leaves the test uncompiled and unrunnable in CI, so regressions in the ROCmFP2/FP3/FP6/FP8 quantization paths could go undetected. Consider adding an add_executable and add_test target for test_rocmfpx.c so it builds and runs alongside the rest of the test suite.

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/CMakeLists.txt, line 211:

<comment>The `rocmfpx` sources are being added to `ggml-base`, but the companion test file `ggml/rocmfpx/test_rocmfpx.c` is not wired into any CMake target or CTest suite. That leaves the test uncompiled and unrunnable in CI, so regressions in the ROCmFP2/FP3/FP6/FP8 quantization paths could go undetected. Consider adding an `add_executable` and `add_test` target for `test_rocmfpx.c` so it builds and runs alongside the rest of the test suite.</comment>

<file context>
@@ -205,6 +205,10 @@ add_library(ggml-base
+            ../rocmfp4/rocmfp4.c
+            ../rocmfp4/rocmfp4.h
+            ../rocmfpx/rocmfpx.c
+            ../rocmfpx/rocmfpx.h
             ggml-quants.h
             gguf.cpp)
</file context>

[cubic-dev-ai]

P2: The documentation for deepseek4_dspark_draft_forward documents a seed_tok parameter that does not exist in the function signature, which is confusing for API consumers. The noise_embed parameter already encompasses the seed embedding, so the seed_tok doc line should be removed or updated to describe how the seed is passed via noise_embed.

Prompt for AI agents

Check if this issue is valid — if so, understand the root cause and fix it. At server/src/deepseek4/deepseek4_dspark.h, line 80:

<comment>The documentation for `deepseek4_dspark_draft_forward` documents a `seed_tok` parameter that does not exist in the function signature, which is confusing for API consumers. The `noise_embed` parameter already encompasses the seed embedding, so the `seed_tok` doc line should be removed or updated to describe how the seed is passed via `noise_embed`.</comment>

<file context>
@@ -0,0 +1,141 @@
+// features. All host-side f32 for a simple v1 (GPU feature-ring plumbing can
+// come later).
+//
+//   seed_tok        : the committed anchor token (block position 0's embedding)
+//   noise_embed     : [n_embd * block_size] embeds of [seed]+[MASK]*(block_size-1)
+//   ctx_features    : [n_target_layers*n_embd * ctx_len] captured features,
</file context>