feat(ds4): Lucebox heterogeneous expert parallel + DSpark verify

[davide221]

Summary

• add the DeepSeek4 DSpark shared worker and retain its stream fallback
• add in-process heterogeneous expert parallelism for R9700 + Strix Halo
• partition routed experts by owner, submit both HIP backends before joining,
  and reduce owner results on device
• add q<=4 ROCmFPX/MMVQ specializations, fused HC/owner join, GPU argmax,
  pinned rollback, and DSpark context-KV reuse
• preserve one boundary-aware q=4 verification graph instead of splitting it
  into serialized q=3 + q=1 work
• add lossless GPU I32 repeat/concat for routing tensors and harden cleanup,
  error propagation, diagnostics, and documentation

Lucebox placement

• R9700 (hip:0): dense target path, selected hot experts, DSpark drafter
• Strix Halo (hip:1): remaining materialized expert weights
• one target KV/cache owner; no host IPC in the fastest expert path
• both GPU owners are submitted before the device-side join

The legacy DSpark IPC worker remains available, but the promoted throughput
profile uses the in-process drafter and expert owners.

Performance

Clean Release build for gfx1151;gfx1201, fixed q=4, deterministic 128-token
integer-list workload:

Path	Acceptance	Verify	Decode
generic boundary split (q=3 + q=1)	1.00	98.5 ms/step	37.5 tok/s
boundary-aware fused q=4	1.00	70.7 ms/step	50.8 tok/s

Steady repetitions after the fix were 50.7 and 50.8 tok/s; draft time was
6.1 ms and output hashes matched. The same 1.00 acceptance on both sides makes
this an execution-path gain, not an acceptance-rate gain.

This profile explicitly uses top-4 routed experts. It is an approximate policy,
and the number is workload-specific; this PR does not claim a held-out quality
or universal 50 tok/s result. Top-6 remains available for quality validation.

Validation

• clean dual-architecture HIP build with ROCm 7.2.4
• 20/20 CTest tests pass, including ROCmFPX, HIP tail, RMSNorm, DS4 HC,
  draft top-k, rollback, DeepSeek4 unit, and server unit tests
• end-to-end q=4 smoke test: four identical output hashes
• benchmark log: /home/lucebox5/pr505_prod_q4_boundary_fix_20260719.log

Stack

Base: codex/deepseek4-dspark-draft.

[cubic-dev-ai]

2 issues found across 39 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/src/common/backend_ipc.cpp">

<violation number="1" location="server/src/common/backend_ipc.cpp:327">
P2: DSpark shared `propose_shared_bidir` can observe stale header/payload data because its daemon still uses non-atomic header accesses while this path now relies on release/acquire publication. Use `backend_ipc_shared_payload_header_matches()` for request validation and `backend_ipc_publish_shared_payload_header()` after writing the response payload.</violation>
</file>

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

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/mmvq.cu:1735">
P0: Disabling the dedicated MoE kernel while shared-ID alignment remains enabled corrupts aligned q<=8 routes: the generic fallback does not decode the metadata format. Couple alignment to the dedicated-kernel dispatch, or leave IDs unencoded when that kernel is disabled.</violation>
</file>

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

Re-trigger cubic</sub>

[cubic-dev-ai]

P0: Disabling the dedicated MoE kernel while shared-ID alignment remains enabled corrupts aligned q<=8 routes: the generic fallback does not decode the metadata format. Couple alignment to the dedicated-kernel dispatch, or leave IDs unencoded when that kernel is disabled.

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/mmvq.cu, line 1735:

<comment>Disabling the dedicated MoE kernel while shared-ID alignment remains enabled corrupts aligned q<=8 routes: the generic fallback does not decode the metadata format. Couple alignment to the dedicated-kernel dispatch, or leave IDs unencoded when that kernel is disabled.</comment>

