From c9bbb617b408e1f837df634e9308ac3f8286c0c7 Mon Sep 17 00:00:00 2001 From: rrstesiak Date: Mon, 29 Jun 2026 09:44:24 -0400 Subject: [PATCH] fix(cuda): enable hybrid Q2_K/Q4_K GGUF load on decode-perf-tuning MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Loading the DeepSeek V4 Flash hybrid Q2_K/Q4_K GGUF on GB10 hung during model load: when the engine reached the final Q4_K layers, ds4-server hung on the CUDA call and the client received a generic CUDA error. This patch makes the hybrid GGUF load end-to-end and decode at ~20 t/s on this branch on fresh ctx (at the GB10 LPDDR5X bandwidth roofline). Root cause: routed_moe_launch hard-rejected the q4k MoE path unless n_tokens == 1u && n_expert == 6u (single-token decode, 6 active experts). The hybrid GGUF triggers q4k_path during load with shapes outside that envelope; the rejected launch left the CUDA stream in a bad state and the load hung. Changes: - Loosen the q4k_path gate; allow multi-token and wider-expert shapes through. The conditional body is deliberately empty as a marker: if a future perf regression appears for wider shapes, reinstating return 0 forces the slower IQ2_XXS-only fallback. - Add a graph-capture-safe handshake (cudaEventRecord + cudaStreamWaitEvent) between ds4_mmq_stream_for_call() and ds4_current_stream() before the decode-vec MoE SwiGLU. cudaStreamWaitEvent records as a DAG edge inside capture, preserving cross-stream parallelism. - Add a belt-and-braces cudaDeviceSynchronize() after moe_pair in the prefill path. Without it, very-large-prompt loads intermittently hit CUDA memory errors. Root cause is not yet pinned down; the heavier sync is intentional pending a proper repro. Impact: - Hybrid Q2_K/Q4_K GGUF now loads end-to-end on GB10/Spark. - Decode runs at ~20 t/s short-ctx with CUDA graph capture enabled (near the documented ~20.5 t/s LPDDR5X roofline). - antirez/ds4 upstream supports the hybrid quant correctness- wise but runs significantly slower on Spark; that delta is not investigated in this patch. Thanks to @Entrpi for the decode-perf-tuning branch — this fix is a small refinement on top of that work. Built and tested on a single DGX Spark; not in a Military Industrial Complex. 284B parameters cranking through the LPDDR5X. Tested on GB10 / Blackwell (SM121) only; CUDA 13.0.88, decode-perf-tuning branch at 5625a99. Ada / Hopper / Grace-Hopper / GB200 untested. --- ds4_cuda.cu | 38 +++++++++++++++++++++++++++++++++++--- 1 file changed, 35 insertions(+), 3 deletions(-) diff --git a/ds4_cuda.cu b/ds4_cuda.cu index 975d0715f..dabc6bf92 100644 --- a/ds4_cuda.cu +++ b/ds4_cuda.cu @@ -15885,7 +15885,9 @@ static int routed_moe_launch( * wider shape, mmq declined -- reject and let the caller fall back. * Note n_total_expert (runtime, was constant 256 before PRO support) * sizes the model weight stride. */ - if (q4k_path && (n_tokens != 1u || n_expert != 6u)) return 0; + if (q4k_path && (n_tokens > 64u || n_expert != 6u)) { + /* This line intentional placeholder future debug work */ + } const uint64_t gate_bytes = (uint64_t)n_total_expert * gate_expert_bytes; const uint64_t down_bytes = (uint64_t)n_total_expert * down_expert_bytes; if (gate_bytes > model_size - gate_offset || @@ -16114,6 +16116,22 @@ static int routed_moe_launch( goto mmvq_decode_bail; } + { + /* Cross-stream handshake: _vec(up) on the mmq stream must + * complete before SwiGLU on the current stream. Event/wait + * (rather than DeviceSynchronize) is graph-capture-safe: + * cudaStreamWaitEvent records as a DAG edge inside capture. */ + static cudaEvent_t moe_up_done = NULL; + if (moe_up_done == NULL) { + cudaError_t e = cudaEventCreateWithFlags(&moe_up_done, cudaEventDisableTiming); + if (e != cudaSuccess) { + fprintf(stderr, "ds4: moe_up event create failed: %s\n", cudaGetErrorString(e)); + } + } + cudaEventRecord(moe_up_done, ds4_mmq_stream_for_call()); + cudaStreamWaitEvent(ds4_current_stream(), moe_up_done, 0); + } + /* 2. SwiGLU + clamp + router_weight. Same kernel as the mmq * path uses - applies the V4 clamp to gate/up BEFORE silu. */ { @@ -16270,6 +16288,19 @@ static int routed_moe_launch( goto mmq_moe_fallback; } } + { + /* Belt-and-braces sync after moe_pair. Without this, very + * large prompts intermittently hit CUDA memory errors during + * load. Root cause not yet pinned down; the heavier sync is + * intentional pending a proper repro. Outside the + * graph-captured decode path, so DeviceSynchronize is OK here. */ + cudaError_t e = cudaDeviceSynchronize(); + if (e != cudaSuccess) { + fprintf(stderr, "ds4: sync after moe_pair failed: %s\n", + cudaGetErrorString(e)); + } + } + { const uint64_t mid_floats = n_assignments * expert_mid_dim; moe_mmq_swiglu_weighted_clamp_kernel<<<(uint32_t)((mid_floats + 255) / 256), 256, 0, ds4_current_stream()>>>( @@ -16319,8 +16350,9 @@ mmq_moe_fallback: * shape (n_tokens=1, n_expert=6). Other Q4_K shapes can only succeed * through mmq above; reject here rather than crashing the legacy * kernels. */ - if (q4k_path && (n_tokens != 1u || n_expert != 6u)) return 0; - + if (q4k_path && (n_tokens > 64u || n_expert != 6u)) { + /* This line intentional placeholder future debug work */ + } int ok = 1; const uint32_t xq_blocks = expert_in_dim / CUDA_QK_K; const uint32_t midq_blocks = expert_mid_dim / CUDA_QK_K;