From 8992588ea9de1dad4b96b7bc65a096b9b5a0e57c Mon Sep 17 00:00:00 2001 From: cheese-cakee Date: Wed, 19 Aug 2026 03:07:26 +0530 Subject: [PATCH 1/4] perf(ds4): elide intermediate prefill logits and tune gfx1151 mmq tiles --- .../deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh | 21 ++++++++++++++++--- server/src/deepseek4/deepseek4_backend.cpp | 12 +++++++---- server/src/deepseek4/deepseek4_graph.cpp | 19 +++++++++++++---- server/src/deepseek4/deepseek4_internal.h | 3 ++- 4 files changed, 43 insertions(+), 12 deletions(-) diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh index dd838f237..b67a65e8c 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh @@ -122,6 +122,9 @@ struct tile_x_sizes { static int get_mmq_x_max_host(const int cc) { if (LUCEBOX_RDNA_TILE_HOST(cc)) { + if (GGML_CUDA_CC_IS_RDNA3(cc)) { + return 48; + } #if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 64; #else @@ -139,7 +142,9 @@ static int get_mmq_x_max_host(const int cc) { static constexpr __device__ int get_mmq_x_max_device() { #if LUCEBOX_RDNA_TILE_DEVICE -#if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) +#if defined(RDNA3) + return 48; +#elif defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 64; #else return 128; @@ -169,6 +174,9 @@ static constexpr __device__ int get_mmq_x_max_device() { static int get_mmq_y_host(const int cc) { if (LUCEBOX_RDNA_TILE_HOST(cc)) { + if (GGML_CUDA_CC_IS_RDNA3(cc)) { + return 64; + } #if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 64; #else @@ -189,7 +197,9 @@ static constexpr __device__ int get_iter_k([[maybe_unused]] const ggml_type type static constexpr __device__ int get_mmq_y_device() { #if LUCEBOX_RDNA_TILE_DEVICE -#if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) +#if defined(RDNA3) + return 64; +#elif defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 64; #else return 128; @@ -342,6 +352,9 @@ static constexpr __device__ int mmq_get_granularity_device(const int /*mmq_x*/) #if defined(GGML_USE_HIP) static int mmq_get_nwarps_host(const int cc, const int warp_size) { if (LUCEBOX_RDNA_TILE_HOST(cc)) { + if (GGML_CUDA_CC_IS_RDNA3(cc)) { + return 4; + } #if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 4; #else @@ -358,7 +371,9 @@ static int mmq_get_nwarps_host(const int /*cc*/, const int warp_size) { static constexpr __device__ int mmq_get_nwarps_device() { #if LUCEBOX_RDNA_TILE_DEVICE -#if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) +#if defined(RDNA3) + return 4; +#elif defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 4; #else return 8; diff --git a/server/src/deepseek4/deepseek4_backend.cpp b/server/src/deepseek4/deepseek4_backend.cpp index d24bb282f..4834cd2d2 100644 --- a/server/src/deepseek4/deepseek4_backend.cpp +++ b/server/src/deepseek4/deepseek4_backend.cpp @@ -1813,6 +1813,7 @@ int DeepSeek4Backend::do_prefill(const std::vector & tokens, DeepSeek4StepTelemetry step_tel; if (timing) step_tel.embed_us = elapsed_us(embed_t0, Clock::now()); + const bool need_logits = (i + n_tok >= n_total); std::vector logits; bool ok = false; std::vector hc_state; @@ -1848,7 +1849,7 @@ int DeepSeek4Backend::do_prefill(const std::vector & tokens, ok = deepseek4_step_layer_range( backend_, cfg_.device.gpu, w_, cache_, hc_state, embed.data(), n_tok, pos, - 0, w_.n_layer, &logits, + 0, w_.n_layer, need_logits ? &logits : nullptr, tokens.data() + i, timing ? &step_tel : nullptr, /*allow_decode_graph_reuse=*/true, hp, @@ -1862,11 +1863,12 @@ int DeepSeek4Backend::do_prefill(const std::vector & tokens, timing ? &step_tel : nullptr, routing_stats_.get(), hp, - expert_runtime_.compute ? &expert_runtime_ : nullptr); + expert_runtime_.compute ? &expert_runtime_ : nullptr, + need_logits); } else { ok = deepseek4_step_layer_range(backend_, cfg_.device.gpu, w_, cache_, hc_state, embed.data(), n_tok, pos, - 0, w_.n_layer, &logits, + 0, w_.n_layer, need_logits ? &logits : nullptr, tokens.data() + i, timing ? &step_tel : nullptr, cfg_.prefill_mode != PrefillAttentionMode::Sparse, hp); @@ -1893,7 +1895,9 @@ int DeepSeek4Backend::do_prefill(const std::vector & tokens, add_step_tel(tel_acc, step_tel); steps++; } - last_logits_ = std::move(logits); + if (need_logits) { + last_logits_ = std::move(logits); + } pos += n_tok; last_logits_pos_ = cache_.cur_pos; i += n_tok; diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index cf4d4a9a2..4ab3418d7 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -4079,7 +4079,8 @@ static bool deepseek4_step_hybrid( MoeHybridStreamEngine * stream_engine, DeepSeek4StepTelemetry * telemetry, MoeHybridRoutingStats * routing_stats, - MoeExpertComputeRuntime * expert_runtime) { + MoeExpertComputeRuntime * expert_runtime, + bool need_logits = true) { const auto step_t0 = Ds4TimingClock::now(); const int n_embd = w.n_embd; const int n_hc = w.n_hc; @@ -4495,6 +4496,15 @@ static bool deepseek4_step_hybrid( if (hot_alloc) ggml_gallocr_free(hot_alloc); if (cold_alloc) ggml_gallocr_free(cold_alloc); + if (!need_logits) { + out_logits.clear(); + cache.cur_pos = kv_start + n_tokens; + if (telemetry) { + telemetry->total_us += ds4_elapsed_us(step_t0, Ds4TimingClock::now()); + } + return true; + } + // ── Output HC pre → norm → logits ─────────────────────────────────── const auto output_t0 = Ds4TimingClock::now(); std::vector final_embd((size_t)n_embd * (size_t)n_tokens); @@ -4576,7 +4586,8 @@ bool deepseek4_step( DeepSeek4StepTelemetry * telemetry, MoeHybridRoutingStats * routing_stats, Ds4VerifyHooks * verify_hooks, - MoeExpertComputeRuntime * expert_runtime) { + MoeExpertComputeRuntime * expert_runtime, + bool need_logits) { if (w.moe_hybrid && moe_hybrid != nullptr) { if (!deepseek4_cuda_hc_set_device(device)) { std::fprintf(stderr, @@ -4587,13 +4598,13 @@ bool deepseek4_step( return deepseek4_step_hybrid(backend, w, cache, *moe_hybrid, embed, n_tokens, kv_start, out_logits, token_ids, stream_engine, telemetry, routing_stats, - expert_runtime); + expert_runtime, need_logits); } std::vector hc_state; return deepseek4_step_layer_range( backend, device, w, cache, hc_state, embed, n_tokens, kv_start, - 0, w.n_layer, &out_logits, token_ids, telemetry, + 0, w.n_layer, need_logits ? &out_logits : nullptr, token_ids, telemetry, /*allow_decode_graph_reuse=*/verify_hooks == nullptr, verify_hooks, /*moe_hybrid=*/nullptr, expert_runtime, routing_stats); } diff --git a/server/src/deepseek4/deepseek4_internal.h b/server/src/deepseek4/deepseek4_internal.h index eea667423..40896cce7 100644 --- a/server/src/deepseek4/deepseek4_internal.h +++ b/server/src/deepseek4/deepseek4_internal.h @@ -377,7 +377,8 @@ bool deepseek4_step( DeepSeek4StepTelemetry * telemetry = nullptr, MoeHybridRoutingStats * routing_stats = nullptr, Ds4VerifyHooks * verify_hooks = nullptr, - MoeExpertComputeRuntime * expert_runtime = nullptr); + MoeExpertComputeRuntime * expert_runtime = nullptr, + bool need_logits = true); // Optional hooks for the DSpark spec-decode batched verify (deepseek4_dspark). // When set on a multi-token deepseek4_step_layer_range call they add: per-layer From b4349778e5d66ca552f2cc249d7102b3c81eea42 Mon Sep 17 00:00:00 2001 From: cheese-cakee Date: Wed, 19 Aug 2026 16:35:21 +0530 Subject: [PATCH 2/4] fix(ds4): address review comments for prefill readout elision and mmq tuning --- .../deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh | 31 +++++-------------- server/src/deepseek4/deepseek4_backend.cpp | 3 +- server/src/deepseek4/deepseek4_graph.cpp | 26 ++++++++-------- 3 files changed, 23 insertions(+), 37 deletions(-) diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh index b67a65e8c..a9f9e836f 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh @@ -113,8 +113,8 @@ struct tile_x_sizes { #ifndef LUCEBOX_RDNA_MMQ_TILE_OVERRIDE #define LUCEBOX_RDNA_MMQ_TILE_OVERRIDE 1 #endif -#define LUCEBOX_RDNA_TILE_HOST(cc) (LUCEBOX_RDNA_MMQ_TILE_OVERRIDE && (GGML_CUDA_CC_IS_RDNA3(cc) || GGML_CUDA_CC_IS_RDNA4(cc))) -#if LUCEBOX_RDNA_MMQ_TILE_OVERRIDE && (defined(RDNA3) || defined(RDNA4)) +#define LUCEBOX_RDNA_TILE_HOST(cc) (LUCEBOX_RDNA_MMQ_TILE_OVERRIDE && (GGML_CUDA_CC_IS_RDNA3_5(cc) || GGML_CUDA_CC_IS_RDNA4(cc))) +#if LUCEBOX_RDNA_MMQ_TILE_OVERRIDE && (defined(RDNA3_5) || defined(RDNA4)) #define LUCEBOX_RDNA_TILE_DEVICE 1 #else #define LUCEBOX_RDNA_TILE_DEVICE 0 @@ -122,7 +122,7 @@ struct tile_x_sizes { static int get_mmq_x_max_host(const int cc) { if (LUCEBOX_RDNA_TILE_HOST(cc)) { - if (GGML_CUDA_CC_IS_RDNA3(cc)) { + if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { return 48; } #if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) @@ -142,7 +142,7 @@ static int get_mmq_x_max_host(const int cc) { static constexpr __device__ int get_mmq_x_max_device() { #if LUCEBOX_RDNA_TILE_DEVICE -#if defined(RDNA3) +#if defined(RDNA3_5) return 48; #elif defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 64; @@ -153,28 +153,13 @@ static constexpr __device__ int get_mmq_x_max_device() { #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) return 128; #else // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) - -#if defined(GGML_USE_HIP) return 64; -#else // defined(GGML_USE_HIP) - -#if __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA -#ifdef GGML_CUDA_FORCE_MMQ - return 128; -#else // GGML_CUDA_FORCE_MMQ - return MMQ_DP4A_MAX_BATCH_SIZE; -#endif // GGML_CUDA_FORCE_MMQ -#else // __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA - return 64; -#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA - -#endif // defined(GGML_USE_HIP) #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } static int get_mmq_y_host(const int cc) { if (LUCEBOX_RDNA_TILE_HOST(cc)) { - if (GGML_CUDA_CC_IS_RDNA3(cc)) { + if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { return 64; } #if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) @@ -197,7 +182,7 @@ static constexpr __device__ int get_iter_k([[maybe_unused]] const ggml_type type static constexpr __device__ int get_mmq_y_device() { #if LUCEBOX_RDNA_TILE_DEVICE -#if defined(RDNA3) +#if defined(RDNA3_5) return 64; #elif defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 64; @@ -352,7 +337,7 @@ static constexpr __device__ int mmq_get_granularity_device(const int /*mmq_x*/) #if defined(GGML_USE_HIP) static int mmq_get_nwarps_host(const int cc, const int warp_size) { if (LUCEBOX_RDNA_TILE_HOST(cc)) { - if (GGML_CUDA_CC_IS_RDNA3(cc)) { + if (GGML_CUDA_CC_IS_RDNA3_5(cc)) { return 4; } #if defined(GGML_CUDA_ROCMFPX_MMQ_TILE) @@ -371,7 +356,7 @@ static int mmq_get_nwarps_host(const int /*cc*/, const int warp_size) { static constexpr __device__ int mmq_get_nwarps_device() { #if LUCEBOX_RDNA_TILE_DEVICE -#if defined(RDNA3) +#if defined(RDNA3_5) return 4; #elif defined(GGML_CUDA_ROCMFPX_MMQ_TILE) return 4; diff --git a/server/src/deepseek4/deepseek4_backend.cpp b/server/src/deepseek4/deepseek4_backend.cpp index 4834cd2d2..7fb49bcc8 100644 --- a/server/src/deepseek4/deepseek4_backend.cpp +++ b/server/src/deepseek4/deepseek4_backend.cpp @@ -1813,7 +1813,8 @@ int DeepSeek4Backend::do_prefill(const std::vector & tokens, DeepSeek4StepTelemetry step_tel; if (timing) step_tel.embed_us = elapsed_us(embed_t0, Clock::now()); - const bool need_logits = (i + n_tok >= n_total); + const bool at_snap_boundary = save_snapshot && !snapshot_saved && (pos + n_tok >= snap_pos); + const bool need_logits = (i + n_tok >= n_total) || at_snap_boundary; std::vector logits; bool ok = false; std::vector hc_state; diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index 4ab3418d7..7ca84ae26 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -6124,7 +6124,7 @@ static int ds4_try_layer_major_prefill( const float * embed, int n_tokens, int kv_start, - std::vector & out_logits, + std::vector * out_logits, const int32_t * token_ids, Ds4VerifyHooks * verify_hooks, DeepSeek4StepTelemetry * telemetry) { @@ -6403,10 +6403,10 @@ static int ds4_try_layer_major_prefill( compute_t0, Ds4TimingClock::now()); } capture_layer(il, (il & 1) == 0 ? state_b : state_a); - if (layer.logits) { - out_logits.resize((size_t) w.n_vocab); + if (layer.logits && out_logits) { + out_logits->resize((size_t) w.n_vocab); ggml_backend_tensor_get( - layer.logits, out_logits.data(), 0, + layer.logits, out_logits->data(), 0, sizeof(float) * (size_t) w.n_vocab); } @@ -6421,7 +6421,7 @@ static int ds4_try_layer_major_prefill( } } cache.cur_pos = next_pos; - return out_logits.empty() ? -1 : 1; + return (out_logits && out_logits->empty()) ? -1 : 1; } ggml_tensor * state_in = state_a; @@ -6540,7 +6540,7 @@ static int ds4_try_layer_major_prefill( ggml_build_forward_expand(gf, state_copy); ggml_tensor * logits = nullptr; - if (il + 1 == w.n_layer) { + if (out_logits && il + 1 == w.n_layer) { ggml_tensor * last_hc = ggml_view_2d( ctx, hc_next, hc_dim, 1, hc_next->nb[1], (size_t) (n_tokens - 1) * hc_next->nb[1]); @@ -6631,9 +6631,9 @@ static int ds4_try_layer_major_prefill( capture_layer(il, state_out); - if (logits) { - out_logits.resize((size_t) w.n_vocab); - ggml_backend_tensor_get(logits, out_logits.data(), 0, + if (logits && out_logits) { + out_logits->resize((size_t) w.n_vocab); + ggml_backend_tensor_get(logits, out_logits->data(), 0, sizeof(float) * (size_t) w.n_vocab); } @@ -6665,7 +6665,7 @@ static int ds4_try_layer_major_prefill( ggml_free(state_ctx); } cache.cur_pos = next_pos; - return out_logits.empty() ? -1 : 1; + return (out_logits && out_logits->empty()) ? -1 : 1; } static bool ds4_hc_layer_weights_ready(const HcWeightsCpu & weights, @@ -6811,7 +6811,7 @@ bool deepseek4_step_layer_range( !fused_verify_candidate && moe_hybrid && cache.prefill_mode == PrefillAttentionMode::Sparse && n_tokens > 4 && n_tokens <= DS4_MAX_LAYER_MAJOR_PREFILL_TOKENS && - layer_begin == 0 && is_last_shard && out_logits && + layer_begin == 0 && is_last_shard && ds4_backend_is_gpu(backend); const bool layer_major_hooks_supported = !verify_hooks || @@ -6824,7 +6824,7 @@ bool deepseek4_step_layer_range( const bool standard_layer_major_prefill = !w.moe_hybrid && cache.prefill_mode != PrefillAttentionMode::Exact && n_tokens > 4 && n_tokens <= DS4_MAX_LAYER_MAJOR_PREFILL_TOKENS && - layer_begin == 0 && is_last_shard && out_logits && + layer_begin == 0 && is_last_shard && ds4_backend_is_gpu(backend) && layer_major_hooks_supported; // These graphs are rebuilt around an owner join on every layer, so tensor // metadata addresses can be recycled for different topologies. Until @@ -7018,7 +7018,7 @@ bool deepseek4_step_layer_range( fused_decode_graph_cache, backend, w, cache, hc_layer_weights_range, hc_output_weights_range, hash_routing_tables_range, scratch.hash_expert_ids, embed, - n_tokens, kv_start, *out_logits, token_ids, verify_hooks, + n_tokens, kv_start, out_logits, token_ids, verify_hooks, telemetry); if (prc < 0) return false; if (prc > 0) { From 92366b47cc3a588064a30d3781685ff204d17723 Mon Sep 17 00:00:00 2001 From: cheese-cakee Date: Wed, 19 Aug 2026 17:05:45 +0530 Subject: [PATCH 3/4] fix(ds4): differentiate layer-major graph cache by logits presence and preserve cuda mmq fallback --- server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh | 15 +++++++++++++++ server/src/deepseek4/deepseek4_graph.cpp | 13 ++++++++++--- 2 files changed, 25 insertions(+), 3 deletions(-) diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh index a9f9e836f..73ad8afdf 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/mmq.cuh @@ -153,7 +153,22 @@ static constexpr __device__ int get_mmq_x_max_device() { #if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) return 128; #else // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) + +#if defined(GGML_USE_HIP) return 64; +#else // defined(GGML_USE_HIP) + +#if __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA +#ifdef GGML_CUDA_FORCE_MMQ + return 128; +#else // GGML_CUDA_FORCE_MMQ + return MMQ_DP4A_MAX_BATCH_SIZE; +#endif // GGML_CUDA_FORCE_MMQ +#else // __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA + return 64; +#endif // __CUDA_ARCH__ >= GGML_CUDA_CC_VOLTA + +#endif // defined(GGML_USE_HIP) #endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) } diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index 7ca84ae26..9f0950487 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -6013,6 +6013,7 @@ struct Ds4LayerMajorGraphCache { PrefillAttentionMode mode = PrefillAttentionMode::Exact; int n_tokens = 0; int kv_start = -1; + bool has_logits = false; bool ready = false; ggml_context * state_ctx = nullptr; ggml_backend_buffer_t state_buf = nullptr; @@ -6021,9 +6022,11 @@ struct Ds4LayerMajorGraphCache { std::vector layers; bool matches(const DeepSeek4Weights & w, ggml_backend_t b, - PrefillAttentionMode m, int tokens, int start) const { + PrefillAttentionMode m, int tokens, int start, + bool logits_needed) const { return ready && owner_ctx == w.ctx && backend == b && mode == m && n_tokens == tokens && kv_start == start && + has_logits == logits_needed && layers.size() == (size_t) w.n_layer; } @@ -6045,6 +6048,7 @@ struct Ds4LayerMajorGraphCache { mode = PrefillAttentionMode::Exact; n_tokens = 0; kv_start = -1; + has_logits = false; ready = false; } }; @@ -6236,10 +6240,11 @@ static int ds4_try_layer_major_prefill( Ds4LayerMajorGraphCache * graph_cache = nullptr; bool cache_hit = false; bool cache_build = false; + const bool logits_needed = (out_logits != nullptr); if (token_ids) { for (auto & candidate : ds4_layer_major_graph_caches) { if (candidate.matches(w, backend, cache.prefill_mode, - n_tokens, kv_start)) { + n_tokens, kv_start, logits_needed)) { graph_cache = &candidate; cache_hit = true; break; @@ -6254,7 +6259,8 @@ static int ds4_try_layer_major_prefill( candidate.backend == backend && candidate.mode == cache.prefill_mode; if (!candidate.ready || !same_owner || - n_tokens > candidate.n_tokens) { + n_tokens > candidate.n_tokens || + (n_tokens == candidate.n_tokens && candidate.has_logits != logits_needed)) { graph_cache = &candidate; graph_cache->destroy(); graph_cache->owner_ctx = w.ctx; @@ -6262,6 +6268,7 @@ static int ds4_try_layer_major_prefill( graph_cache->mode = cache.prefill_mode; graph_cache->n_tokens = n_tokens; graph_cache->kv_start = kv_start; + graph_cache->has_logits = logits_needed; graph_cache->layers.resize((size_t) w.n_layer); cache_build = true; } From a2878db6ec0d03a6252a6f8222169425e8c62298 Mon Sep 17 00:00:00 2001 From: cheese-cakee Date: Wed, 19 Aug 2026 17:29:57 +0530 Subject: [PATCH 4/4] fix(ds4): preserve dominant no-logits graph cache across transient terminal steps --- server/src/deepseek4/deepseek4_graph.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index 9f0950487..c427fe724 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -6255,12 +6255,14 @@ static int ds4_try_layer_major_prefill( // Do not evict a full/larger chunk for an equal-size graph at a // later position or for a short tail. Both execute with the shared // scratch arena below, but only the dominant topology stays cached. + // When logits are needed on a tail/terminal step, execute it + // transiently rather than evicting the dominant no-logits graph. const bool same_owner = candidate.owner_ctx == w.ctx && candidate.backend == backend && candidate.mode == cache.prefill_mode; - if (!candidate.ready || !same_owner || - n_tokens > candidate.n_tokens || - (n_tokens == candidate.n_tokens && candidate.has_logits != logits_needed)) { + const bool can_cache_dominant = !logits_needed; + if (can_cache_dominant && (!candidate.ready || !same_owner || + n_tokens > candidate.n_tokens)) { graph_cache = &candidate; graph_cache->destroy(); graph_cache->owner_ctx = w.ctx; @@ -6268,7 +6270,7 @@ static int ds4_try_layer_major_prefill( graph_cache->mode = cache.prefill_mode; graph_cache->n_tokens = n_tokens; graph_cache->kv_start = kv_start; - graph_cache->has_logits = logits_needed; + graph_cache->has_logits = false; graph_cache->layers.resize((size_t) w.n_layer); cache_build = true; }