<file context>
@@ -1289,25 +1720,133 @@ static void mul_mat_vec_q_moe_launch(
+        const char * e = std::getenv("DFLASH_CUDA_MMVQ_MOE_COMPACT_MASKED_IDS");
+        return e && e[0] == '1' && e[1] == '\0';
+    }();
+    static const bool aligned_shared_ids = []() {
+        const char * e = std::getenv("DFLASH_CUDA_MMVQ_MOE_ALIGN_SHARED_IDS");
+        return e && e[0] == '1' && e[1] == '\0';
</file context>

[cubic-dev-ai]

All reported issues were addressed across 6 files (changes from recent commits).

<sub>Tip: Review your code locally with the cubic CLI to iterate faster.

Re-trigger cubic</sub>

[cubic-dev-ai]

4 issues found across 42 files (changes from recent commits).

Not reviewed (too large): server/src/deepseek4/deepseek4_graph.cpp (~2,716 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/ggml-cuda/ds4-indexer.cu">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-indexer.cu:185">
P2: F32 score inputs lose precision on the WMMA path because this conversion rounds every query value to F16 before the dot product. Either constrain the op to QAT/F16-exact queries or retain a full-F32 fallback for general callers so its documented F32 contract matches CPU.</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:479">
P2: Prefills with 65,536+ tokens fail to launch because `n_tokens` is used as `grid.y`, whose HIP limit is 65,535. Chunk tokens or flatten the token/tile indexing into `grid.x` so large-context prefill remains valid.</violation>
</file>

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

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml.c:5514">
P3: `keep_rows == INT_MIN` creates a DS4 attention op that the CUDA backend rejects during graph placement, since its indexed-mask path must negate this value. Reject this sentinel in the setter so invalid sparse configuration fails at construction.</violation>
</file>

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

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/quantize.cu:213">
P3: Ordinary MMQ activation quantization now executes the grouped-layout branch check and carries grouped address state in its hot kernel despite every normal launch passing zero. A separate grouped specialization/kernel would preserve the prior lean path and avoid possible occupancy/throughput regression outside the new grouped flow.</violation>
</file>

<sub>Tip: Review your code locally with the cubic CLI to iterate faster.

Re-trigger cubic</sub>

[cubic-dev-ai]

P2: F32 score inputs lose precision on the WMMA path because this conversion rounds every query value to F16 before the dot product. Either constrain the op to QAT/F16-exact queries or retain a full-F32 fallback for general callers so its documented F32 contract matches CPU.

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-indexer.cu, line 185:

<comment>F32 score inputs lose precision on the WMMA path because this conversion rounds every query value to F16 before the dot product. Either constrain the op to QAT/F16-exact queries or retain a full-F32 fallback for general callers so its documented F32 contract matches CPU.</comment>

<file context>
@@ -0,0 +1,409 @@
+            if (token < n_tokens) {
+                const float2 q_value = *reinterpret_cast<const float2 *>(
+                    q + ((size_t) token * n_head + h) * 128 + d);
+                value = __floats2half2_rn(q_value.x, q_value.y);
+            }
+            *reinterpret_cast<half2 *>(a_sh + row * 128 + d) = value;
</file context>

[cubic-dev-ai]

P2: Prefills with 65,536+ tokens fail to launch because n_tokens is used as grid.y, whose HIP limit is 65,535. Chunk tokens or flatten the token/tile indexing into grid.x so large-context prefill remains valid.

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

<comment>Prefills with 65,536+ tokens fail to launch because `n_tokens` is used as `grid.y`, whose HIP limit is 65,535. Chunk tokens or flatten the token/tile indexing into `grid.x` so large-context prefill remains valid.</comment>

<file context>
@@ -364,33 +476,54 @@ void ggml_cuda_op_ds4_hc(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
             const int pre_blocks = (n_embd + 255) / 256;
-            if (n_hc == 4) {
-                ds4_hc_pre_kernel_t<4><<<pre_blocks, 256, 0, stream>>>(
+            const dim3 grid(pre_blocks, n_tokens, 1);
+            if (n_hc == 4 && n_tokens >= 64) {
+                ds4_hc_pre_split_kernel_t<4><<<n_tokens, 256, 0, stream>>>(
</file context>

[cubic-dev-ai]

P3: keep_rows == INT_MIN creates a DS4 attention op that the CUDA backend rejects during graph placement, since its indexed-mask path must negate this value. Reject this sentinel in the setter so invalid sparse configuration fails at construction.

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

<comment>`keep_rows == INT_MIN` creates a DS4 attention op that the CUDA backend rejects during graph placement, since its indexed-mask path must negate this value. Reject this sentinel in the setter so invalid sparse configuration fails at construction.</comment>

<file context>
@@ -5453,6 +5500,58 @@ void ggml_flash_attn_ext_set_prec(
+    GGML_ASSERT(raw_window > 0 && raw_window <= 0xffff);
+    GGML_ASSERT(block_size > 0 && block_size <= 0xffff);
+    ggml_set_op_params_i32(a, 4, raw_rows);
+    ggml_set_op_params_i32(a, 5, keep_rows);
+    // Both values are DS4-local and bounded well below 16 bits. Packing them
+    // keeps the remaining op-parameter slots available for the exact RoPE
</file context>

Suggested change

	ggml_set_op_params_i32(a, 5, keep_rows);
	GGML_ASSERT(keep_rows != INT_MIN);
	ggml_set_op_params_i32(a, 5, keep_rows);

[cubic-dev-ai]

P3: Ordinary MMQ activation quantization now executes the grouped-layout branch check and carries grouped address state in its hot kernel despite every normal launch passing zero. A separate grouped specialization/kernel would preserve the prior lean path and avoid possible occupancy/throughput regression outside the new grouped flow.

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/quantize.cu, line 213:

<comment>Ordinary MMQ activation quantization now executes the grouped-layout branch check and carries grouped address state in its hot kernel despite every normal launch passing zero. A separate grouped specialization/kernel would preserve the prior lean path and avoid possible occupancy/throughput regression outside the new grouped flow.</comment>

<file context>
@@ -204,8 +205,21 @@ static __global__ void quantize_mmq_q8_1(
+    // kernel emits the same flattened-K Q8 layout as a materialized permute.
+    int64_t src_i00 = i00;
+    int64_t src_group_offset = 0;
+    if (group_width > 0) {
+        const int64_t group = i00 / group_width;
+        src_i00 = i00 - group * group_width;
</file context>

[cubic-dev-ai]

8 issues found across 42 files (changes from recent commits).

Not reviewed (too large): server/src/deepseek4/deepseek4_graph.cpp (~2,716 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/ggml-cuda/common.cuh">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/common.cuh:116">
P3: Long-context HIP argsort still uses its existing path because `GGML_CUDA_USE_HIPCUB` has no consumer or hipCUB include anywhere in llama.cpp. Wire this macro into the sort implementation (and its hipCUB include), or omit it until that implementation lands.</violation>
</file>

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

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cu:14">
P3: Owner-route compaction hint is currently dead: no producer attaches this magic payload to an ID tensor, so every call retains the full `ne_get_rows` allocation and quantization path. Add the matching owner-builder attachment/lifetime management, or remove this unreachable contract until both sides land together.</violation>
</file>

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

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-indexer.cu:318">
P1: Default CUDA builds retain Maxwell/Pascal targets, but these devices enter this WMMA path because they also have 32-thread warps; WMMA requires compute capability 7.0+. Gate the WMMA launch on NVIDIA capability as well, and use the scalar kernel below that threshold.</violation>
</file>

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

<violation number="1" location="server/src/deepseek4/deepseek4_backend.cpp:1311">
P1: Dense prefill with mixed hot/cold experts now processes default 512-token chunks through a hybrid path that does not perform per-token HC post-mixing, producing incorrect hidden states/logits. Keep hybrid prefill tokenwise until its batched HC path is selected unconditionally (the nearby comment already describes this fallback).</violation>
</file>

<file name="server/src/common/moe_hybrid_ffn_eval.cpp">

<violation number="1" location="server/src/common/moe_hybrid_ffn_eval.cpp:3332">
P1: Device-resident joins can consume stale cold partials when a layer/batch routes entirely to hot experts: no cold graph runs, leaving `device_outputs->cold` unchanged before HC post combines it. Initialize/zero the cold destination when `has_cold` is false, or make the downstream join omit it for this case.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml-backend-meta.cpp">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-backend-meta.cpp:868">
P1: Heterogeneously owned grouped-expert weights split on `src[0]` axis 1 abort during meta tensor initialization: `scalar_only=true` rejects that split. This op still has MUL_MAT's logical split semantics when grouped activations are mirrored, so route it through `handle_mul_mat` to preserve owner-local output slices.</violation>
</file>

<file name="server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py">

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/template-instances/generate_cu_files.py:36">
P2: SOURCE_MMQ template lacks the #define GGML_CUDA_ROCMFPX_MMQ_TILE 1 guard that the existing ROCmFP instance files require. Re-running generate_cu_files.py would produce files without the guard, silently changing RDNA4 launch bounds from minBlocks=2 to minBlocks=1 for these quantization types. Add a conditional define in the generation loop for ROCMFP types.</violation>
</file>

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

<violation number="1" location="server/deps/llama.cpp/ggml/src/ggml-cuda/vecdotq.cuh:432">
P1: The 2-byte load in rocmfpx_pack4_fp3_vec_cuda truncates the 4th 3-bit code when (3*base)&7 >= 5 (the 12-bit span crosses into a 3rd byte). Replace the 2-byte load with a 3-byte load to cover the full span, or guard the fast path so the fallback handles problematic bit offsets.</violation>
</file>

<sub>Tip: Review your code locally with the cubic CLI to iterate faster.

Re-trigger cubic</sub>

[cubic-dev-ai]

P1: Default CUDA builds retain Maxwell/Pascal targets, but these devices enter this WMMA path because they also have 32-thread warps; WMMA requires compute capability 7.0+. Gate the WMMA launch on NVIDIA capability as well, and use the scalar kernel below that threshold.

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-indexer.cu, line 318:

<comment>Default CUDA builds retain Maxwell/Pascal targets, but these devices enter this WMMA path because they also have 32-thread warps; WMMA requires compute capability 7.0+. Gate the WMMA launch on NVIDIA capability as well, and use the scalar kernel below that threshold.</comment>

<file context>
@@ -0,0 +1,409 @@
+    const int warp_size =
+        ggml_cuda_info().devices[ggml_cuda_get_device()].warp_size;
+#if DS4_INDEXER_WMMA_AVAILABLE
+    if (warp_size == 32) {
+        const dim3 grid((unsigned) ((n_comp + 127) / 128),
+                        (unsigned) ((n_tokens + 15) / 16), 1);
</file context>

[cubic-dev-ai]

P1: Dense prefill with mixed hot/cold experts now processes default 512-token chunks through a hybrid path that does not perform per-token HC post-mixing, producing incorrect hidden states/logits. Keep hybrid prefill tokenwise until its batched HC path is selected unconditionally (the nearby comment already describes this fallback).

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

<comment>Dense prefill with mixed hot/cold experts now processes default 512-token chunks through a hybrid path that does not perform per-token HC post-mixing, producing incorrect hidden states/logits. Keep hybrid prefill tokenwise until its batched HC path is selected unconditionally (the nearby comment already describes this fallback).</comment>

<file context>
@@ -1205,17 +1288,30 @@ bool DeepSeek4Backend::unpark(const std::string & what) {
+    // kernels. Smaller tail chunks use the same scheduler or its reference
+    // fallback.
+    const int layer_major_cap = DS4_MAX_LAYER_MAJOR_PREFILL_TOKENS;
+    const int chunk = !prefill_attention_mode_is_approximate(cfg_.prefill_mode)
+        ? 1
+        : std::max(1, std::min(requested_chunk,
</file context>

[cubic-dev-ai]

P1: Device-resident joins can consume stale cold partials when a layer/batch routes entirely to hot experts: no cold graph runs, leaving device_outputs->cold unchanged before HC post combines it. Initialize/zero the cold destination when has_cold is false, or make the downstream join omit it for this case.

Prompt for AI agents

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

<comment>Device-resident joins can consume stale cold partials when a layer/batch routes entirely to hot experts: no cold graph runs, leaving `device_outputs->cold` unchanged before HC post combines it. Initialize/zero the cold destination when `has_cold` is false, or make the downstream join omit it for this case.</comment>

<file context>
@@ -2448,6 +3184,190 @@ bool eval_moe_hybrid_ffn_batched(
+                n_tokens, cold_partial, &cold_err,
+                nullptr, p_cold_alloc, nullptr, nullptr, nullptr,
+                /*skip_cold=*/false, /*skip_hot=*/true, cur_backend,
+                device_join ? device_outputs->cold : nullptr,
+                device_join ? device_outputs->backend : nullptr);
+        });
</file context>

[cubic-dev-ai]

P1: Heterogeneously owned grouped-expert weights split on src[0] axis 1 abort during meta tensor initialization: scalar_only=true rejects that split. This op still has MUL_MAT's logical split semantics when grouped activations are mirrored, so route it through handle_mul_mat to preserve owner-local output slices.

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-backend-meta.cpp, line 868:

<comment>Heterogeneously owned grouped-expert weights split on `src[0]` axis 1 abort during meta tensor initialization: `scalar_only=true` rejects that split. This op still has MUL_MAT's logical split semantics when grouped activations are mirrored, so route it through `handle_mul_mat` to preserve owner-local output slices.</comment>

<file context>
@@ -846,6 +861,12 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(co
+                // The logical activation rows are assembled from a private
+                // grouped physical layout. Keep the operation local; only the
+                // CPU and CUDA/HIP backends implement that layout contract.
+                split_state = handle_generic(src_ss, /*scalar_only =*/ true);
+            } break;
             case GGML_OP_OUT_PROD: {
</file context>

[cubic-dev-ai]

P1: The 2-byte load in rocmfpx_pack4_fp3_vec_cuda truncates the 4th 3-bit code when (3*base)&7 >= 5 (the 12-bit span crosses into a 3rd byte). Replace the 2-byte load with a 3-byte load to cover the full span, or guard the fast path so the fallback handles problematic bit offsets.

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/vecdotq.cuh, line 432:

<comment>The 2-byte load in rocmfpx_pack4_fp3_vec_cuda truncates the 4th 3-bit code when (3*base)&7 >= 5 (the 12-bit span crosses into a 3rd byte). Replace the 2-byte load with a 3-byte load to cover the full span, or guard the fast path so the fallback handles problematic bit offsets.</comment>

<file context>
@@ -425,12 +425,14 @@ static __device__ __forceinline__ int rocmfpx_pack4_fp3_bits12_vec_cuda(const ui
+    // adjacent bytes. Decode them with the same lookup used by MMVQ.
+    const int bit_pos = 3 * base;
+    const int byte_pos = bit_pos >> 3;
+    const uint32_t packed = (uint32_t) qs[byte_pos] |
+                            ((uint32_t) qs[byte_pos + 1] << 8);
+    return rocmfpx_pack4_fp3_bits12_vec_cuda(
</file context>

[cubic-dev-ai]

P2: SOURCE_MMQ template lacks the #define GGML_CUDA_ROCMFPX_MMQ_TILE 1 guard that the existing ROCmFP instance files require. Re-running generate_cu_files.py would produce files without the guard, silently changing RDNA4 launch bounds from minBlocks=2 to minBlocks=1 for these quantization types. Add a conditional define in the generation loop for ROCMFP types.

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/template-instances/generate_cu_files.py, line 36:

<comment>SOURCE_MMQ template lacks the #define GGML_CUDA_ROCMFPX_MMQ_TILE 1 guard that the existing ROCmFP instance files require. Re-running generate_cu_files.py would produce files without the guard, silently changing RDNA4 launch bounds from minBlocks=2 to minBlocks=1 for these quantization types. Add a conditional define in the generation loop for ROCMFP types.</comment>

<file context>
@@ -33,6 +33,7 @@
 
 TYPES_MMQ = [
     "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0",
+    "GGML_TYPE_Q4_0_ROCMFP4_FAST", "GGML_TYPE_Q2_0_ROCMFP2", "GGML_TYPE_Q3_0_ROCMFPX",
     "GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K",
     "GGML_TYPE_IQ2_XXS", "GGML_TYPE_IQ2_XS", "GGML_TYPE_IQ2_S", "GGML_TYPE_IQ3_XXS", "GGML_TYPE_IQ3_S",
</file context>

[cubic-dev-ai]

P3: Long-context HIP argsort still uses its existing path because GGML_CUDA_USE_HIPCUB has no consumer or hipCUB include anywhere in llama.cpp. Wire this macro into the sort implementation (and its hipCUB include), or omit it until that implementation lands.

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/common.cuh, line 116:

<comment>Long-context HIP argsort still uses its existing path because `GGML_CUDA_USE_HIPCUB` has no consumer or hipCUB include anywhere in llama.cpp. Wire this macro into the sort implementation (and its hipCUB include), or omit it until that implementation lands.</comment>

<file context>
@@ -109,6 +109,13 @@
+// single-block bitonic argsort cannot represent DS4's long-context indexer once
+// the number of compressed rows exceeds 1024, so use hipCUB for that case.
+#if defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+#    define GGML_CUDA_USE_HIPCUB
+#endif  // defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
+
</file context>

[cubic-dev-ai]

P3: Owner-route compaction hint is currently dead: no producer attaches this magic payload to an ID tensor, so every call retains the full ne_get_rows allocation and quantization path. Add the matching owner-builder attachment/lifetime management, or remove this unreachable contract until both sides land together.

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/mmq.cu, line 14:

<comment>Owner-route compaction hint is currently dead: no producer attaches this magic payload to an ID tensor, so every call retains the full `ne_get_rows` allocation and quantization path. Add the matching owner-builder attachment/lifetime management, or remove this unreachable contract until both sides land together.</comment>

<file context>
@@ -3,6 +3,16 @@
+    int32_t  max_expert_rows;
+    int32_t  live_routes;
+};
+static constexpr uint32_t MMID_OWNER_GRID_HINT_MAGIC = 0x4D4F4752u; // "MOGR"
+
 static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) {
</file context>