fix(server): each turn reuses the previous one, long prompts share time, late system messages work

server/deps/llama.cpp/ggml/src/ggml-cuda/paged-attn.cu

@@ -3,6 +3,7 @@
 #include "fattn-common.cuh"
 #include "mma.cuh"
 
+#include <algorithm>
 #include <atomic>
 #include <cfloat>
 #include <cstdlib>
@@ -18,6 +19,20 @@ static __host__ __device__ __forceinline__ int32_t paged_attn_ceil_div(
     return (int32_t) (dividend / divisor + (dividend % divisor != 0));
 }
 
+// Most partitions a launch may split the context into: each one keeps a
+// head-dim row of halves and a float2 per output row, and the context alone
+// asks for one per 1024 tokens. A 2048-row prompt chunk 118K tokens deep ran
+// the R9700 out of memory on those partials, beside the weights and the KV
+// cache. Its rows give such a chunk its parallelism, so it covers the
+// context with fewer, longer partitions instead.
+static int32_t paged_attn_partials_cap(int64_t output_rows, int32_t head_dim) {
+    constexpr int64_t budget = 256ll << 20;
+    const int64_t per_partition =
+        output_rows * ((int64_t) head_dim * (int64_t) sizeof(half) + (int64_t) sizeof(float2));
+    const int64_t cap = per_partition > 0 ? budget / per_partition : PAGED_ATTN_MAX_PARTITIONS;
+    return (int32_t) std::max<int64_t>(1, std::min<int64_t>(cap, PAGED_ATTN_MAX_PARTITIONS));
+}
+
 // Partition count for n_blocks of context: enough partitions to cover the
 // blocks and to reach min_partitions for occupancy, but never more than the
 // blocks themselves or the cap. Host and device must agree on this so the
@@ -978,6 +993,10 @@ static bool try_launch_paged_attn(
     if (min_partitions > partition_limit) {
         min_partitions = partition_limit;
     }
+    const int32_t partition_cap = paged_attn_partials_cap(output_rows, D);
+    if (min_partitions > partition_cap) {
+        min_partitions = partition_cap;
+    }
     const int32_t tree_blocks = paged_attn_ceil_div(tree_width, block_size);
     const int64_t partitionable_blocks =
         block_table->ne[0] + (parent_ids ? tree_blocks : 0);
@@ -993,7 +1012,7 @@ static bool try_launch_paged_attn(
     const int32_t live_blocks =
         paged_attn_ceil_div(live_tokens, block_size);
     int32_t n_partitions = paged_attn_partitions(
-        live_blocks, min_partitions, PAGED_ATTN_MAX_PARTITIONS);
+        live_blocks, min_partitions, partition_cap);
 
     // Test/debug override used to exercise both the direct and partials paths
     // independently of device-specific occupancy. Values outside the valid
@@ -2162,11 +2181,15 @@ static bool try_launch_paged_attn_wmma(ggml_backend_cuda_context & ctx, ggml_ten
     if (min_partitions > partition_limit) {
         min_partitions = partition_limit;
     }
+    const int32_t partition_cap = paged_attn_partials_cap((int64_t) n_rows * n_head, 256);
+    if (min_partitions > partition_cap) {
+        min_partitions = partition_cap;
+    }
     if (min_partitions > partitionable_blocks) {
         min_partitions = partitionable_blocks;
     }
     int32_t n_partitions =
-        paged_attn_partitions(n_logical_blocks, min_partitions, PAGED_ATTN_MAX_PARTITIONS);
+        paged_attn_partitions(n_logical_blocks, min_partitions, partition_cap);
     // Test/debug override, applied after the occupancy floor like the decode
     // launcher: the floor above would otherwise silently raise any forced
     // value <= 32 back to 32 and defeat partition bisection experiments.

server/src/common/concurrency/seq_engine.h

@@ -78,6 +78,13 @@ struct StepPlanLimits {
     int max_prefill_tokens_per_sequence = 512;
     int max_prefill_tokens_total = 512;
     int prefill_allocation_quantum = 512;
+    // While requests decode, the share of wall time a long prompt may take:
+    // after each of its chunks, the decoding requests step alone for the
+    // rest. 1: a prompt chunk rides every step.
+    double prefill_time_share = 1.0;
+    // A prompt is long when it brings more new tokens than this; 0: more
+    // than one step's chunk (max_prefill_tokens_per_sequence).
+    int time_share_prompt_tokens = 0;
 };
 
 // Select the oldest eligible sequences, then distribute the step's token

server/src/common/model_backend.h

@@ -269,9 +269,9 @@ struct ModelBackend {
     virtual int restore_point_spacing() const { return 0; }
 
     // True when the backend ends a prompt forward at every restore point
-    // (GenerateRequest::restore_points). A tool-result request then keeps its
-    // own last boundary among them, so a hit that skips the capture and a
-    // miss that captures the tools head cut prefill alike.
+    // (GenerateRequest::restore_points). A request then keeps its own last
+    // boundary among them, where its snapshot goes, so a hit that skips the
+    // capture and a miss that captures cut prefill alike.
     virtual bool prefill_cuts_at_restore_points() const { return false; }
 
     // System-memory bytes a snapshot of the first `tokens` positions would

server/src/qwen35/concurrency/qwen35_seq_engine.h

@@ -88,6 +88,10 @@ class Qwen35SeqEngine final : public SeqEngine {
         const PrefixStorePlan & plan) override;
 
     StepResult step(const StepPlan & plan) override;
+    // A step that carries a prompt chunk decodes without the drafter, one
+    // token per request. Beside a prompt of more than 8192 new tokens (a
+    // document, not an agent's next step), the decoding requests get half
+    // the time back in steps of their own, with it.
     StepPlanLimits step_plan_limits(int decode_rows) const override {
         const bool mixed = decode_rows > 0;
         const int per_sequence = mixed ? 512 : 2048;
@@ -103,6 +107,8 @@ class Qwen35SeqEngine final : public SeqEngine {
             per_sequence,
             std::min(max_prefills_ * per_sequence, total_cap),
             prefill_quantum_,
+            /*prefill_time_share=*/mixed ? 0.5 : 1.0,
+            /*time_share_prompt_tokens=*/8192,
         };
     }
 

server/src/qwen4exp/qwen4exp_seq_engine.cpp

@@ -355,12 +355,17 @@ int Qwen4ExpSeqEngine::prefill_segment(int slot, int max_tokens) const {
 }
 
 // One granule per prompt forward, whatever the step carries. While slots
-// decode, one granule in total per step bounds the hold on the live streams.
+// decode, one granule in total per step bounds the hold on the live streams,
+// and beside a prompt of more than 8192 new tokens (a document, not an agent's
+// next step) they then decode alone for as long as the granule took: a granule
+// pays the whole expert read, so a shorter one would not let them run faster,
+// and a long prompt would otherwise give them one token per granule.
 StepPlanLimits Qwen4ExpSeqEngine::step_plan_limits(int decode_rows) const {
     const int prefill_slots = slot_count() - std::clamp(decode_rows, 0, slot_count());
     if (decode_rows > 0) {
         const int sequences = std::min(prefill_slots, 1);
-        return {sequences, prefill_granule_, sequences * prefill_granule_, prefill_granule_};
+        return {sequences, prefill_granule_, sequences * prefill_granule_, prefill_granule_,
+                /*prefill_time_share=*/0.5, /*time_share_prompt_tokens=*/8192};
     }
     return {prefill_slots, prefill_granule_, prefill_slots * prefill_granule_, prefill_granule_};
 }

server/src/server/chat_template.cpp

@@ -1046,7 +1046,7 @@ static std::shared_ptr<jinja::program> get_or_parse(const std::string & template
 
 }  // namespace
 
-std::string render_chat_template_jinja(
+static std::string render_jinja(
     const std::string & template_src,
     const std::vector<ChatMessage> & messages,
     const std::string & bos_token,
@@ -1055,7 +1055,8 @@ std::string render_chat_template_jinja(
     bool enable_thinking,
     const std::string & tools_json,
     const std::string & reasoning_effort,
-    int preserve_thinking)
+    int preserve_thinking,
+    bool late_system_as_user)
 {
     if (template_src.empty()) {
         throw std::runtime_error("render_chat_template_jinja: template_src is empty");
@@ -1077,7 +1078,9 @@ std::string render_chat_template_jinja(
                 "normalize with tool-memory replay disabled");
         }
         nlohmann::ordered_json mj;
-        mj["role"]    = m.role;
+        const bool late_system = late_system_as_user && !messages_j.empty() &&
+            (m.role == "system" || m.role == "developer");
+        mj["role"]    = late_system ? std::string("user") : m.role;
         mj["content"] = m.content;
         if (!m.tool_call_id.empty()) {
             mj["tool_call_id"] = m.tool_call_id;
@@ -1164,4 +1167,33 @@ std::string render_chat_template_jinja(
     }
 }
 
+std::string render_chat_template_jinja(
+    const std::string & template_src,
+    const std::vector<ChatMessage> & messages,
+    const std::string & bos_token,
+    const std::string & eos_token,
+    bool add_generation_prompt,
+    bool enable_thinking,
+    const std::string & tools_json,
+    const std::string & reasoning_effort,
+    int preserve_thinking)
+{
+    try {
+        return render_jinja(template_src, messages, bos_token, eos_token, add_generation_prompt,
+                            enable_thinking, tools_json, reasoning_effort, preserve_thinking, false);
+    } catch (const std::runtime_error &) {
+        // Some templates take a system message only first (Qwen's raise).
+        // Claude Code sends its environment as a system message after the
+        // user's first one; for such a template it reads as a user turn
+        // there, in place, so the prompt before it stays the same and cached.
+        const bool late_system = messages.size() > 1 &&
+            std::any_of(messages.begin() + 1, messages.end(), [](const ChatMessage & m) {
+                return m.role == "system" || m.role == "developer";
+            });
+        if (!late_system) throw;
+        return render_jinja(template_src, messages, bos_token, eos_token, add_generation_prompt,
+                            enable_thinking, tools_json, reasoning_effort, preserve_thinking, true);
+    }
+}
+
 }  // namespace luce::common

server/src/server/http_server.cpp

@@ -1470,9 +1470,9 @@ HttpServer::HttpServer(luce::engine::LuceEngine & engine,
     curl_global_init(CURL_GLOBAL_DEFAULT);
     #endif
     prefix_cache_.init_full_cache(config.prefill_cache_cap);
-    // Single-sequence commits prune superseded snapshots (see
-    // trim_snapshots_after_commit); the paged scheduler does not.
-    prefix_cache_.set_prunes_superseded(!config.concurrent_prefix_cache);
+    // Commits prune the snapshots they supersede: single-sequence serving in
+    // trim_snapshots_after_commit, the concurrent scheduler as each capture lands.
+    prefix_cache_.set_prunes_superseded(true);
     // Fold model+config identity into the layout fingerprint BEFORE init()
     // so compute_layout_id sees it on every learn/verify call. Prevents stale
     // KV hits when the server restarts over the same --kv-cache-dir with a
@@ -2415,7 +2415,9 @@ void apply_request_reasoning(
     }
     if (body.contains("thinking") && body["thinking"].is_object()) {
         const auto & thinking = body["thinking"];
-        if (thinking.contains("type")) {
+        // "adaptive" (Claude Code's default) leaves thinking to the model: the
+        // same as no thinking field.
+        if (thinking.contains("type") && thinking.value("type", "") != "adaptive") {
             const bool enabled = thinking.value("type", "") == "enabled";
             enable_thinking = enabled;
             req.thinking_opt_in = enabled;
@@ -2560,7 +2562,9 @@ bool HttpServer::render_and_tokenize_request(
     if (!render_messages_to_text(chat_messages, req,
                                  /*add_generation_prompt=*/true,
                                  req.rendered_prompt, error)) {
-        send_error(fd, 500, error);
+        // The template rejected these messages: a client error, which a
+        // client must not retry as it would a server failure.
+        send_error(fd, 400, error);
         return false;
     }
     req.started_in_thinking = prompt_ends_in_open_think(req.rendered_prompt);
@@ -2973,6 +2977,9 @@ json build_openai_completion_response(
         {"completion_tokens_details", {
             {"reasoning_tokens", counts.reasoning},
         }},
+        {"prompt_tokens_details", {
+            {"cached_tokens", timings.cached_prefix_tokens},
+        }},
         {"timings", build_timings_json(timings, counts.total)},
         {"accept_rate", result.accept_rate},
         {"spec_decode_ran", result.spec_decode_ran},
@@ -3073,6 +3080,9 @@ json build_responses_api_response(
         {"input_tokens", prompt_tokens},
         {"output_tokens", counts.total},
         {"total_tokens", prompt_tokens + counts.total},
+        {"input_tokens_details", {
+            {"cached_tokens", timings.cached_prefix_tokens},
+        }},
         {"timings", build_timings_json(timings, counts.total)},
         {"accept_rate", result.accept_rate},
         {"spec_decode_ran", result.spec_decode_ran},
@@ -3667,10 +3677,9 @@ bool HttpServer::forward_upstream(
 }
 
 std::vector<int> HttpServer::request_restore_points(const std::vector<int32_t> & prompt,
-                                                    bool ends_with_tool_result,
                                                     std::initializer_list<int> cuts) const {
     return prefix_restore_points(prompt, prefix_cache_.chat_markers(), cuts,
-        /*drop_last_boundary=*/!(backend_.prefill_cuts_at_restore_points() && ends_with_tool_result),
+        /*drop_last_boundary=*/!backend_.prefill_cuts_at_restore_points(),
         backend_.restore_point_spacing());
 }
 
@@ -4033,8 +4042,7 @@ HttpServer::GenerationCacheState HttpServer::prepare_generation_cache(
             [this](int target_cut) {
                 return backend_.snapshot_bytes_estimate(target_cut);
             },
-            req.ends_with_tool_result, reachable_from,
-            backend_.restore_point_spacing());
+            reachable_from, backend_.restore_point_spacing());
         cache.snap_slot = cache.snap_reservation.slot();
         cache.snap_cut = cache.snap_reservation.target_cut();
     };
@@ -4114,7 +4122,7 @@ HttpServer::GenerationCacheState HttpServer::prepare_generation_cache(
 
     if (!prefix_cache_.disabled() || !disk_cache_.disabled()) {
         generate_request.restore_points = request_restore_points(
-            effective_prompt, req.ends_with_tool_result,
+            effective_prompt,
             {forced_cut, selected_boundary,
              cache.snap_prepared ? cache.snap_cut : 0});
     }

server/src/server/http_server.h

@@ -378,8 +378,8 @@ struct ParsedRequest {
     DiskPrefixCachePolicy     disk_cache_policy;
     // PPP: stable pin cut for tool-heavy requests (0 = use default boundary).
     int                       pin_end_token = 0;
-    // The prompt ends with tool results. Clients only append after them, so
-    // the inline snapshot may cover them (see select_inline_snapshot_boundary).
+    // The prompt ends with tool results: an agent's step, which clients only
+    // append to.
     bool                      ends_with_tool_result = false;
 };
 
@@ -542,10 +542,9 @@ class HttpServer {
 
     // A request's restore points (prefix_restore_points) on both serving
     // paths: its chat boundaries spaced for the backend, plus `cuts`. The
-    // generation prompt's own boundary stays only for a tool-result request
-    // on a backend that cuts prefill at every restore point.
+    // generation prompt's own boundary, where the request's snapshot goes,
+    // stays only on a backend that cuts prefill at every restore point.
     std::vector<int> request_restore_points(const std::vector<int32_t> & prompt,
-                                            bool ends_with_tool_result,
                                             std::initializer_list<int> cuts) const;
     GenerationCacheState prepare_generation_cache(
         const ParsedRequest & req, PreparedPrompt & prepared,