# Applies to llama.cpp commit c6824a9e
# Includes eaman commit ddddf03d (cumulative Eaman ROCm/MTP patch merged onto c6824a9e mainline)
# Tested with llama.cpp version 1676
# Tested on ROCm core-10.0 / HIP 7.15.26333 / Vulkan
diff --git a/common/arg.cpp b/common/arg.cpp
index c4c4e143..9f24ad28 100644
--- a/common/arg.cpp
+++ b/common/arg.cpp
@@ -1291,6 +1291,9 @@ bool common_params_parse(int argc, char ** argv, common_params & params, llama_e
             ctx_arg.params = params_org;
             return false;
         }
+        if (ctx_arg.params.speculative.draft.n_rs_seq > ctx_arg.params.speculative.draft.n_max) {
+            throw std::invalid_argument("--spec-mtp-cr-depth must not exceed --spec-draft-n-max");
+        }
         if (ctx_arg.params.usage) {
             common_params_print_usage(ctx_arg);
             if (ctx_arg.print_usage) {
@@ -1763,6 +1766,34 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
                                    string_format("error: unknown value for --flash-attn: '%s'\n", value.c_str()));
                            }
                        }).set_env("LLAMA_ARG_FLASH_ATTN"));
+    add_opt(common_arg({ "--hip-fa-force-vec" }, "[on|off]",
+                       string_format("force the HIP quantized-KV Flash Attention VEC path when supported (default: '%s')",
+                                     params.hip_fa_force_vec ? "on" : "off"),
+                       [](common_params & params, const std::string & value) {
+                           if (is_truthy(value)) {
+                               params.hip_fa_force_vec = true;
+                           } else if (is_falsey(value)) {
+                               params.hip_fa_force_vec = false;
+                           } else {
+                               throw std::runtime_error(
+                                   string_format("error: unknown value for --hip-fa-force-vec: '%s'\n", value.c_str()));
+                           }
+                       }).set_env("LLAMA_ARG_HIP_FA_FORCE_VEC"));
+    add_opt(common_arg({ "--pipeline-parallel" }, "[on|off|auto]",
+                       string_format("set pipeline parallelism ('on', 'off', or 'auto', default: '%s')",
+                                     llama_pipeline_parallel_type_name(params.pipeline_parallel_type)),
+                       [](common_params & params, const std::string & value) {
+                           if (is_truthy(value)) {
+                               params.pipeline_parallel_type = LLAMA_PIPELINE_PARALLEL_TYPE_ENABLED;
+                           } else if (is_falsey(value)) {
+                               params.pipeline_parallel_type = LLAMA_PIPELINE_PARALLEL_TYPE_DISABLED;
+                           } else if (is_autoy(value)) {
+                               params.pipeline_parallel_type = LLAMA_PIPELINE_PARALLEL_TYPE_AUTO;
+                           } else {
+                               throw std::runtime_error(
+                                   string_format("error: unknown value for --pipeline-parallel: '%s'\n", value.c_str()));
+                           }
+                       }).set_env("LLAMA_ARG_PIPELINE_PARALLEL"));
     add_opt(common_arg(
         {"-p", "--prompt"}, "PROMPT",
         "prompt to start generation with; for system message, use -sys",
@@ -4133,6 +4164,23 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
             params.speculative.draft.n_max = value;
         }
     ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_LOOKUP, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_DRAFT_N_MAX"));
+    add_opt(common_arg(
+        {"--spec-mtp-cr-depth"}, "N",
+        "MTP Compact Rollback depth; lower values save memory but replay accepted tokens after deep rejection (default: --spec-draft-n-max)",
+        [](common_params & params, int value) {
+            if (value < 1) {
+                throw std::invalid_argument("--spec-mtp-cr-depth must be at least 1");
+            }
+            params.speculative.draft.n_rs_seq = value;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_MTP_CR_DEPTH"));
+    add_opt(common_arg(
+        {"--spec-draft-adaptive"},
+        string_format("size each draft from measured acceptance rather than always drafting --spec-draft-n-max (default: %s)", params.speculative.draft.adaptive ? "true" : "false"),
+        [](common_params & params) {
+            params.speculative.draft.adaptive = true;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_LOOKUP, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_DRAFT_ADAPTIVE"));
     add_opt(common_arg(
         {"--spec-draft-n-min"}, "N",
         string_format("minimum number of draft tokens to use for speculative decoding (default: %d)", params.speculative.draft.n_min),
diff --git a/common/common.cpp b/common/common.cpp
index d8319cd9..4ba9dad9 100644
--- a/common/common.cpp
+++ b/common/common.cpp
@@ -1307,6 +1307,7 @@ common_init_result::common_init_result(common_params & params, bool model_only)
         auto cparams_dft = common_context_params_to_llama(params_dft);
         if (spec_mtp) {
             cparams_dft.ctx_type = LLAMA_CONTEXT_TYPE_MTP;
+            cparams_dft.pipeline_parallel_type = LLAMA_PIPELINE_PARALLEL_TYPE_DISABLED;
         }
         cparams_dft.n_rs_seq = 0;
 
@@ -1740,6 +1741,8 @@ struct llama_context_params common_context_params_to_llama(const common_params &
     cparams.pooling_type      = params.pooling_type;
     cparams.attention_type    = params.attention_type;
     cparams.flash_attn_type   = params.flash_attn_type;
+    cparams.hip_fa_force_vec  = params.hip_fa_force_vec;
+    cparams.pipeline_parallel_type = params.pipeline_parallel_type;
     cparams.cb_eval           = params.cb_eval;
     cparams.cb_eval_user_data = params.cb_eval_user_data;
     cparams.offload_kqv       = !params.no_kv_offload;
@@ -2268,6 +2271,8 @@ void common_prompt_checkpoint::clear() {
     data_tgt.clear();
     data_dft.clear();
     data_spec.clear();
+    data_tgt_on_device = false;
+    data_dft_on_device = false;
 }
 
 void common_prompt_checkpoint::update_pos(
@@ -2279,6 +2284,50 @@ void common_prompt_checkpoint::update_pos(
     this->pos_max  = pos_max;
 }
 
+static void common_prompt_checkpoint_save(
+        std::vector<uint8_t> & data,
+        bool & on_device,
+        llama_context * ctx,
+        llama_seq_id seq_id,
+        llama_state_seq_flags flags,
+        const char * label) {
+    const bool requested_on_device = flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+    on_device = false;
+
+    auto save = [&](llama_state_seq_flags save_flags) {
+        const size_t ckpt_size = llama_state_seq_get_size_ext(ctx, seq_id, save_flags);
+        if (ckpt_size == 0) {
+            return false;
+        }
+
+        std::vector<uint8_t> saved(ckpt_size);
+        const size_t n = llama_state_seq_get_data_ext(ctx, saved.data(), ckpt_size, seq_id, save_flags);
+        if (n != ckpt_size) {
+            return false;
+        }
+
+        data.swap(saved);
+        return true;
+    };
+
+    bool saved = save(flags);
+    if (requested_on_device && !saved) {
+        COM_WRN("%s: ON_DEVICE %s checkpoint save failed; retrying with host storage\n",
+                __func__, label);
+        flags &= ~LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+        saved = save(flags);
+    }
+
+    if (!saved) {
+        if (flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) {
+            GGML_ABORT("checkpoint device save failed for %s\n", label);
+        }
+        GGML_ABORT("checkpoint size mismatch while saving %s\n", label);
+    }
+
+    on_device = flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+}
+
 void common_prompt_checkpoint::update_tgt(
         llama_context * ctx,
         llama_seq_id seq_id,
@@ -2287,14 +2336,7 @@ void common_prompt_checkpoint::update_tgt(
         return;
     }
 
-    const size_t ckpt_size = llama_state_seq_get_size_ext(ctx, seq_id, flags);
-
-    data_tgt.resize(ckpt_size);
-
-    const size_t n = llama_state_seq_get_data_ext(ctx, data_tgt.data(), ckpt_size, seq_id, flags);
-    if (n != ckpt_size) {
-        GGML_ABORT("checkpoint size mismatch: expected %zu, got %zu\n", ckpt_size, n);
-    }
+    common_prompt_checkpoint_save(data_tgt, data_tgt_on_device, ctx, seq_id, flags, "target");
 }
 
 void common_prompt_checkpoint::update_dft(
@@ -2305,14 +2347,7 @@ void common_prompt_checkpoint::update_dft(
         return;
     }
 
-    const size_t ckpt_size = llama_state_seq_get_size_ext(ctx, seq_id, flags);
-
-    data_dft.resize(ckpt_size);
-
-    const size_t n = llama_state_seq_get_data_ext(ctx, data_dft.data(), ckpt_size, seq_id, flags);
-    if (n != ckpt_size) {
-        GGML_ABORT("checkpoint size mismatch: expected %zu, got %zu\n", ckpt_size, n);
-    }
+    common_prompt_checkpoint_save(data_dft, data_dft_on_device, ctx, seq_id, flags, "draft");
 }
 
 void common_prompt_checkpoint::load_tgt(
@@ -2327,6 +2362,9 @@ void common_prompt_checkpoint::load_tgt(
         return;
     }
 
+    flags = (flags & ~LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) |
+            (data_tgt_on_device ? LLAMA_STATE_SEQ_FLAGS_ON_DEVICE : 0);
+
     const size_t n = llama_state_seq_set_data_ext(ctx, data_tgt.data(), data_tgt.size(), seq_id, flags);
     if (n != data_tgt.size()) {
         GGML_ABORT("checkpoint size mismatch: expected %zu, got %zu\n", data_tgt.size(), n);
@@ -2345,6 +2383,9 @@ void common_prompt_checkpoint::load_dft(
         return;
     }
 
+    flags = (flags & ~LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) |
+            (data_dft_on_device ? LLAMA_STATE_SEQ_FLAGS_ON_DEVICE : 0);
+
     const size_t n = llama_state_seq_set_data_ext(ctx, data_dft.data(), data_dft.size(), seq_id, flags);
     if (n != data_dft.size()) {
         GGML_ABORT("checkpoint size mismatch: expected %zu, got %zu\n", data_dft.size(), n);
@@ -2353,9 +2394,11 @@ void common_prompt_checkpoint::load_dft(
 
 void common_prompt_checkpoint::clear_tgt() {
     data_tgt.clear();
+    data_tgt_on_device = false;
 }
 
 void common_prompt_checkpoint::clear_dft() {
     data_dft.clear();
     data_spec.clear();
+    data_dft_on_device = false;
 }
diff --git a/common/common.h b/common/common.h
index 63d0badd..da8a243c 100644
--- a/common/common.h
+++ b/common/common.h
@@ -325,6 +325,7 @@ struct common_params_model {
 struct common_params_speculative_draft {
     int32_t n_max = 3; // maximum number of tokens to draft during speculative decoding
     int32_t n_min = 0; // minimum number of draft tokens to use for speculative decoding
+    int32_t n_rs_seq = -1; // MTP Compact Rollback depth (-1 = n_max)
 
     float p_split = 0.1f; // speculative decoding split probability
     float p_min   = 0.0f; // minimum speculative decoding probability (greedy)
@@ -341,6 +342,9 @@ struct common_params_speculative_draft {
     ggml_type cache_type_k = GGML_TYPE_F16; // KV cache data type for the K
     ggml_type cache_type_v = GGML_TYPE_F16; // KV cache data type for the V
 
+    // size each draft from measured acceptance instead of always drafting n_max
+    bool adaptive = false;
+
     common_cpu_params cpuparams;
     common_cpu_params cpuparams_batch;
 
@@ -392,11 +396,18 @@ struct common_params_speculative {
     }
 
     uint32_t need_n_rs_seq() const {
-        bool needs_rs_seq = std::any_of(types.begin(), types.end(), [&](auto t) {
-            return t == COMMON_SPECULATIVE_TYPE_DRAFT_MTP || t == COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3 || t == COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH || t == COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK;
+        const bool needs_mtp = std::find(types.begin(), types.end(), COMMON_SPECULATIVE_TYPE_DRAFT_MTP) != types.end();
+        const bool needs_other_rs = std::any_of(types.begin(), types.end(), [&](auto t) {
+            return t == COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3 || t == COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH || t == COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK;
         });
 
-        return needs_rs_seq ? draft.n_max : 0u;
+        if (needs_other_rs) {
+            return draft.n_max;
+        }
+        if (needs_mtp) {
+            return draft.n_rs_seq >= 0 ? draft.n_rs_seq : draft.n_max;
+        }
+        return 0u;
     }
 };
 
@@ -497,6 +508,8 @@ struct common_params {
     enum llama_pooling_type      pooling_type      = LLAMA_POOLING_TYPE_UNSPECIFIED; // pooling type for embeddings
     enum llama_attention_type    attention_type    = LLAMA_ATTENTION_TYPE_UNSPECIFIED; // attention type for embeddings
     enum llama_flash_attn_type   flash_attn_type   = LLAMA_FLASH_ATTN_TYPE_AUTO; // whether to use Flash Attention
+    bool                         hip_fa_force_vec  = false; // force HIP quantized-KV FA onto VEC when supported
+    enum llama_pipeline_parallel_type pipeline_parallel_type = LLAMA_PIPELINE_PARALLEL_TYPE_AUTO; // whether to use pipeline parallelism
 
     struct common_params_sampling    sampling;
     struct common_params_speculative speculative;
@@ -1174,6 +1187,11 @@ struct common_prompt_checkpoint {
     std::vector<uint8_t> data_tgt;
     std::vector<uint8_t> data_dft;
 
+    // The actual storage mode used for each checkpoint.  This can differ from
+    // the flags passed to update_* when an on-device save is not possible.
+    bool data_tgt_on_device = false;
+    bool data_dft_on_device = false;
+
     // (optional) speculative-decoding implementation state stashed with the checkpoint
     // (e.g. eagle3's deferred-boundary g_embd row)
     std::vector<uint8_t> data_spec;
diff --git a/common/fit.cpp b/common/fit.cpp
index faa595f8..7118c1c8 100644
--- a/common/fit.cpp
+++ b/common/fit.cpp
@@ -35,7 +35,8 @@ static std::vector<llama_device_memory_data> common_get_device_memory_data_impl(
         uint32_t & hp_ngl,
         uint32_t & hp_n_ctx_train,
         uint32_t & hp_n_expert,
-        ggml_log_level log_level) {
+        ggml_log_level log_level,
+        std::vector<size_t> * checkpoint_sizes = nullptr) {
     struct user_data_t {
         struct {
             ggml_log_callback callback;
@@ -97,6 +98,28 @@ static std::vector<llama_device_memory_data> common_get_device_memory_data_impl(
         }
     }
 
+    if (checkpoint_sizes) {
+        checkpoint_sizes->assign(nd + 1, 0);
+        const auto checkpoint_breakdown = llama_get_state_seq_device_buffer_sizes(ctx);
+        for (const auto & [buft, size] : checkpoint_breakdown) {
+            if (ggml_backend_buft_is_host(buft)) {
+                checkpoint_sizes->back() += size;
+                continue;
+            }
+
+            ggml_backend_dev_t dev = ggml_backend_buft_get_device(buft);
+            if (!dev) {
+                continue;
+            }
+            for (size_t i = 0; i < nd; ++i) {
+                if (dev == llama_model_get_device(model, i)) {
+                    (*checkpoint_sizes)[i] += size;
+                    break;
+                }
+            }
+        }
+    }
+
     {
         ggml_backend_dev_t cpu_dev = ggml_backend_dev_by_type(GGML_BACKEND_DEVICE_TYPE_CPU);
         if (cpu_dev == nullptr) {
@@ -162,8 +185,10 @@ common_device_memory_data_vec common_get_device_memory_data(
         uint32_t & hp_n_ctx_train,
         uint32_t & hp_n_expert,
         ggml_log_level log_level) {
+    std::vector<size_t> checkpoint_sizes;
     std::vector<llama_device_memory_data> impl = common_get_device_memory_data_impl(
-            path_model, mparams, cparams, devs, hp_ngl, hp_n_ctx_train, hp_n_expert, log_level);
+            path_model, mparams, cparams, devs, hp_ngl, hp_n_ctx_train, hp_n_expert, log_level,
+            &checkpoint_sizes);
 
     common_device_memory_data_vec ret(impl.size());
     for (size_t i = 0; i < impl.size(); i++) {
@@ -172,6 +197,7 @@ common_device_memory_data_vec common_get_device_memory_data(
         ret[i].model   = impl[i].mb.model;
         ret[i].context = impl[i].mb.context;
         ret[i].compute = impl[i].mb.compute;
+        ret[i].checkpoint = checkpoint_sizes[i];
     }
     return ret;
 }
@@ -186,6 +212,7 @@ static void common_params_fit_impl(
     constexpr int64_t MiB = 1024*1024;
     typedef std::vector<llama_device_memory_data> dmds_t;
     const llama_model_params default_mparams = llama_model_default_params();
+    const bool fixed_gpu_layers = mparams->n_gpu_layers != default_mparams.n_gpu_layers;
 
     std::vector<ggml_backend_dev_t> devs;
     uint32_t hp_ngl = 0; // hparams.n_gpu_layers
@@ -400,7 +427,7 @@ static void common_params_fit_impl(
                             sum_used_target -= margins[id];
                         }
                     }
-                    if (nd > 1) {
+                    if (nd > 1 && !fixed_gpu_layers) {
                         // for multiple devices we need to be more conservative in terms of how much context we think can fit:
                         //   - for dense models only whole layers can be assigned to devices
                         //   - for MoE models only whole tensors can be assigned to devices, which we estimate to be <= 1/3 of a layer
@@ -421,6 +448,31 @@ static void common_params_fit_impl(
                             sum_projected_used_min_ctx += dmds_min_ctx[id].mb.total();
                         }
                     }
+                    if (nd > 1 && fixed_gpu_layers) {
+                        uint32_t n_ctx_fit = n_ctx_max;
+                        bool fits = true;
+                        for (size_t id = 0; id < nd; id++) {
+                            const int64_t used_full = dmds_full[id].mb.total();
+                            const int64_t used_min  = dmds_min_ctx[id].mb.total();
+                            const int64_t target    = dmds_full[id].free - margins[id];
+                            if (used_full <= target) {
+                                continue;
+                            }
+                            uint32_t n_ctx_device = n_ctx_min_total;
+                            fits = fits && target >= used_min;
+                            if (target > used_min && used_full > used_min) {
+                                n_ctx_device += (n_ctx_max - n_ctx_min_total) * (target - used_min) / (used_full - used_min);
+                            }
+                            n_ctx_fit = std::min(n_ctx_fit, n_ctx_device);
+                        }
+                        const uint32_t align = 256 * n_streams;
+                        cparams->n_ctx = std::max(n_ctx_fit - n_ctx_fit % align, n_ctx_min_total);
+                        LOG_TRC("%s: context size reduced from %" PRIu32 " to %" PRIu32 " using per-device limits\n",
+                            __func__, n_ctx_max, cparams->n_ctx);
+                        if (fits) {
+                            return;
+                        }
+                    }
                     if (sum_used_target > sum_projected_used_min_ctx) {
                         // linear interpolation between minimum and maximum context size:
                         cparams->n_ctx += (n_ctx_max - n_ctx_min_total) * (sum_used_target - sum_projected_used_min_ctx)
@@ -460,7 +512,7 @@ static void common_params_fit_impl(
         throw common_params_fit_exception("was unable to fit model into system memory by reducing context, abort");
     }
 
-    if (mparams->n_gpu_layers != default_mparams.n_gpu_layers) {
+    if (fixed_gpu_layers) {
         throw common_params_fit_exception("n_gpu_layers already set by user to " + std::to_string(mparams->n_gpu_layers) + ", abort");
     }
     if (nd > 1) {
diff --git a/common/fit.h b/common/fit.h
index 824d386b..7022746a 100644
--- a/common/fit.h
+++ b/common/fit.h
@@ -51,6 +51,7 @@ struct common_device_memory_data {
     size_t  model;
     size_t  context;
     size_t  compute;
+    size_t  checkpoint;
 };
 
 using common_device_memory_data_vec = std::vector<common_device_memory_data>;
diff --git a/common/speculative.cpp b/common/speculative.cpp
index 77dfe953..f7139a42 100644
--- a/common/speculative.cpp
+++ b/common/speculative.cpp
@@ -159,8 +159,97 @@ struct common_speculative_impl {
     int64_t t_draft_us  = 0; // total time spent in generating drafts in this implementation in microseconds.
     int64_t t_accept_us = 0; // total time spent in accumulation of this implementation in microseconds.
 
-    common_speculative_impl(common_speculative_type type, uint32_t n_seq, int32_t n_max) : type(type), n_seq(n_seq), n_max(n_max) {}
+    // Adaptive draft length. The configured n_max is the known-good steady
+    // state. Ordinary tail rejection does not justify shortening because a
+    // wider verification batch is usually cheaper than another target decode
+    // round. Two consecutive drafts accepting fewer than three tokens identify
+    // a real phase change and step the limit down by one. A fully accepted draft
+    // steps it back up by one.
+    std::vector<int32_t> adaptive_limit;
+    std::vector<int32_t> adaptive_floor;
+    std::vector<int32_t> adaptive_ceiling;
+    std::vector<int32_t> n_last_draft;
+    std::vector<uint8_t> early_reject_streak;
+    std::vector<uint8_t> adaptive_limit_reached;
+    bool adaptive_n = false;
+
+    static constexpr int32_t adaptive_productive = 3;
+    static constexpr uint8_t adaptive_patience   = 2;
+
+    void update_adaptive_limit(llama_seq_id seq_id, int32_t n_accepted) {
+        if (!adaptive_n || seq_id < 0 || (size_t) seq_id >= adaptive_limit.size()) {
+            return;
+        }
+        const int32_t n_drafted = n_last_draft[seq_id];
+        if (n_drafted <= 0) {
+            return;
+        }
+        const bool limit_reached = adaptive_limit_reached[seq_id] != 0;
+        adaptive_limit_reached[seq_id] = 0;
+        if (!limit_reached) {
+            // Confidence stopping (p_min), a per-slot cap, or a decode failure
+            // already terminated this draft. Do not reinterpret that shorter
+            // observation as evidence for changing the phase-level ceiling.
+            early_reject_streak[seq_id] = 0;
+            n_last_draft[seq_id] = 0;
+            return;
+        }
+
+        if (n_accepted >= n_drafted) {
+            early_reject_streak[seq_id] = 0;
+            adaptive_limit[seq_id] = std::min(adaptive_ceiling[seq_id], adaptive_limit[seq_id] + 1);
+        } else if (n_accepted >= std::min(adaptive_productive, n_drafted)) {
+            early_reject_streak[seq_id] = 0;
+        } else if (++early_reject_streak[seq_id] >= adaptive_patience) {
+            early_reject_streak[seq_id] = 0;
+            adaptive_limit[seq_id] = std::max(adaptive_floor[seq_id], adaptive_limit[seq_id] - 1);
+        }
+        n_last_draft[seq_id] = 0;
+    }
+
+    // Reset on a new prompt / reused server slot; never mid-generation, since
+    // tracking phase changes within a response is the point of adaptation.
+    void reset_adaptive_limit(llama_seq_id seq_id) {
+        if (seq_id >= 0 && (size_t) seq_id < adaptive_limit.size()) {
+            adaptive_limit[seq_id]       = 0;
+            adaptive_floor[seq_id]       = 1;
+            adaptive_ceiling[seq_id]     = 0;
+            n_last_draft[seq_id]         = 0;
+            early_reject_streak[seq_id]  = 0;
+            adaptive_limit_reached[seq_id] = 0;
+        }
+    }
+
+    void record_adaptive_draft(llama_seq_id seq_id, size_t n_drafted) {
+        if (adaptive_n && seq_id >= 0 && (size_t) seq_id < n_last_draft.size()) {
+            n_last_draft[seq_id] = static_cast<int32_t>(n_drafted);
+            adaptive_limit_reached[seq_id] = n_drafted >= (size_t) adaptive_limit[seq_id];
+        }
+    }
 
+    // effective draft length for this step, never above the configured n_max
+    int32_t adaptive_n_draft(llama_seq_id seq_id, int32_t n_cfg, int32_t n_min) {
+        if (!adaptive_n || n_cfg <= 0 || seq_id < 0 || (size_t) seq_id >= adaptive_limit.size()) {
+            return n_cfg;
+        }
+        adaptive_floor[seq_id]   = std::clamp(n_min, 1, n_cfg);
+        adaptive_ceiling[seq_id] = n_cfg;
+        if (adaptive_limit[seq_id] <= 0) {
+            adaptive_limit[seq_id] = n_cfg;
+        }
+        adaptive_limit[seq_id] = std::clamp(adaptive_limit[seq_id], adaptive_floor[seq_id], n_cfg);
+        return adaptive_limit[seq_id];
+    }
+
+    common_speculative_impl(common_speculative_type type, uint32_t n_seq, int32_t n_max)
+        : type(type), n_seq(n_seq), n_max(n_max) {
+        adaptive_limit.assign(n_seq, 0);
+        adaptive_floor.assign(n_seq, 1);
+        adaptive_ceiling.assign(n_seq, 0);
+        n_last_draft.assign(n_seq, 0);
+        early_reject_streak.assign(n_seq, 0);
+        adaptive_limit_reached.assign(n_seq, 0);
+    }
     virtual ~common_speculative_impl() = default;
 
     virtual void begin(llama_seq_id seq_id, const llama_tokens & prompt) = 0;
@@ -195,6 +284,7 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
         }
 
         SPC_TRC("%s", "adding speculative implementation 'draft-simple'\n");
+        adaptive_n = this->params.adaptive;
         SPC_TRC("- n_max=%d, n_min=%d, p_min=%f\n", this->params.n_max, this->params.n_min, this->params.p_min);
         SPC_TRC("- gpu_layers=%d, cache_k=%s, cache_v=%s, ctx_tgt=%s, ctx_dft=%s, devices=[%s]\n",
                 this->params.n_gpu_layers,
@@ -255,8 +345,8 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
         llama_batch_free(batch);
     }
 
-    void begin(llama_seq_id /*seq_id*/, const llama_tokens & /*prompt*/) override {
-        // noop
+    void begin(llama_seq_id seq_id, const llama_tokens & /*prompt*/) override {
+        reset_adaptive_limit(seq_id);
     }
 
     bool process(const llama_batch & batch) override {
@@ -319,6 +409,8 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
 
                 auto * smpl = smpls[seq_id].get();
 
+                const int32_t n_draft_eff = adaptive_n_draft(seq_id, params.n_max, params.n_min);
+
                 common_sampler_sample(smpl, ctx_dft, i_batch, true);
                 ++i_batch;
 
@@ -348,7 +440,7 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl {
 
                 result.push_back(id);
 
-                if ((params.n_max <= (int) result.size()) ||
+                if ((n_draft_eff <= (int) result.size()) ||
                     (dp.n_max > 0 && dp.n_max <= (int) result.size())) {
                     drafting[seq_id] = false;
                     n_drafting--;
@@ -458,6 +550,7 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl {
         , params(params.draft)
     {
         SPC_TRC("%s", "adding speculative implementation 'draft-eagle3'\n");
+        adaptive_n = this->params.adaptive;
         SPC_TRC("- n_max=%d, n_min=%d, p_min=%f, backend_sampling=%d\n", params.draft.n_max, params.draft.n_min, params.draft.p_min, (int) params.draft.backend_sampling);
 
         auto * ctx_tgt = this->params.ctx_tgt;
@@ -554,6 +647,8 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl {
     }
 
     void begin(llama_seq_id seq_id, const llama_tokens & prompt) override {
+        reset_adaptive_limit(seq_id);
+
         const int32_t N = (int32_t) prompt.size();
         if (N <= 0) {
             return;
@@ -779,6 +874,8 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl {
 
                 auto * smpl = smpls[seq_id].get();
 
+                const int32_t n_draft_eff = adaptive_n_draft(seq_id, params.n_max, params.n_min);
+
                 common_sampler_sample(smpl, ctx_dft, i_batch, true);
                 // pre-norm hidden state of this position becomes g_embd for the next step
                 const float * prenorm = llama_get_embeddings_nextn_ith(ctx_dft, i_batch);
@@ -810,7 +907,7 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl {
 
                 result.push_back(id);
 
-                if (params.n_max <= (int) result.size()) {
+                if (n_draft_eff <= (int) result.size()) {
                     drafting[seq_id] = false;
                     n_drafting--;
                     continue;
@@ -992,6 +1089,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
         }
 
         LOG_INF("%s: adding speculative implementation '%s'\n", __func__, common_speculative_type_to_str(type).c_str());
+        adaptive_n = this->params.adaptive;
         LOG_INF("%s: - n_max=%d, n_min=%d, p_min=%.2f\n", __func__, this->params.n_max, this->params.n_min, this->params.p_min);
         LOG_INF("%s: - block_size=%d, mask_token_id=%d, n_extract=%u, sample_from_anchor=%s\n", __func__,
                 block_size, mask_token_id, target_layer_ids_n, sample_from_anchor ? "true" : "false");
@@ -1070,6 +1168,8 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
     }
 
     void begin(llama_seq_id seq_id, const llama_tokens & prompt) override {
+
+        reset_adaptive_limit(seq_id);
         if (seq_id < 0 || seq_id >= (llama_seq_id) n_seq) {
             return;
         }
@@ -1199,7 +1299,10 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
 
             const int32_t n = (int32_t) dp.pos0;
 
-            const int32_t n_draft = params.n_max;
+            // DFlash decodes the whole block in one pass, so a shorter block does
+            // not save draft time -- but it does shrink the target's verification
+            // batch, which is where the cost actually is.
+            const int32_t n_draft = adaptive_n_draft(seq_id, params.n_max, params.n_min);
 
             const int32_t n_block_tokens = n_draft + (is_dspark && sample_from_anchor ? 0 : 1);
             i_block_beg[seq_id] = batch.n_tokens;
@@ -1377,6 +1480,7 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
         n_mtp_layers = std::max(1, (int) llama_model_n_layer_nextn(llama_get_model(ctx_dft)));
 
         SPC_TRC("%s", "adding speculative implementation 'draft-mtp'\n");
+        adaptive_n = this->params.adaptive;
         SPC_TRC("- n_max=%d, n_min=%d, p_min=%.2f, n_embd=%d, backend_sampling=%d\n", this->params.n_max, this->params.n_min, this->params.p_min, n_embd, (int) this->params.backend_sampling);
         SPC_TRC("- gpu_layers=%d, cache_k=%s, cache_v=%s, ctx_tgt=%s, ctx_dft=%s, devices=[%s]\n",
                 this->params.n_gpu_layers,
@@ -1464,6 +1568,8 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
     }
 
     void begin(llama_seq_id seq_id, const llama_tokens & prompt) override {
+
+        reset_adaptive_limit(seq_id);
         const int32_t N = (int32_t) prompt.size();
         if (N <= 0) {
             return;
@@ -1668,6 +1774,10 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
 
                 auto * smpl = smpls[seq_id].get();
 
+                // MTP drafts sequentially, so a shorter draft saves draft passes
+                // as well as target verification work.
+                const int32_t n_draft_eff = adaptive_n_draft(seq_id, params.n_max, params.n_min);
+
                 common_sampler_sample(smpl, ctx_dft, i_last[seq_id], true);
                 const float * h_row = llama_get_embeddings_nextn_ith(ctx_dft, i_last[seq_id]);
 
@@ -1697,7 +1807,7 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl {
 
                 result.push_back(id);
 
-                if (params.n_max <= (int) result.size()) {
+                if (n_draft_eff <= (int) result.size()) {
                     drafting[seq_id] = false;
                     n_drafting--;
                     continue;
@@ -2544,6 +2654,10 @@ common_speculative_init_result::common_speculative_init_result(
 
     if (spec_mtp) {
         cparams.ctx_type = LLAMA_CONTEXT_TYPE_MTP;
+        cparams.pipeline_parallel_type = LLAMA_PIPELINE_PARALLEL_TYPE_DISABLED;
+        // use the same (possibly fit-reduced) context size as the target
+        // avoids MTP draft reserving flash attention workspace for full n_ctx
+        cparams.n_ctx = llama_n_ctx(ctx_tgt);
     }
 
     // the draft context holds as many tokens per sequence as the target context
@@ -2863,6 +2977,7 @@ void common_speculative_draft(common_speculative * spec) {
 
                     // remember which implementation was used
                     spec->impl_last[seq_id] = impl.get();
+                    impl->record_adaptive_draft(seq_id, result.size());
 
                     impl->n_gen_drafts++;
                     impl->n_gen_tokens += result.size();
@@ -2889,7 +3004,11 @@ void common_speculative_draft(common_speculative * spec) {
     }
 }
 
-void common_speculative_accept(common_speculative * spec, llama_seq_id seq_id, uint16_t n_accepted) {
+void common_speculative_accept(
+        common_speculative * spec,
+        llama_seq_id         seq_id,
+        uint16_t             n_accepted,
+        int32_t              n_accepted_observed) {
     common_speculative_impl * impl = spec->impl_last[seq_id];
 
     if (impl == nullptr) {
@@ -2913,6 +3032,7 @@ void common_speculative_accept(common_speculative * spec, llama_seq_id seq_id, u
             impl->n_acc_tokens += n_accepted;
         }
 
+        impl->update_adaptive_limit(seq_id, n_accepted_observed >= 0 ? n_accepted_observed : n_accepted);
         impl->accept(seq_id, n_accepted, false);
         impl->n_call_accept++;
     }
diff --git a/common/speculative.h b/common/speculative.h
index c968750e..efdd4cd2 100644
--- a/common/speculative.h
+++ b/common/speculative.h
@@ -82,8 +82,14 @@ bool common_speculative_process(common_speculative * spec, const llama_batch & b
 // generate drafts for the sequences specified with `common_speculative_get_draft_params`
 void common_speculative_draft(common_speculative * spec);
 
-// informs the speculative context that n_accepted tokens were accepted by the target model
-void common_speculative_accept(common_speculative * spec, llama_seq_id, uint16_t n_accepted);
+// informs the speculative context that n_accepted tokens were accepted by the target model.
+// n_accepted_observed can preserve the original acceptance measurement when a
+// rollback implementation has to replay a different number of physical tokens.
+void common_speculative_accept(
+        common_speculative * spec,
+        llama_seq_id         seq_id,
+        uint16_t             n_accepted,
+        int32_t              n_accepted_observed = -1);
 
 // (optional) get/set internal state
 bool common_speculative_get_state(common_speculative * spec, llama_seq_id seq_id, std::vector<uint8_t> & data);
diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h
index 224bdef9..318b9b84 100644
--- a/ggml/include/ggml.h
+++ b/ggml/include/ggml.h
@@ -2502,6 +2502,13 @@ extern "C" {
     GGML_API enum ggml_prec ggml_flash_attn_ext_get_prec(
             const struct ggml_tensor * a);
 
+    GGML_API void ggml_flash_attn_ext_set_force_vec(
+            struct ggml_tensor * a,
+            bool                 force_vec);
+
+    GGML_API bool ggml_flash_attn_ext_get_force_vec(
+            const struct ggml_tensor * a);
+
     // Use finite mask entries as a sparse K/V set. Set 0 to disable.
     // n_kv_max must bound the number of finite entries in every mask row.
     GGML_API void ggml_flash_attn_ext_set_n_kv_max(
diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp
index 3ec40fb1..87667beb 100644
--- a/ggml/src/ggml-backend-meta.cpp
+++ b/ggml/src/ggml-backend-meta.cpp
@@ -1712,7 +1712,12 @@ static ggml_backend_buffer_t ggml_backend_meta_buffer_type_alloc_buffer(ggml_bac
 struct ggml_backend_buffer * ggml_backend_meta_alloc_ctx_tensors_from_buft(struct ggml_context * ctx, ggml_backend_buffer_type_t buft) {
     const size_t n_simple_bufts = ggml_backend_meta_buft_n_bufts(buft);
 
-    constexpr size_t compute_headroom = 16; // Maximum number of views per statically allocated tensor that can be created between evals.
+    // Views of the static tensors that are created between graph evals are stored in the compute
+    // containers. The number of such views is proportional to the number of tensors in the graph
+    // that share the buffer, which for hybrid recurrent models with n_rs_seq snapshotting can be
+    // much larger than 16 per static tensor (e.g. Qwen35: ~2*(n_rs_seq+1) views per recurrent
+    // layer are created for the conv-state snapshot copies). Size the headroom accordingly.
+    constexpr size_t compute_headroom = 128;
     const ggml_init_params params_static = {
         /*.mem_size   =*/ ggml_get_mem_size(ctx),
         /*.mem_buffer =*/ nullptr,
diff --git a/ggml/src/ggml-cuda/fattn.cu b/ggml/src/ggml-cuda/fattn.cu
index ceb47279..3ef68e70 100644
--- a/ggml/src/ggml-cuda/fattn.cu
+++ b/ggml/src/ggml-cuda/fattn.cu
@@ -610,6 +610,15 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
     // 192 satisfies % 64 == 0 but has no vec instance (DKQ != DV); force it onto the MMA path.
     const bool can_use_vector_kernel = Q->ne[0] <= 256 && Q->ne[0] % 64 == 0 && Q->ne[0] != 192 && K->ne[1] % FATTN_KQ_STRIDE == 0;
 
+#ifdef GGML_USE_HIP
+    // HIP quantized-KV TILE/MMA paths materialize large F16 temporary buffers;
+    // VEC dequantizes in-register and is also the safe path on RDNA2.
+    if (ggml_flash_attn_ext_get_force_vec(dst) &&
+            (ggml_is_quantized(K->type) || ggml_is_quantized(V->type)) && can_use_vector_kernel) {
+        return BEST_FATTN_KERNEL_VEC;
+    }
+#endif
+
     // If Turing tensor cores are available, use them:
     if (turing_mma_available(cc) && Q->ne[0] != 40 && Q->ne[0] != 72) {
         if (can_use_vector_kernel) {
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index a9038f1f..200313b5 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -4426,13 +4426,16 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend,
     bool cuda_graph_update_required = false;
     const void * graph_key = nullptr;
 
+    // op timing instruments each node with stream events, which is not possible during capture
+    const bool op_timing = getenv("GGML_CUDA_OP_TIMING") != nullptr;
+
 #ifdef USE_CUDA_GRAPH
     graph_key = ggml_cuda_graph_get_key(cgraph);
 
     ggml_cuda_graph_set_enabled(cuda_ctx, graph_key);
 
     ggml_cuda_graph * graph = cuda_ctx->cuda_graph(graph_key);
-    if (graph->is_enabled()) {
+    if (!op_timing && graph->is_enabled()) {
         const bool graph_compatible = ggml_cuda_graph_check_compability(cgraph);
         if (graph_compatible) {
             const bool properties_changed = ggml_cuda_graph_update_required(cuda_ctx, cgraph);
diff --git a/ggml/src/ggml-cuda/mmq-vec-dot.cuh b/ggml/src/ggml-cuda/mmq-vec-dot.cuh
index 4d1c398f..864d161d 100644
--- a/ggml/src/ggml-cuda/mmq-vec-dot.cuh
+++ b/ggml/src/ggml-cuda/mmq-vec-dot.cuh
@@ -340,6 +340,17 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
             load_ldmatrix(A[n], x_qs + (i0 + n*tile_A::I)*sram_stride + k0, sram_stride);
         }
 
+        // Row scale pair (d, m) is invariant over the j0 loop; load it once per element.
+        float2 dmA_reg[ntx][tile_C::ne];
+#pragma unroll
+        for (int n = 0; n < ntx; ++n) {
+#pragma unroll
+            for (int l = 0; l < tile_C::ne; ++l) {
+                const int i = i0 + n*tile_A::I + tile_C::get_i(l);
+                dmA_reg[n][l] = __half22float2(x_dm[i*sram_stride + k0/QI8_1]);
+            }
+        }
+
 #pragma unroll
         for (int j0 = 0; j0 < J; j0 += ntx*tile_C::J) {
             tile_B B;
@@ -355,10 +366,8 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
 
 #pragma unroll
                 for (int l = 0; l < tile_C::ne; ++l) {
-                    const int i = i0 + n*tile_A::I + tile_C::get_i(l);
-                    float2 dmA = __half22float2(x_dm[i*sram_stride + k0/QI8_1]);
-                    sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA.x*dsB.x*C.x[l];
-                    sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA.y*dsB.y;
+                    sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA_reg[n][l].x*dsB.x*C.x[l];
+                    sum[(j0/tile_C::J + n)*tile_C::ne + l] += dmA_reg[n][l].y*dsB.y;
                 }
             }
         }
diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu
index 6305230b..9a7f6c67 100644
--- a/ggml/src/ggml-cuda/mmvq.cu
+++ b/ggml/src/ggml-cuda/mmvq.cu
@@ -50,9 +50,9 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type)
         case GGML_TYPE_NVFP4:   return vec_dot_nvfp4_q8_1;
         case GGML_TYPE_Q2_K:    return vec_dot_q2_K_q8_1;
         case GGML_TYPE_Q3_K:    return vec_dot_q3_K_q8_1;
-        case GGML_TYPE_Q4_K:    return vec_dot_q4_K_q8_1;
-        case GGML_TYPE_Q5_K:    return vec_dot_q5_K_q8_1;
-        case GGML_TYPE_Q6_K:    return vec_dot_q6_K_q8_1;
+        case GGML_TYPE_Q4_K:    return vec_dot_q4_K_q8_1_vdr4;
+        case GGML_TYPE_Q5_K:    return vec_dot_q5_K_q8_1_vdr4;
+        case GGML_TYPE_Q6_K:    return vec_dot_q6_K_q8_1_vdr2;
         case GGML_TYPE_IQ2_XXS: return vec_dot_iq2_xxs_q8_1;
         case GGML_TYPE_IQ2_XS:  return vec_dot_iq2_xs_q8_1;
         case GGML_TYPE_IQ2_S:   return vec_dot_iq2_s_q8_1;
@@ -99,6 +99,7 @@ enum mmvq_parameter_table_id {
     MMVQ_PARAMETERS_GCN,
     MMVQ_PARAMETERS_RDNA2,
     MMVQ_PARAMETERS_RDNA3_0,
+    MMVQ_PARAMETERS_RDNA3_5,
     MMVQ_PARAMETERS_RDNA4,
     MMVQ_PARAMETERS_GB10
 };
@@ -108,7 +109,9 @@ static constexpr __device__ mmvq_parameter_table_id get_device_table_id() {
     return MMVQ_PARAMETERS_RDNA4;
 #elif defined(RDNA3_0)
     return MMVQ_PARAMETERS_RDNA3_0;
-#elif defined(RDNA2) || defined(RDNA3_5)
+#elif defined(RDNA3_5)
+    return MMVQ_PARAMETERS_RDNA3_5;
+#elif defined(RDNA2)
     return MMVQ_PARAMETERS_RDNA2;
 #elif defined(GCN) || defined(CDNA)
     return MMVQ_PARAMETERS_GCN;
@@ -128,7 +131,10 @@ static __host__ mmvq_parameter_table_id get_device_table_id(int cc) {
     if (GGML_CUDA_CC_IS_RDNA3_0(cc)) {
         return MMVQ_PARAMETERS_RDNA3_0;
     }
-    if (GGML_CUDA_CC_IS_RDNA2(cc) || GGML_CUDA_CC_IS_RDNA3_5(cc)) {
+    if (GGML_CUDA_CC_IS_RDNA3_5(cc)) {
+        return MMVQ_PARAMETERS_RDNA3_5;
+    }
+    if (GGML_CUDA_CC_IS_RDNA2(cc)) {
         return MMVQ_PARAMETERS_RDNA2;
     }
     if (GGML_CUDA_CC_IS_GCN(cc) || GGML_CUDA_CC_IS_CDNA(cc)) {
@@ -271,7 +277,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type
         case GGML_TYPE_Q3_K:    return 4;
         case GGML_TYPE_Q4_0:    return 7;
         case GGML_TYPE_Q4_1:    return 7;
-        case GGML_TYPE_Q4_K:    return 4;
+        case GGML_TYPE_Q4_K:    return 7;
         case GGML_TYPE_Q5_0:    return 7;
         case GGML_TYPE_Q5_1:    return 7;
         case GGML_TYPE_Q5_K:    return 5;
@@ -509,6 +515,24 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d
         }
         return 1;
     }
+    if (table_id == MMVQ_PARAMETERS_RDNA3_5) {
+        // gfx1151 (Strix Halo iGPU): nwarps=1 (the RDNA2 table) underutilizes the
+        // wave32 datapath on the large-K decode matmuls; nwarps=8 (the RDNA3_0
+        // table) over-parallelizes the small ones. Swept 2025-08: nwarps=2 wins
+        // (~+0.6% decode on Qwen3.6-35B-A3B Q8_0), nwarps=4 regresses.
+        // Apply to the whole mmvq range (ncols_dst 1..8), not just decode: the
+        // speculative verify batch (n_draft+1 tokens) must use the same nwarps
+        // as decode so its per-row dot-product accumulation is bit-identical.
+        if (ncols_dst <= MMVQ_MAX_BATCH_SIZE) {
+            switch (type) {
+                case GGML_TYPE_Q8_0:
+                    return 2;
+                default:
+                    return 1;
+            }
+        }
+        return 1;
+    }
     if (table_id == MMVQ_PARAMETERS_TURING) {
         if (ncols_dst == 1) {
             switch (type) {
diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh
index f2a6f200..d42feb23 100644
--- a/ggml/src/ggml-cuda/vecdotq.cuh
+++ b/ggml/src/ggml-cuda/vecdotq.cuh
@@ -240,7 +240,16 @@ template <int vdr> static __device__ __forceinline__ float vec_dot_q5_1_q8_1_imp
     return sumi*d5d8 + m5s8 / (QI5_1 / vdr);
 }
 
+#if defined(RDNA4) || defined(RDNA3_0)
+// VDR=4 measured on gfx1200/gfx1201 (RX 9000): decode -10..-33% on the
+// compute-bound shapes, neutral on the DRAM-bound lm_head. RDNA3_0 (gfx1100-
+// gfx1103, RX 7900 XTX) verified 2026-08-28: tg128 123.74 -> 127.7x (+3.x%),
+// PPL 24.4430 vs 24.44xx (near-lossless), greedy byte-identical. RDNA3_5
+// (gfx115x) keeps VDR=2 pending verification on those GPUs.
+#define VDR_Q8_0_Q8_1_MMVQ 4
+#else
 #define VDR_Q8_0_Q8_1_MMVQ 2
+#endif
 #define VDR_Q8_0_Q8_1_MMQ 8
 
 template <typename T, int vdr> static __device__ __forceinline__ T vec_dot_q8_0_q8_1_impl(
@@ -501,7 +510,7 @@ static __device__ __forceinline__ float vec_dot_q3_K_q8_1_impl_mmq(
     return d3*d8 * sumi;
 }
 
-#define VDR_Q4_K_Q8_1_MMVQ 2
+#define VDR_Q4_K_Q8_1_MMVQ 4
 #define VDR_Q4_K_Q8_1_MMQ  8
 
 // contiguous v/x values
@@ -529,6 +538,34 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq(
     return dm4f.x*sumf_d - dm4f.y*sumf_m;
 }
 
+// VDR=4 variant: processes two adjacent 16-element chunks (32 elements).
+// Both chunks share the q8_1 block pair, the two sub-scales, the two mins and
+// the d8 values, so the loads are amortized over twice the dp4a work of the
+// VDR=2 kernel.
+static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq4(
+    const int * __restrict__ v, const int * __restrict__ u, const uint8_t * __restrict__ sc,
+    const uint8_t * __restrict__ m, const half2 & dm4, const float * __restrict__ d8) {
+
+    float sumf_d = 0.0f;
+    float sumf_m = 0.0f;
+
+#pragma unroll
+    for (int i = 0; i < 2*QR4_K; ++i) {
+        const int v0i = (v[2*(i>>1)]     >> (4*(i&1))) & 0x0F0F0F0F;
+        const int v1i = (v[2*(i>>1) + 1] >> (4*(i&1))) & 0x0F0F0F0F;
+
+        const int dot1 = ggml_cuda_dp4a(v1i, u[2*i+1], ggml_cuda_dp4a(v0i, u[2*i+0], 0)); // SIMD dot product
+        const int dot2 = ggml_cuda_dp4a(0x01010101, u[2*i+1], ggml_cuda_dp4a(0x01010101, u[2*i+0], 0)); // sum of u
+
+        sumf_d += d8[i&1] * (dot1 * sc[i&1]);
+        sumf_m += d8[i&1] * (dot2 * m[i&1]);  // multiply constant part of q4_K with sum of q8_1 values
+    }
+
+    const float2 dm4f = __half22float2(dm4);
+
+    return dm4f.x*sumf_d - dm4f.y*sumf_m;
+}
+
 // contiguous v/x + u/y values
 static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_mmq(
     const int * __restrict__ v, const int * __restrict__ u, const uint8_t * __restrict__ sc,
@@ -557,7 +594,7 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_mmq(
     return dm4f.x*sumf_d - dm4f.y*sumf_m;
 }
 
-#define VDR_Q5_K_Q8_1_MMVQ 2
+#define VDR_Q5_K_Q8_1_MMVQ 4
 #define VDR_Q5_K_Q8_1_MMQ  8
 
 // contiguous v/x values
@@ -592,6 +629,41 @@ static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_vmmq(
     return dm5f.x*sumf_d - dm5f.y*sumf_m;
 }
 
+// VDR=4 variant: processes two adjacent 16-element chunks (32 elements).
+// Both chunks share the q8_1 block pair, the two sub-scales, the two mins and
+// the d8 values, so the loads are amortized over twice the dp4a work of the
+// VDR=2 kernel.
+static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_vmmq4(
+    const int * __restrict__ vl, const int * __restrict__ vh, const int * __restrict__ u, const uint8_t * __restrict__ sc,
+    const uint8_t * __restrict__ m, const half2 & dm5, const float * __restrict__ d8) {
+
+    float sumf_d = 0.0f;
+    float sumf_m = 0.0f;
+
+#pragma unroll
+    for (int i = 0; i < 2*QR5_K; ++i) {
+        const int vl0i = (vl[2*(i>>1)]     >> (4*(i&1))) & 0x0F0F0F0F;
+        const int vl1i = (vl[2*(i>>1) + 1] >> (4*(i&1))) & 0x0F0F0F0F;
+
+        const int vh0i = ((vh[2*(i>>1)]     >> (i&1)) << 4) & 0x10101010;
+        const int vh1i = ((vh[2*(i>>1) + 1] >> (i&1)) << 4) & 0x10101010;
+
+        const int v0i = vl0i | vh0i;
+        const int v1i = vl1i | vh1i;
+
+        const int dot1 = ggml_cuda_dp4a(v0i, u[2*i+0], ggml_cuda_dp4a(v1i, u[2*i+1], 0)); // SIMD dot product
+        const int dot2 = ggml_cuda_dp4a(0x01010101, u[2*i+0], ggml_cuda_dp4a(0x01010101, u[2*i+1], 0)); // sum of u
+
+        sumf_d += d8[i&1] * (dot1 * sc[i&1]);
+        sumf_m += d8[i&1] * (dot2 * m[i&1]);
+
+    }
+
+    const float2 dm5f = __half22float2(dm5);
+
+    return dm5f.x*sumf_d - dm5f.y*sumf_m;
+}
+
 // contiguous v/x + u/y values
 static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_mmq(
     const int * __restrict__ v, const int * __restrict__ u, const uint8_t * __restrict__ sc,
@@ -620,7 +692,7 @@ static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_mmq(
     return dm4f.x*sumf_d - dm4f.y*sumf_m;
 }
 
-#define VDR_Q6_K_Q8_1_MMVQ 1
+#define VDR_Q6_K_Q8_1_MMVQ 2
 #define VDR_Q6_K_Q8_1_MMQ  8
 
 // contiguous v/x values
@@ -646,6 +718,35 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmvq(
     return d*sumf;
 }
 
+// VDR=2 variant: processes two adjacent 8-element chunks (16 elements).
+// Both chunks share the q8_1 block pair, the two sub-scales and the d8 values,
+// so the loads are amortized over twice the dp4a work of the VDR=1 kernel.
+static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmvq_vdr2(
+    const int & vl0, const int & vl1, const int & vh0, const int & vh1,
+    const int * __restrict__ u, const int8_t * __restrict__ scales,
+    const float & d, const float * __restrict__ d8) {
+
+    float sumf = 0.0f;
+
+#pragma unroll
+    for (int i = 0; i < 2*QR6_K; ++i) {
+        const int sc = scales[4*(i&1)];
+
+        const int vl = (i < QR6_K) ? vl0 : vl1;
+        const int vh = (i < QR6_K) ? vh0 : vh1;
+
+        const int vil = (vl >> (4*(i&1))) & 0x0F0F0F0F;
+
+        const int vih = ((vh >> (4*(i&1))) << 4) & 0x30303030;
+
+        const int vi = __vsubss4((vil | vih), 0x20202020); // vi = (vil | vih) - 32
+
+        sumf += d8[i&1] * (ggml_cuda_dp4a(vi, u[i], 0) * sc); // SIMD dot product
+    }
+
+    return d*sumf;
+}
+
 // contiguous v/x + u/y values
 static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmq(
     const int * __restrict__ v, const int * __restrict__ u, const int8_t * __restrict__ sc,
@@ -965,6 +1066,59 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1(
     return vec_dot_q4_K_q8_1_impl_vmmq(v, u, sc, m, bq4_K->dm, d8);
 }
 
+// VDR=4 entry point: iqs must be a multiple of 4 (the mmvq kernel strides kqs
+// by VDR). Processes 32 elements per call, splitting the ql/u loads over two
+// 16-element chunks that share the q8_1 block pair, the sub-scale/min pair and
+// the d8 values.
+static __device__ __forceinline__ float vec_dot_q4_K_q8_1_vdr4(
+    const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
+
+    const block_q4_K * bq4_K = (const block_q4_K *) vbq + kbx;
+
+    int v[4];
+    int u[4*QR4_K];
+    float d8[2*QR4_K];
+
+    // iqs is in 0,4..28. The two 16-element chunks at iqs and iqs+2 land in the
+    // same (iqs/4, (iqs%4)/2) cell, so bq8_offset and the scale pair are shared.
+    const int bq8_offset = QR4_K * ((iqs/2) / (QI8_1/2));
+    const int i8 = (iqs/2) % 4;
+
+    const int * q4 = (const int *)(bq4_K->qs + 16 * bq8_offset + 4 * i8);
+    v[0] = q4[0];
+    v[1] = q4[4];
+    v[2] = q4[1];
+    v[3] = q4[5];
+
+    const uint16_t * scales = (const uint16_t *)bq4_K->scales;
+    uint16_t aux[2];
+    const int j = bq8_offset/2;
+    if (j < 2) {
+        aux[0] = scales[j+0] & 0x3f3f;
+        aux[1] = scales[j+2] & 0x3f3f;
+    } else {
+        aux[0] = ((scales[j+2] >> 0) & 0x0f0f) | ((scales[j-2] & 0xc0c0) >> 2);
+        aux[1] = ((scales[j+2] >> 4) & 0x0f0f) | ((scales[j-0] & 0xc0c0) >> 2);
+    }
+    const uint8_t * sc = (const uint8_t *)aux;
+    const uint8_t * m  = sc + 2;
+
+    // i8 is even, so i8 and i8+1 are adjacent int32 groups in each q8_1 block
+    d8[0] = __low2float(bq8_1[bq8_offset + 0].ds);
+    d8[1] = __low2float(bq8_1[bq8_offset + 1].ds);
+
+#pragma unroll
+    for (int i = 0; i < 2*QR4_K; ++i) {
+        const block_q8_1 * bq8i = bq8_1 + bq8_offset + (i&1);
+
+        const int * q8 = (const int *)bq8i->qs + i8 + (i>>1);
+        u[2*i+0] = q8[0];
+        u[2*i+1] = q8[4];
+    }
+
+    return vec_dot_q4_K_q8_1_impl_vmmq4(v, u, sc, m, bq4_K->dm, d8);
+}
+
 static __device__ __forceinline__ float vec_dot_q5_K_q8_1(
     const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
 
@@ -1016,6 +1170,64 @@ static __device__ __forceinline__ float vec_dot_q5_K_q8_1(
     return vec_dot_q5_K_q8_1_impl_vmmq(vl, vh, u, sc, m, bq5_K->dm, d8);
 }
 
+// VDR=4 entry point: iqs must be a multiple of 4 (the mmvq kernel strides kqs
+// by VDR). Processes 32 elements per call, splitting the ql/qh/u loads over two
+// 16-element chunks that share the q8_1 block pair, the sub-scale/min pair and
+// the d8 values.
+static __device__ __forceinline__ float vec_dot_q5_K_q8_1_vdr4(
+    const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
+
+    const block_q5_K * bq5_K = (const block_q5_K *) vbq + kbx;
+
+    int   vl[4];
+    int   vh[4];
+    int    u[4*QR5_K];
+    float d8[2*QR5_K];
+
+    const int bq8_offset = QR5_K * ((iqs/2) / (QI8_1/2));
+    const int i8 = (iqs/2) % 4;
+
+    const int * ql = (const int *)(bq5_K->qs + 16 * bq8_offset + 4 * i8);
+    const int * qh = (const int *)(bq5_K->qh + 4 * i8);
+
+    vl[0] = ql[0];
+    vl[1] = ql[4];
+    vl[2] = ql[1];
+    vl[3] = ql[5];
+
+    vh[0] = qh[0] >> bq8_offset;
+    vh[1] = qh[4] >> bq8_offset;
+    vh[2] = qh[1] >> bq8_offset;
+    vh[3] = qh[5] >> bq8_offset;
+
+    const uint16_t * scales = (const uint16_t *)bq5_K->scales;
+    uint16_t aux[2];
+    const int j = bq8_offset/2;
+    if (j < 2) {
+        aux[0] = scales[j+0] & 0x3f3f;
+        aux[1] = scales[j+2] & 0x3f3f;
+    } else {
+        aux[0] = ((scales[j+2] >> 0) & 0x0f0f) | ((scales[j-2] & 0xc0c0) >> 2);
+        aux[1] = ((scales[j+2] >> 4) & 0x0f0f) | ((scales[j-0] & 0xc0c0) >> 2);
+    }
+    const uint8_t * sc = (const uint8_t *)aux;
+    const uint8_t * m  = sc + 2;
+
+    d8[0] = __low2float(bq8_1[bq8_offset + 0].ds);
+    d8[1] = __low2float(bq8_1[bq8_offset + 1].ds);
+
+#pragma unroll
+    for (int i = 0; i < 2*QR5_K; ++i) {
+        const block_q8_1 * bq8i = bq8_1 + bq8_offset + (i&1);
+
+        const int * q8 = (const int *)bq8i->qs + i8 + (i>>1);
+        u[2*i+0] = q8[0];
+        u[2*i+1] = q8[4];
+    }
+
+    return vec_dot_q5_K_q8_1_impl_vmmq4(vl, vh, u, sc, m, bq5_K->dm, d8);
+}
+
 static __device__ __forceinline__ float vec_dot_q6_K_q8_1(
     const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
 
@@ -1042,6 +1254,41 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1(
     return vec_dot_q6_K_q8_1_impl_mmvq(vl, vh, u, scales, bq6_K->d, d8);
 }
 
+// VDR=2 entry point: iqs must be even (the mmvq kernel strides kqs by VDR).
+// Processes 16 elements per call, splitting the ql/qh/u loads over two chunks.
+static __device__ __forceinline__ float vec_dot_q6_K_q8_1_vdr2(
+    const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
+
+    const block_q6_K * bq6_K = (const block_q6_K *) vbq + kbx;
+
+    // both chunks share these offsets: iqs+1 lands in the same (iqs/16, (iqs%16)/8) cell
+    const int bq8_offset = 2 * QR6_K * (iqs / (QI6_K/2)) + (iqs % (QI6_K/2)) / (QI6_K/4);
+    const int scale_offset = (QI6_K/4) * (iqs / (QI6_K/2)) + (iqs % (QI6_K/2)) / (QI6_K/8);
+    const int vh_shift = 2 * ((iqs % (QI6_K/2)) / (QI6_K/4));
+    const int vh_idx = (QI6_K/4) * (iqs / (QI6_K/2)) + iqs % (QI6_K/4);
+
+    const int vl0 = get_int_b4(bq6_K->ql, iqs);
+    const int vl1 = get_int_b4(bq6_K->ql, iqs + 1);
+    const int vh0 = get_int_b2(bq6_K->qh, vh_idx) >> vh_shift;
+    const int vh1 = get_int_b2(bq6_K->qh, vh_idx + 1) >> vh_shift;
+
+    const int8_t * scales = bq6_K->scales + scale_offset;
+
+    int    u[2*QR6_K];
+    float d8[QR6_K];
+
+    // iqs is even, so iqs%QI8_1 and (iqs%QI8_1)+1 are adjacent int32 groups in the block
+    const int i8 = iqs % QI8_1;
+    u[0] = get_int_b4(bq8_1[bq8_offset + 0].qs, i8);
+    u[1] = get_int_b4(bq8_1[bq8_offset + 2].qs, i8);
+    u[2] = get_int_b4(bq8_1[bq8_offset + 0].qs, i8 + 1);
+    u[3] = get_int_b4(bq8_1[bq8_offset + 2].qs, i8 + 1);
+    d8[0] = __low2float(bq8_1[bq8_offset + 0].ds);
+    d8[1] = __low2float(bq8_1[bq8_offset + 2].ds);
+
+    return vec_dot_q6_K_q8_1_impl_mmvq_vdr2(vl0, vl1, vh0, vh1, u, scales, bq6_K->d, d8);
+}
+
 #define VDR_IQ2_XXS_Q8_1_MMVQ 2
 #define VDR_IQ2_XXS_Q8_1_MMQ  2
 
diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c
index 17528140..05b370ff 100644
--- a/ggml/src/ggml.c
+++ b/ggml/src/ggml.c
@@ -5558,6 +5558,18 @@ enum ggml_prec ggml_flash_attn_ext_get_prec(
     return (enum ggml_prec) prec_i32;
 }
 
+void ggml_flash_attn_ext_set_force_vec(
+        struct ggml_tensor * a,
+        bool                 force_vec) {
+    GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT);
+    ggml_set_op_params_i32(a, 5, force_vec ? 1 : 0);
+}
+
+bool ggml_flash_attn_ext_get_force_vec(const struct ggml_tensor * a) {
+    GGML_ASSERT(a->op == GGML_OP_FLASH_ATTN_EXT);
+    return ggml_get_op_params_i32(a, 5) != 0;
+}
+
 void ggml_flash_attn_ext_set_n_kv_max(
         struct ggml_tensor * a,
         int32_t              n_kv_max) {
diff --git a/include/llama.h b/include/llama.h
index 3ab93593..586fddd6 100644
--- a/include/llama.h
+++ b/include/llama.h
@@ -195,6 +195,14 @@ extern "C" {
 
     LLAMA_API const char * llama_flash_attn_type_name(enum llama_flash_attn_type flash_attn_type);
 
+    enum llama_pipeline_parallel_type {
+        LLAMA_PIPELINE_PARALLEL_TYPE_AUTO     = -1,
+        LLAMA_PIPELINE_PARALLEL_TYPE_DISABLED = 0,
+        LLAMA_PIPELINE_PARALLEL_TYPE_ENABLED  = 1,
+    };
+
+    LLAMA_API const char * llama_pipeline_parallel_type_name(enum llama_pipeline_parallel_type pipeline_parallel_type);
+
     enum llama_split_mode {
         LLAMA_SPLIT_MODE_NONE   = 0, // single GPU
         LLAMA_SPLIT_MODE_LAYER  = 1, // split layers and KV across GPUs
@@ -372,6 +380,7 @@ extern "C" {
         enum llama_pooling_type      pooling_type;      // whether to pool (sum) embedding results by sequence id
         enum llama_attention_type    attention_type;    // attention type to use for embeddings
         enum llama_flash_attn_type   flash_attn_type;   // when to enable Flash Attention
+        enum llama_pipeline_parallel_type pipeline_parallel_type; // when to enable pipeline parallelism
 
         // ref: https://github.com/ggml-org/llama.cpp/pull/2054
         float    rope_freq_base;   // RoPE base frequency, 0 = from model
@@ -416,6 +425,8 @@ extern "C" {
         // a source/target/parent context
         // can be utilized in various ways, for example by sharing results or llama_memory between 2 contexts
         struct llama_context * ctx_other;
+
+        bool hip_fa_force_vec; // force HIP quantized-KV FA onto VEC when supported
     };
 
     struct llama_model_tensor_override {
@@ -929,6 +940,11 @@ extern "C" {
                     llama_seq_id   seq_id,
            llama_state_seq_flags   flags);
 
+    // Allocate and pin all configured per-sequence device checkpoint buffers.
+    // Returns false when the context has no device checkpoint layout or an
+    // allocation fails.
+    LLAMA_API bool llama_state_seq_reserve_device_buffers(struct llama_context * ctx);
+
     LLAMA_API size_t llama_state_seq_set_data_ext(
             struct llama_context * ctx,
                    const uint8_t * src,
diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index ef53728d..3e6097c8 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -228,6 +228,7 @@ llama_context::llama_context(
 
     cparams.flash_attn = params.flash_attn_type != LLAMA_FLASH_ATTN_TYPE_DISABLED;
     cparams.auto_fa    = params.flash_attn_type == LLAMA_FLASH_ATTN_TYPE_AUTO;
+    cparams.hip_fa_force_vec = params.hip_fa_force_vec;
 
     cparams.fused_gdn_ar = true;
     cparams.fused_gdn_ch = true;
@@ -310,6 +311,8 @@ llama_context::llama_context(
     LLAMA_LOG_INFO("%s: n_ubatch              = %u\n",   __func__, cparams.n_ubatch);
     LLAMA_LOG_INFO("%s: causal_attn           = %d\n",   __func__, cparams.causal_attn);
     LLAMA_LOG_INFO("%s: flash_attn            = %s\n",   __func__, llama_flash_attn_type_name(params.flash_attn_type));
+    LLAMA_LOG_INFO("%s: hip_fa_force_vec      = %s\n",   __func__, cparams.hip_fa_force_vec ? "true" : "false");
+    LLAMA_LOG_INFO("%s: pipeline mode         = %s\n",   __func__, llama_pipeline_parallel_type_name(params.pipeline_parallel_type));
     LLAMA_LOG_INFO("%s: kv_unified            = %s\n",   __func__, cparams.kv_unified ? "true" : "false");
     LLAMA_LOG_INFO("%s: freq_base             = %.1f\n", __func__, cparams.rope_freq_base);
     LLAMA_LOG_INFO("%s: freq_scale            = %g\n",   __func__, cparams.rope_freq_scale);
@@ -426,6 +429,8 @@ llama_context::llama_context(
         // TODO: move these checks to ggml_backend_sched
         // enabling pipeline parallelism in the scheduler increases memory usage, so it is only done when necessary
         bool pipeline_parallel =
+            cparams.ctx_type != LLAMA_CONTEXT_TYPE_MTP &&
+            params.pipeline_parallel_type != LLAMA_PIPELINE_PARALLEL_TYPE_DISABLED &&
             model.n_devices() > 1 &&
             model.n_gpu_layers() > model.hparams.n_layer_all &&
             model.split_mode() == LLAMA_SPLIT_MODE_LAYER &&
@@ -452,6 +457,11 @@ llama_context::llama_context(
             }
         }
 
+        if (cparams.ctx_type != LLAMA_CONTEXT_TYPE_MTP &&
+            params.pipeline_parallel_type == LLAMA_PIPELINE_PARALLEL_TYPE_ENABLED && !pipeline_parallel) {
+            LLAMA_LOG_WARN("%s: pipeline parallelism requested but unavailable\n", __func__);
+        }
+
         cparams.pipeline_parallel = pipeline_parallel;
 
         if (cparams.pipeline_parallel) {
@@ -1815,6 +1825,9 @@ int llama_context::decode(const llama_batch & batch_inp) {
     int64_t n_outputs_prev = 0;
     int64_t n_tokens_prev  = 0;
 
+    bool has_next_ubatch = false;
+    bool mtp_multi_ubatch = false;
+
     do {
         const auto & ubatch = mctx->get_ubatch();
 
@@ -1990,7 +2003,15 @@ int llama_context::decode(const llama_batch & batch_inp) {
 
         n_outputs_prev += n_outputs;
         n_tokens_prev  += ubatch.n_tokens;
-    } while (mctx->next());
+
+        has_next_ubatch = mctx->next();
+        mtp_multi_ubatch |= has_next_ubatch;
+
+        // MTP ubatches update the same KV cache and must complete in order.
+        if (cparams.ctx_type == LLAMA_CONTEXT_TYPE_MTP && mtp_multi_ubatch) {
+            synchronize();
+        }
+    } while (has_next_ubatch);
 
     // set to total number of outputs in the batch, for use in llama_get_logits_ith
     n_outputs = n_outputs_all;
@@ -2760,10 +2781,12 @@ private:
 
 class llama_io_write_device : public llama_io_write_i {
 public:
-    llama_io_write_device(uint8_t * p, size_t len, llama_memory_buffers & mbufs) : ptr(p), buf_size(len), mbufs(mbufs)  {
+    llama_io_write_device(
+            uint8_t * p, size_t len, llama_memory_buffers & mbufs, bool copy_tensors = true) :
+        ptr(p), buf_size(len), mbufs(mbufs), copy_tensors(copy_tensors) {
     }
 
-    ~llama_io_write_device() {
+    void finish() {
         llama_memory_buffers mbufs_new;
 
         for (const auto & winfo : winfos) {
@@ -2825,11 +2848,21 @@ public:
 
             if (need_alloc) {
                 if (!mbuf_cur.buf || mbuf_cur.total_size != mbuf.total_size) {
-                    mbuf_cur = std::move(mbuf);
+                    ggml_backend_buffer_ptr buf {
+                        ggml_backend_alloc_ctx_tensors_from_buft(mbuf.ctx.get(), buft)
+                    };
+                    if (!buf) {
+                        throw std::runtime_error(
+                            std::string("failed to allocate device checkpoint buffer from '") +
+                            ggml_backend_buft_name(buft) + "'");
+                    }
 
-                    mbuf_cur.buf.reset(ggml_backend_alloc_ctx_tensors_from_buft(mbuf_cur.ctx.get(), buft));
+                    const size_t allocated_size = ggml_backend_buffer_get_size(buf.get());
+                    mbuf.buf = std::move(buf);
+                    mbuf_cur = std::move(mbuf);
 
-                    LLAMA_LOG_INFO("%s: allocated '%s' buffer %.3f MiB\n", __func__, ggml_backend_buft_name(buft), mbuf.total_size/1024.0/1024.0);
+                    LLAMA_LOG_INFO("%s: allocated '%s' buffer %.3f MiB\n", __func__,
+                            ggml_backend_buft_name(buft), allocated_size/1024.0/1024.0);
                 } else {
                     //LLAMA_LOG_INFO("%s: reallocating tensors in '%s' buffer %.3f MiB\n", __func__, ggml_backend_buft_name(buft), mbuf.total_size/1024.0/1024.0);
 
@@ -2849,8 +2882,10 @@ public:
                 }
             }
 
-            for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
-                ggml_backend_tensor_copy(mbuf_cur.org[i], mbuf_cur.cpy[i]);
+            if (copy_tensors) {
+                for (size_t i = 0; i < mbuf_cur.org.size(); ++i) {
+                    ggml_backend_tensor_copy(mbuf_cur.org[i], mbuf_cur.cpy[i]);
+                }
             }
         }
     }
@@ -2859,14 +2894,16 @@ public:
         if (size > buf_size) {
             throw std::runtime_error("unexpectedly reached end of buffer");
         }
-        memcpy(ptr, src, size);
-        ptr += size;
+        if (ptr != nullptr) {
+            memcpy(ptr, src, size);
+            ptr += size;
+        }
         size_written += size;
         buf_size -= size;
     }
 
     void write_tensor(ggml_tensor * tensor, size_t offset, size_t size) override {
-        // save the write for later during destruction
+        // save the write until finish(), after all state tensors are known
         winfos.push_back({tensor, ptr, size, offset});
     }
 
@@ -2888,6 +2925,7 @@ private:
     std::vector<write_info> winfos;
 
     llama_memory_buffers & mbufs;
+    const bool copy_tensors;
 };
 
 class llama_io_read_device : public llama_io_read_i {
@@ -3107,8 +3145,11 @@ size_t llama_context::state_seq_get_size(llama_seq_id seq_id, llama_state_seq_fl
 
 size_t llama_context::state_seq_get_data(llama_seq_id seq_id, uint8_t * dst, size_t size, llama_state_seq_flags flags) {
     std::unique_ptr<llama_io_write_i> io;
+    llama_io_write_device * io_device = nullptr;
     if (flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) {
-        io = std::make_unique<llama_io_write_device>(dst, size, mem_storage[seq_id]);
+        auto device = std::make_unique<llama_io_write_device>(dst, size, mem_storage[seq_id]);
+        io_device = device.get();
+        io = std::move(device);
     } else {
         io = std::make_unique<llama_io_write_host>(dst, size);
     }
@@ -3117,7 +3158,11 @@ size_t llama_context::state_seq_get_data(llama_seq_id seq_id, uint8_t * dst, siz
         io->write(&io_magic, sizeof(io_magic));
         io->write(&seq_id, sizeof(seq_id));
 
-        return state_seq_write_data(*io, seq_id, flags);
+        const size_t n = state_seq_write_data(*io, seq_id, flags);
+        if (io_device) {
+            io_device->finish();
+        }
+        return n;
     } catch (const std::exception & err) {
         LLAMA_LOG_ERROR("%s: error saving state: %s\n", __func__, err.what());
         return 0;
@@ -3410,6 +3455,30 @@ llama_memory_breakdown llama_context::memory_breakdown() const {
     return ret;
 }
 
+std::map<ggml_backend_buffer_type_t, size_t> llama_context::state_seq_device_buffer_sizes() const {
+    return memory ? memory->state_seq_device_buffer_sizes()
+                  : std::map<ggml_backend_buffer_type_t, size_t> {};
+}
+
+bool llama_context::state_seq_reserve_device_buffers() {
+    if (!memory || memory->state_seq_device_buffer_sizes().empty()) {
+        return false;
+    }
+
+    try {
+        for (llama_seq_id seq_id = 0; seq_id < static_cast<llama_seq_id>(cparams.n_seq_max); ++seq_id) {
+            llama_io_write_device io(nullptr, std::numeric_limits<size_t>::max(), mem_storage[seq_id], false);
+            memory->state_seq_write_device_layout(io);
+            io.finish();
+        }
+        return true;
+    } catch (const std::exception & err) {
+        mem_storage.clear();
+        LLAMA_LOG_ERROR("%s: failed to reserve device checkpoint buffers: %s\n", __func__, err.what());
+        return false;
+    }
+}
+
 //
 // training
 //
@@ -3659,6 +3728,7 @@ llama_context_params llama_context_default_params() {
         /*.pooling_type                =*/ LLAMA_POOLING_TYPE_UNSPECIFIED,
         /*.attention_type              =*/ LLAMA_ATTENTION_TYPE_UNSPECIFIED,
         /*.flash_attn_type             =*/ LLAMA_FLASH_ATTN_TYPE_AUTO,
+        /*.pipeline_parallel_type      =*/ LLAMA_PIPELINE_PARALLEL_TYPE_AUTO,
         /*.rope_freq_base              =*/ 0.0f,
         /*.rope_freq_scale             =*/ 0.0f,
         /*.yarn_ext_factor             =*/ -1.0f,
@@ -3682,6 +3752,7 @@ llama_context_params llama_context_default_params() {
         /*.sampler                     =*/ nullptr,
         /*.n_sampler                   =*/ 0,
         /*.ctx_other                   =*/ nullptr,
+        /*.hip_fa_force_vec            =*/ false,
     };
 
     return result;
@@ -4358,6 +4429,15 @@ llama_memory_breakdown llama_get_memory_breakdown(const struct llama_context * c
     return ctx->memory_breakdown();
 }
 
+std::map<ggml_backend_buffer_type_t, size_t> llama_get_state_seq_device_buffer_sizes(
+        const struct llama_context * ctx) {
+    return ctx->state_seq_device_buffer_sizes();
+}
+
+bool llama_state_seq_reserve_device_buffers(struct llama_context * ctx) {
+    return ctx->state_seq_reserve_device_buffers();
+}
+
 llama_context * llama_get_ctx_other(struct llama_context * ctx) {
     return ctx->get_cparams().ctx_other;
 }
diff --git a/src/llama-context.h b/src/llama-context.h
index b7a9db59..6c722194 100644
--- a/src/llama-context.h
+++ b/src/llama-context.h
@@ -189,6 +189,8 @@ struct llama_context {
     void perf_reset();
 
     llama_memory_breakdown memory_breakdown() const;
+    std::map<ggml_backend_buffer_type_t, size_t> state_seq_device_buffer_sizes() const;
+    bool state_seq_reserve_device_buffers();
 
     //
     // training
diff --git a/src/llama-cparams.h b/src/llama-cparams.h
index b592de18..ba833f9e 100644
--- a/src/llama-cparams.h
+++ b/src/llama-cparams.h
@@ -39,6 +39,7 @@ struct llama_cparams {
     bool offload_kqv;
     bool flash_attn;
     bool auto_fa;
+    bool hip_fa_force_vec;
     bool fused_gdn_ar;       // use fused gated delta net (autoregressive)
     bool fused_gdn_ch;       // use fused gated delta net (chunked)
     bool auto_fgdn;
diff --git a/src/llama-ext.h b/src/llama-ext.h
index 92a759b7..e011d6f3 100644
--- a/src/llama-ext.h
+++ b/src/llama-ext.h
@@ -90,6 +90,11 @@ LLAMA_API ggml_backend_dev_t llama_model_get_device(const struct llama_model * m
 
 LLAMA_API llama_memory_breakdown llama_get_memory_breakdown(const struct llama_context * ctx);
 
+// Exact backend-buffer sizes needed for one device-resident sequence-state
+// checkpoint per configured sequence.
+LLAMA_API std::map<ggml_backend_buffer_type_t, size_t> llama_get_state_seq_device_buffer_sizes(
+        const struct llama_context * ctx);
+
 // Set whether the context outputs nextn embeddings or not
 // If masked == true,  output the embeddings only for the tokens with batch.logits != 0
 // If masked == false, output the embeddings for all tokens in the batch regardless of batch.logits
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index fd4290cf..b208bfd6 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -2637,6 +2637,7 @@ ggml_tensor * llm_graph_context::build_attn_mha(
         GGML_ASSERT(n_kv_max >= 0 && n_kv_max <= INT32_MAX);
         ggml_flash_attn_ext_set_n_kv_max(cur, static_cast<int32_t>(n_kv_max));
         ggml_prec_set_acc(cur, GGML_PREC_F32);
+        ggml_flash_attn_ext_set_force_vec(cur, cparams.hip_fa_force_vec);
 
         if (v_mla) {
 #if 0
diff --git a/src/llama-memory-hybrid-iswa.cpp b/src/llama-memory-hybrid-iswa.cpp
index 06f7fd54..408982b7 100644
--- a/src/llama-memory-hybrid-iswa.cpp
+++ b/src/llama-memory-hybrid-iswa.cpp
@@ -192,6 +192,14 @@ std::map<ggml_backend_buffer_type_t, size_t> llama_memory_hybrid_iswa::memory_br
     return mb;
 }
 
+std::map<ggml_backend_buffer_type_t, size_t> llama_memory_hybrid_iswa::state_seq_device_buffer_sizes() const {
+    return mem_recr->state_seq_device_buffer_sizes();
+}
+
+void llama_memory_hybrid_iswa::state_seq_write_device_layout(llama_io_write_i & io) const {
+    mem_recr->state_seq_write_device_layout(io);
+}
+
 void llama_memory_hybrid_iswa::state_write(llama_io_write_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) const {
     mem_attn->state_write(io, seq_id, flags);
     mem_recr->state_write(io, seq_id, flags);
diff --git a/src/llama-memory-hybrid-iswa.h b/src/llama-memory-hybrid-iswa.h
index c9d3f9f5..02e5ce33 100644
--- a/src/llama-memory-hybrid-iswa.h
+++ b/src/llama-memory-hybrid-iswa.h
@@ -70,6 +70,8 @@ public:
     llama_pos seq_pos_max(llama_seq_id seq_id) const override;
 
     std::map<ggml_backend_buffer_type_t, size_t> memory_breakdown() const override;
+    std::map<ggml_backend_buffer_type_t, size_t> state_seq_device_buffer_sizes() const override;
+    void state_seq_write_device_layout(llama_io_write_i & io) const override;
 
     // state write/load
 
diff --git a/src/llama-memory-hybrid.cpp b/src/llama-memory-hybrid.cpp
index 42c7381a..b05cb1a9 100644
--- a/src/llama-memory-hybrid.cpp
+++ b/src/llama-memory-hybrid.cpp
@@ -187,6 +187,14 @@ std::map<ggml_backend_buffer_type_t, size_t> llama_memory_hybrid::memory_breakdo
     return mb;
 }
 
+std::map<ggml_backend_buffer_type_t, size_t> llama_memory_hybrid::state_seq_device_buffer_sizes() const {
+    return mem_recr->state_seq_device_buffer_sizes();
+}
+
+void llama_memory_hybrid::state_seq_write_device_layout(llama_io_write_i & io) const {
+    mem_recr->state_seq_write_device_layout(io);
+}
+
 void llama_memory_hybrid::state_write(llama_io_write_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) const {
     if ((flags & LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY) == 0) {
         mem_attn->state_write(io, seq_id, flags);
diff --git a/src/llama-memory-hybrid.h b/src/llama-memory-hybrid.h
index 484eafb7..5818fab4 100644
--- a/src/llama-memory-hybrid.h
+++ b/src/llama-memory-hybrid.h
@@ -70,6 +70,8 @@ public:
     llama_pos seq_pos_max(llama_seq_id seq_id) const override;
 
     std::map<ggml_backend_buffer_type_t, size_t> memory_breakdown() const override;
+    std::map<ggml_backend_buffer_type_t, size_t> state_seq_device_buffer_sizes() const override;
+    void state_seq_write_device_layout(llama_io_write_i & io) const override;
 
     // state write/load
 
diff --git a/src/llama-memory-recurrent.cpp b/src/llama-memory-recurrent.cpp
index 57919acc..b8db6875 100644
--- a/src/llama-memory-recurrent.cpp
+++ b/src/llama-memory-recurrent.cpp
@@ -1,5 +1,6 @@
 #include "llama-memory-recurrent.h"
 
+#include "ggml-alloc.h"
 #include "ggml-backend.h"
 #include "llama-impl.h"
 #include "llama-io.h"
@@ -425,6 +426,64 @@ std::map<ggml_backend_buffer_type_t, size_t> llama_memory_recurrent::memory_brea
     return ret;
 }
 
+std::map<ggml_backend_buffer_type_t, size_t> llama_memory_recurrent::state_seq_device_buffer_sizes() const {
+    struct layout {
+        size_t n_tensors = 0;
+        std::vector<std::pair<ggml_type, int64_t>> tensors;
+    };
+
+    std::map<ggml_backend_buffer_type_t, layout> layouts;
+    const uint32_t n_layer = hparams.n_layer();
+
+    auto add_row = [&](ggml_tensor * tensor, uint32_t n_embd) {
+        if (tensor == nullptr) {
+            return;
+        }
+        const uint64_t row_size = ggml_row_size(tensor->type, n_embd);
+        auto * buft = ggml_backend_buffer_get_type(tensor->buffer);
+        auto & cur = layouts[buft];
+        cur.n_tensors++;
+        cur.tensors.emplace_back(tensor->type, row_size / ggml_element_size(tensor));
+    };
+
+    for (uint32_t il = 0; il < n_layer; ++il) {
+        add_row(r_l[il], hparams.n_embd_r());
+    }
+    for (uint32_t il = 0; il < n_layer; ++il) {
+        add_row(s_l[il], hparams.n_embd_s());
+    }
+
+    std::map<ggml_backend_buffer_type_t, size_t> ret;
+    for (const auto & [buft, layout] : layouts) {
+        ggml_init_params params = {
+            /*.mem_size   =*/ layout.n_tensors * ggml_tensor_overhead(),
+            /*.mem_buffer =*/ nullptr,
+            /*.no_alloc   =*/ true,
+        };
+        ggml_context_ptr ctx { ggml_init(params) };
+        if (!ctx) {
+            throw std::runtime_error("failed to create recurrent checkpoint sizing context");
+        }
+        for (const auto & [type, n] : layout.tensors) {
+            ggml_new_tensor_1d(ctx.get(), type, n);
+        }
+
+        const size_t one_seq = ggml_backend_alloc_ctx_tensors_from_buft_size(ctx.get(), buft);
+        if (n_seq_max != 0 && one_seq > std::numeric_limits<size_t>::max() / n_seq_max) {
+            throw std::runtime_error("recurrent checkpoint size overflow");
+        }
+        ret[buft] = one_seq * n_seq_max;
+    }
+    return ret;
+}
+
+void llama_memory_recurrent::state_seq_write_device_layout(llama_io_write_i & io) const {
+    // One logical recurrent row is the complete per-sequence checkpoint.  The
+    // actual save may select a rollback plane, but its tensor shapes and buffer
+    // sizes are identical to row zero.
+    state_write_data(io, {{0, 1}});
+}
+
 llama_memory_context_ptr llama_memory_recurrent::init_batch(llama_batch_allocr & balloc, uint32_t n_ubatch, bool embd_all) {
     do {
         balloc.split_reset();
@@ -822,7 +881,7 @@ void llama_memory_recurrent::state_write(llama_io_write_i & io, llama_seq_id seq
     }
 
     if ((flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) && cell_ranges.size() > 1) {
-        GGML_ABORT("cannot save/load multiple ranges of cells to/from device memory\n");
+        throw std::runtime_error("cannot save/load multiple ranges of cells to/from device memory");
     }
 
     // DEBUG CHECK: Sum of cell counts in ranges should equal the total cell count
diff --git a/src/llama-memory-recurrent.h b/src/llama-memory-recurrent.h
index 4abb3f5c..945e179d 100644
--- a/src/llama-memory-recurrent.h
+++ b/src/llama-memory-recurrent.h
@@ -53,6 +53,8 @@ public:
     llama_pos seq_pos_max(llama_seq_id seq_id) const override;
 
     std::map<ggml_backend_buffer_type_t, size_t> memory_breakdown() const override;
+    std::map<ggml_backend_buffer_type_t, size_t> state_seq_device_buffer_sizes() const override;
+    void state_seq_write_device_layout(llama_io_write_i & io) const override;
 
     bool prepare(const std::vector<llama_ubatch> & ubatches);
 
diff --git a/src/llama-memory.h b/src/llama-memory.h
index db825396..374d49d4 100644
--- a/src/llama-memory.h
+++ b/src/llama-memory.h
@@ -118,6 +118,15 @@ struct llama_memory_i {
 
     virtual std::map<ggml_backend_buffer_type_t, size_t> memory_breakdown() const = 0;
 
+    // Exact backend-buffer sizes needed to keep one sequence-state checkpoint
+    // per configured sequence. Memory types without device checkpoint data
+    // return an empty map.
+    virtual std::map<ggml_backend_buffer_type_t, size_t> state_seq_device_buffer_sizes() const { return {}; }
+
+    // Describe the tensor data for one device-resident sequence checkpoint.
+    // This is used to allocate the real checkpoint buffers before evaluation.
+    virtual void state_seq_write_device_layout(llama_io_write_i & io) const { GGML_UNUSED(io); }
+
     //
     // state write/read
     //
diff --git a/src/llama.cpp b/src/llama.cpp
index ad8e4438..e51fd108 100644
--- a/src/llama.cpp
+++ b/src/llama.cpp
@@ -47,6 +47,18 @@ const char * llama_flash_attn_type_name(enum llama_flash_attn_type flash_attn_ty
     GGML_ABORT("fatal error");
 }
 
+const char * llama_pipeline_parallel_type_name(enum llama_pipeline_parallel_type pipeline_parallel_type) {
+    switch (pipeline_parallel_type) {
+        case LLAMA_PIPELINE_PARALLEL_TYPE_AUTO:
+            return "auto";
+        case LLAMA_PIPELINE_PARALLEL_TYPE_DISABLED:
+            return "disabled";
+        case LLAMA_PIPELINE_PARALLEL_TYPE_ENABLED:
+            return "enabled";
+    }
+    GGML_ABORT("fatal error");
+}
+
 const char * llama_load_mode_name(enum llama_load_mode load_mode) {
     switch (load_mode) {
         case LLAMA_LOAD_MODE_AUTO:
@@ -617,4 +629,3 @@ const char * llama_print_system_info(void) {
 
     return s.c_str();
 }
-
diff --git a/src/models/minimax-m3.cpp b/src/models/minimax-m3.cpp
index f3b64b21..b006a08a 100644
--- a/src/models/minimax-m3.cpp
+++ b/src/models/minimax-m3.cpp
@@ -192,6 +192,7 @@ ggml_tensor * llama_model_minimax_m3::graph::build_attn_msa_fa(
     ggml_tensor * o = ggml_flash_attn_ext(ctx0, q, k, v, mask, kq_scale,
                                           hparams.f_max_alibi_bias, 0.0f);
     ggml_prec_set_acc(o, GGML_PREC_F32);
+    ggml_flash_attn_ext_set_force_vec(o, cparams.hip_fa_force_vec);
     cb(o, "msa_fattn", il);
 
     // [D, Gp, R, C] -> [D, Gp, C, R] -> [n_embd, T]
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index a398344c..8f35b629 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -194,6 +194,13 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
 
     llama_build(test-llama-archs.cpp)
 
+    llama_test(
+        test-llama-archs
+        NAME test-mtp-ubatch-sync
+        LABEL main
+        ARGS --test-mtp-ubatch-sync
+    )
+
     set(MODEL_DIR "${CMAKE_CURRENT_BINARY_DIR}/test-models/")
     file(MAKE_DIRECTORY "${MODEL_DIR}")
 
diff --git a/tests/test-arg-parser.cpp b/tests/test-arg-parser.cpp
index e0907631..8ac51a97 100644
--- a/tests/test-arg-parser.cpp
+++ b/tests/test-arg-parser.cpp
@@ -3,6 +3,7 @@
 #include "download.h"
 #include "llama.h"
 #include "speculative.h"
+#include "ggml-backend.h"
 
 #include <cmath>
 #include <limits>
@@ -186,6 +187,23 @@ static void test(void) {
     argv = {"binary_name", "-sm", "hello"};
     assert(false == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_COMMON));
 
+    {
+        common_params hip_fa_params;
+        assert(!hip_fa_params.hip_fa_force_vec);
+        assert(!llama_context_default_params().hip_fa_force_vec);
+
+        argv = {"binary_name", "-m", "model.gguf", "--hip-fa-force-vec", "off"};
+        assert(common_params_parse(argv.size(), list_str_to_char(argv).data(), hip_fa_params, LLAMA_EXAMPLE_COMMON));
+        assert(!hip_fa_params.hip_fa_force_vec);
+
+        argv = {"binary_name", "-m", "model.gguf", "--hip-fa-force-vec", "on"};
+        assert(common_params_parse(argv.size(), list_str_to_char(argv).data(), hip_fa_params, LLAMA_EXAMPLE_COMMON));
+        assert(hip_fa_params.hip_fa_force_vec);
+
+        argv = {"binary_name", "-m", "model.gguf", "--hip-fa-force-vec", "maybe"};
+        assert(!common_params_parse(argv.size(), list_str_to_char(argv).data(), hip_fa_params, LLAMA_EXAMPLE_COMMON));
+    }
+
     {
         common_params penalty_params;
         assert(penalty_params.sampling.penalty_last_n == 64);
@@ -254,6 +272,136 @@ static void test(void) {
     assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_SPECULATIVE));
     assert(params.speculative.draft.n_max == 123);
 
+    argv = {"binary_name", "--spec-draft-adaptive"};
+    assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params, LLAMA_EXAMPLE_SPECULATIVE));
+    assert(params.speculative.draft.adaptive);
+
+    {
+        common_params params_mtp;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-draft-n-max", "5", "--spec-mtp-cr-depth", "1"};
+        assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params_mtp, LLAMA_EXAMPLE_SPECULATIVE));
+        assert(params_mtp.speculative.draft.n_max == 5);
+        assert(params_mtp.speculative.draft.n_rs_seq == 1);
+
+        params_mtp.speculative.types = { COMMON_SPECULATIVE_TYPE_DRAFT_MTP };
+        assert(params_mtp.speculative.need_n_rs_seq() == 1);
+
+        params_mtp.speculative.types.push_back(COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3);
+        assert(params_mtp.speculative.need_n_rs_seq() == 5);
+    }
+
+    {
+        common_params params_mtp;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-draft-n-max", "5"};
+        assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params_mtp, LLAMA_EXAMPLE_SPECULATIVE));
+        assert(params_mtp.speculative.draft.n_rs_seq == -1);
+
+        params_mtp.speculative.types = { COMMON_SPECULATIVE_TYPE_DRAFT_MTP };
+        assert(params_mtp.speculative.need_n_rs_seq() == 5);
+    }
+
+    {
+        common_params params_mtp;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-mtp-cr-depth", "1", "--spec-draft-n-max", "5"};
+        assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), params_mtp, LLAMA_EXAMPLE_SPECULATIVE));
+        assert(params_mtp.speculative.draft.n_rs_seq == 1);
+    }
+
+    {
+        common_params params_mtp;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-draft-n-max", "5", "--spec-mtp-cr-depth", "0"};
+        assert(false == common_params_parse(argv.size(), list_str_to_char(argv).data(), params_mtp, LLAMA_EXAMPLE_SPECULATIVE));
+    }
+
+    {
+        common_params params_mtp;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-draft-n-max", "5", "--spec-mtp-cr-depth=-1"};
+        assert(false == common_params_parse(argv.size(), list_str_to_char(argv).data(), params_mtp, LLAMA_EXAMPLE_SPECULATIVE));
+    }
+
+    {
+        common_params params_mtp;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-mtp-cr-depth", "6", "--spec-draft-n-max", "5"};
+        assert(false == common_params_parse(argv.size(), list_str_to_char(argv).data(), params_mtp, LLAMA_EXAMPLE_SPECULATIVE));
+    }
+
+    {
+        printf("test-arg-parser: test recurrent checkpoint reservation inputs\n\n");
+
+        common_params fit_params;
+        assert(fit_params.fit_params);
+        assert(fit_params.fit_params_target.size() == llama_max_devices());
+        for (const size_t target : fit_params.fit_params_target) {
+            assert(target == 1024 * 1024 * 1024ULL);
+        }
+
+        argv = {"binary_name", "-m", "abc.gguf", "--fit", "off", "--split-mode", "none", "--main-gpu", "1",
+                "--fit-target", "256,64", "--tensor-split", "3,1"};
+        assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), fit_params, LLAMA_EXAMPLE_COMMON));
+        assert(!fit_params.fit_params);
+        assert(fit_params.split_mode == LLAMA_SPLIT_MODE_NONE);
+        assert(fit_params.main_gpu == 1);
+        assert(fit_params.fit_params_target[0] == 256 * 1024 * 1024ULL);
+        assert(fit_params.fit_params_target[1] == 64 * 1024 * 1024ULL);
+        assert(fit_params.tensor_split[0] == 3.0f);
+        assert(fit_params.tensor_split[1] == 1.0f);
+
+        auto checkpoint_reservation_needed = [](const common_params & p) {
+            return p.speculative.need_n_rs_seq() < static_cast<uint32_t>(p.speculative.draft.n_max);
+        };
+
+        common_params mtp_full;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-draft-n-max", "5"};
+        assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), mtp_full, LLAMA_EXAMPLE_SPECULATIVE));
+        mtp_full.speculative.types = { COMMON_SPECULATIVE_TYPE_DRAFT_MTP };
+        assert(mtp_full.speculative.need_n_rs_seq() == 5);
+        assert(!checkpoint_reservation_needed(mtp_full));
+
+        common_params mtp_reduced;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-draft-n-max", "5", "--spec-mtp-cr-depth", "2"};
+        assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), mtp_reduced, LLAMA_EXAMPLE_SPECULATIVE));
+        mtp_reduced.speculative.types = { COMMON_SPECULATIVE_TYPE_DRAFT_MTP };
+        assert(mtp_reduced.speculative.need_n_rs_seq() == 2);
+        assert(checkpoint_reservation_needed(mtp_reduced));
+
+        common_params mtp_with_other_rs;
+        argv = {"binary_name", "-m", "abc.gguf", "--spec-draft-n-max", "5", "--spec-mtp-cr-depth", "2"};
+        assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), mtp_with_other_rs, LLAMA_EXAMPLE_SPECULATIVE));
+        mtp_with_other_rs.speculative.types = {
+            COMMON_SPECULATIVE_TYPE_DRAFT_MTP,
+            COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3,
+        };
+        assert(mtp_with_other_rs.speculative.need_n_rs_seq() == 5);
+        assert(!checkpoint_reservation_needed(mtp_with_other_rs));
+
+        ggml_backend_load_all();
+        std::vector<ggml_backend_dev_t> gpu_devices;
+        for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
+            ggml_backend_dev_t dev = ggml_backend_dev_get(i);
+            if (ggml_backend_dev_type(dev) != GGML_BACKEND_DEVICE_TYPE_CPU) {
+                gpu_devices.push_back(dev);
+            }
+        }
+        if (gpu_devices.size() >= 2) {
+            const std::string device_arg = string_format("%s,%s",
+                    ggml_backend_dev_name(gpu_devices[0]), ggml_backend_dev_name(gpu_devices[1]));
+            common_params multi_gpu;
+            argv = {"binary_name", "-m", "abc.gguf", "--device", device_arg, "--split-mode", "layer",
+                    "--fit-target", "128,256", "--tensor-split", "3,1"};
+            assert(true == common_params_parse(argv.size(), list_str_to_char(argv).data(), multi_gpu, LLAMA_EXAMPLE_COMMON));
+            assert(multi_gpu.devices.size() == 3);
+            assert(multi_gpu.devices[0] == gpu_devices[0]);
+            assert(multi_gpu.devices[1] == gpu_devices[1]);
+            assert(multi_gpu.devices[2] == nullptr);
+            assert(multi_gpu.fit_params_target[0] == 128 * 1024 * 1024ULL);
+            assert(multi_gpu.fit_params_target[1] == 256 * 1024 * 1024ULL);
+            assert(multi_gpu.tensor_split[0] == 3.0f);
+            assert(multi_gpu.tensor_split[1] == 1.0f);
+        } else {
+            printf("test-arg-parser: skip asymmetric multi-GPU input test (fewer than two non-CPU devices)\n\n");
+        }
+    }
+
     {
         common_params synth_params;
         argv = {"binary_name", "--spec-synth-len", "3.4"};
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index dc529a35..5c9d01dd 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -11170,9 +11170,33 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
         }
     }
 
+    // Qwen3.6-27B decode shapes (n=1, mmvq path), Q6_K vs Q8_0:
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 10240, 1, 5120, {1, 1}, {1, 1})); // ffn
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 248320, 1, 5120, {1, 1}, {1, 1})); // lm_head
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 12288, 1, 5120, {1, 1}, {1, 1})); // qkv
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 6144, 1, 5120, {1, 1}, {1, 1})); // ssm z
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q6_K, GGML_TYPE_F32, 1024, 1, 5120, {1, 1}, {1, 1})); // kv proj
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 10240, 1, 5120, {1, 1}, {1, 1})); // ffn
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 248320, 1, 5120, {1, 1}, {1, 1})); // lm_head
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 12288, 1, 5120, {1, 1}, {1, 1})); // qkv
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 6144, 1, 5120, {1, 1}, {1, 1})); // ssm z
+    test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 1024, 1, 5120, {1, 1}, {1, 1})); // kv proj
+
+    // same shapes, Q4_K/Q5_K decode (n=1, mmvq) and prefill (mmq) rows:
+    for (ggml_type type_a : {GGML_TYPE_Q4_K, GGML_TYPE_Q5_K}) {
+        test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 10240, 1, 5120, {1, 1}, {1, 1})); // ffn
+        test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 248320, 1, 5120, {1, 1}, {1, 1})); // lm_head
+        test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 12288, 1, 5120, {1, 1}, {1, 1})); // qkv
+        test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 6144, 1, 5120, {1, 1}, {1, 1})); // ssm z
+        test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 1024, 1, 5120, {1, 1}, {1, 1})); // kv proj
+        for (int bs : {16, 128, 512}) {
+            test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 10240, bs, 5120, {1, 1}, {1, 1})); // prefill
+        }
+    }
+
     // qwen3-30b-a3b
-    for (int bs : {1, 4, 8, 32, 64, 128, 256, 512}) {
-        for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) {
+    for (int bs : {1, 4, 5, 6, 7, 8, 32, 64, 128, 256, 512}) {
+        for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) {
             for (ggml_type type_b : {GGML_TYPE_F32}) {
                 test_cases.emplace_back(new test_mul_mat_id(type_a, type_b, 128, 8, false, 768, bs, 2048));
                 test_cases.emplace_back(new test_mul_mat_id_fusion(type_a, type_b, 128, 8, false, 768, bs, 2048, 1));
@@ -11180,8 +11204,8 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
         }
     }
 
-    for (int bs : {1, 4, 8, 32, 64, 128, 256, 512}) {
-        for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) {
+    for (int bs : {1, 4, 5, 6, 7, 8, 32, 64, 128, 256, 512}) {
+        for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) {
             for (ggml_type type_b : {GGML_TYPE_F32}) {
                 test_cases.emplace_back(new test_mul_mat_id(type_a, type_b, 32, 4, false, 1792, bs, 2048));
                 test_cases.emplace_back(new test_mul_mat_id_fusion(type_a, type_b, 32, 4, false, 1792, bs, 2048, 1));
diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp
index 568f7234..ed617bf7 100644
--- a/tests/test-llama-archs.cpp
+++ b/tests/test-llama-archs.cpp
@@ -65,7 +65,7 @@ static void set_tensor_data(struct ggml_tensor * tensor, void * userdata) {
 }
 
 static void usage(char ** argv) {
-    printf("Usage: %s [-a/--arch arch] [-s/--seed seed] [-o/--out dir] [-v N] [-h/--help]\n", argv[0]);
+    printf("Usage: %s [-a/--arch arch] [-s/--seed seed] [-o/--out dir] [-v N] [-h/--help] [--test-mtp-ubatch-sync]\n", argv[0]);
 }
 
 static std::vector<llama_token> get_tokens(const uint32_t n_tokens, const uint32_t n_vocab, const size_t seed){
@@ -79,7 +79,7 @@ static std::vector<llama_token> get_tokens(const uint32_t n_tokens, const uint32
     return ret;
 }
 
-static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
+static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe, const bool mtp = false) {
     gguf_context_ptr ret(gguf_init_empty());
     llama_model_saver ms(arch, ret.get());
     const uint32_t n_ctx = 256;
@@ -149,6 +149,9 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
     ms.add_kv(LLM_KV_FEATURES_LENGTH,           n_embd);
     ms.add_kv(LLM_KV_BLOCK_COUNT,               n_layer);
     ms.add_kv(LLM_KV_LEADING_DENSE_BLOCK_COUNT, uint32_t(1));
+    if (mtp) {
+        ms.add_kv(LLM_KV_NEXTN_PREDICT_LAYERS, uint32_t(1));
+    }
 
     if (arch == LLM_ARCH_NEMOTRON_H || arch == LLM_ARCH_NEMOTRON_H_MOE) {
         std::vector<uint32_t> n_ff_per_layer;
@@ -436,6 +439,83 @@ static std::pair<llama_model_ptr, llama_context_ptr> get_model_and_ctx(
     return std::make_pair(std::move(model), std::move(lctx));
 }
 
+static bool mtp_sync_test_decode(llama_model * model, uint32_t n_ubatch) {
+    const int32_t n_tokens = 4;
+    const int32_t n_embd = llama_model_n_embd_out(model);
+
+    llama_context_params ctx_params = llama_context_default_params();
+    ctx_params.n_ctx = 8;
+    ctx_params.n_batch = n_tokens;
+    ctx_params.n_ubatch = n_ubatch;
+    ctx_params.n_threads = 4;
+    ctx_params.n_threads_batch = 4;
+    ctx_params.ctx_type = LLAMA_CONTEXT_TYPE_MTP;
+
+    llama_context_ptr ctx(llama_init_from_model(model, ctx_params));
+    if (!ctx) {
+        throw std::runtime_error("failed to create MTP context");
+    }
+
+    std::vector<llama_token> token(n_tokens);
+    std::vector<float> embd((size_t) n_tokens * n_embd, 1.0e-2f);
+    std::vector<llama_pos> pos(n_tokens);
+    std::vector<int32_t> n_seq_id(n_tokens, 1);
+    std::vector<llama_seq_id> seq_id_data(n_tokens, 0);
+    std::vector<llama_seq_id *> seq_id(n_tokens);
+    std::vector<int8_t> logits(n_tokens, 0);
+
+    for (int32_t i = 0; i < n_tokens; ++i) {
+        token[i] = i;
+        pos[i] = i;
+        seq_id[i] = &seq_id_data[i];
+    }
+    logits.back() = 1;
+
+    llama_batch batch = {
+        /*.n_tokens =*/ n_tokens,
+        /*.token =*/ token.data(),
+        /*.embd =*/ embd.data(),
+        /*.pos =*/ pos.data(),
+        /*.n_seq_id =*/ n_seq_id.data(),
+        /*.seq_id =*/ seq_id.data(),
+        /*.logits =*/ logits.data(),
+    };
+
+    llama_perf_context_reset(ctx.get());
+    const int32_t ret = llama_decode(ctx.get(), batch);
+    if (ret != 0) {
+        throw std::runtime_error("failed to decode MTP batch");
+    }
+
+    // synchronize() accounts the complete queued batch immediately. Without a
+    // synchronization, llama_perf_context() reports its minimum placeholder of 1.
+    return llama_perf_context(ctx.get()).n_p_eval >= n_tokens;
+}
+
+static int test_mtp_ubatch_sync(const size_t seed) {
+    gguf_context_ptr gguf_ctx = get_gguf_ctx(LLM_ARCH_QWEN35, false, true);
+    llama_model_params model_params = llama_model_default_params();
+    model_params.progress_callback = silent_model_load_progress;
+    model_params.load_mtp = true;
+
+    size_t tmp = seed;
+    llama_model_ptr model(llama_model_init_from_user(gguf_ctx.get(), set_tensor_data, &tmp, model_params));
+    if (!model) {
+        throw std::runtime_error("failed to create MTP model");
+    }
+
+    if (!mtp_sync_test_decode(model.get(), 2)) {
+        fprintf(stderr, "MTP ubatches were not synchronized\n");
+        return 1;
+    }
+    if (mtp_sync_test_decode(model.get(), 4)) {
+        fprintf(stderr, "single MTP ubatch was synchronized\n");
+        return 1;
+    }
+
+    return 0;
+}
+
 static std::vector<float> get_logits(
         llama_model * model, llama_context * lctx, const std::vector<llama_token> & tokens, bool encode = false) {
     const uint32_t n_vocab  = llama_vocab_n_tokens(llama_model_get_vocab(model));
@@ -835,6 +915,7 @@ int main(int argc, char ** argv) {
     llm_arch arch = LLM_ARCH_UNKNOWN;
     size_t seed = rd();
     std::string out;
+    bool test_mtp_sync = false;
 
     int verbosity = LOG_LEVEL_ERROR;
 
@@ -880,6 +961,9 @@ int main(int argc, char ** argv) {
                 return 1;
             }
         }
+        if (strcmp(argv[i], "--test-mtp-ubatch-sync") == 0) {
+            test_mtp_sync = true;
+        }
     }
     printf("%s: using seed %zu\n", __func__, seed);
 
@@ -887,6 +971,9 @@ int main(int argc, char ** argv) {
         if (!out.empty()) {
             return save_models(arch, seed, verbosity, out);
         }
+        if (test_mtp_sync) {
+            return test_mtp_ubatch_sync(seed);
+        }
         return test_backends(arch, seed, verbosity);
     } catch (const std::exception & err) {
         fprintf(stderr, "encountered runtime error: %s\n", err.what());
diff --git a/tests/test-recurrent-state-rollback.cpp b/tests/test-recurrent-state-rollback.cpp
index ef05de67..7a514d63 100644
--- a/tests/test-recurrent-state-rollback.cpp
+++ b/tests/test-recurrent-state-rollback.cpp
@@ -86,6 +86,19 @@ static llama_context * make_ctx(const common_params & params, llama_model * mode
     return init_ctx(model, cparams, fill);
 }
 
+static llama_context * make_ctx_rs(
+        const common_params & params,
+        llama_model * model,
+        uint32_t n_rs_seq,
+        uint32_t n_seq_max = 1) {
+    auto cparams = common_context_params_to_llama(params);
+    cparams.n_seq_max = n_seq_max;
+    cparams.n_rs_seq  = n_rs_seq;
+    cparams.n_batch   = std::max(cparams.n_batch,  (uint32_t) (n_rs_seq + 1));
+    cparams.n_ubatch  = std::max(cparams.n_ubatch, (uint32_t) (n_rs_seq + 1));
+    return llama_init_from_model(model, cparams);
+}
+
 static float logit_diff(float a, float b) {
     return std::isfinite(a) && std::isfinite(b) ? std::fabs(a - b) : std::numeric_limits<float>::infinity();
 }
@@ -262,6 +275,249 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode
     return true;
 }
 
+static bool test_fragmented_on_device_fallback(
+        const common_params & params,
+        llama_model * model,
+        const std::vector<llama_token> & tokens) {
+    llama_context * ctx = make_ctx_rs(params, model, 1, 3);
+    if (ctx == nullptr) {
+        fprintf(stderr, "%s : failed to initialize fragmented context\n", __func__);
+        return false;
+    }
+
+    llama_batch batch = llama_batch_init(3, 0, 3);
+    for (llama_seq_id seq_id = 0; seq_id < 3; ++seq_id) {
+        common_batch_add(batch, tokens[seq_id % tokens.size()], 0, { seq_id }, seq_id == 2);
+    }
+    bool ok = llama_decode(ctx, batch) == 0;
+    llama_batch_free(batch);
+
+    if (ok) {
+        ok = llama_memory_seq_rm(llama_get_memory(ctx), 1, -1, -1);
+    }
+
+    common_prompt_checkpoint ckpt;
+    constexpr llama_state_seq_flags flags =
+        LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+    if (ok) {
+        ckpt.update_tgt(ctx, -1, flags);
+        if (ckpt.data_tgt_on_device) {
+            // Some recurrent layouts compact the surviving sequences into one
+            // device range.  That is valid and does not require host fallback.
+            fprintf(stderr, "%s : checkpoint remained device-contiguous\n", __func__);
+        } else {
+            fprintf(stderr, "%s : fragmented checkpoint fell back to host storage\n", __func__);
+        }
+    }
+    if (ok) {
+        // load_tgt() must use the actual storage mode recorded when the
+        // checkpoint was saved, whether the layout stayed contiguous or fell
+        // back to host storage.
+        ckpt.load_tgt(ctx, -1, flags);
+    }
+
+    llama_free(ctx);
+    return ok;
+}
+
+static bool test_preallocated_device_checkpoints(
+        const common_params & params,
+        llama_model * model,
+        const std::vector<llama_token> & tokens) {
+    constexpr int32_t n_seq = 2;
+    llama_context * ctx = make_ctx_rs(params, model, 1, n_seq);
+    if (ctx == nullptr) {
+        fprintf(stderr, "%s : failed to initialize multi-sequence context\n", __func__);
+        return false;
+    }
+
+    bool ok = llama_state_seq_reserve_device_buffers(ctx);
+    if (!ok) {
+        // Architectures without recurrent tensor rows have no device layout to
+        // reserve.  The server handles this by using host checkpoints.
+        fprintf(stderr, "%s : no device checkpoint layout; using host checkpoints\n", __func__);
+        llama_free(ctx);
+        return true;
+    }
+
+    llama_batch batch = llama_batch_init(n_seq, 0, n_seq);
+    for (llama_seq_id seq_id = 0; seq_id < n_seq; ++seq_id) {
+        common_batch_add(batch, tokens[seq_id % tokens.size()], 0, { seq_id }, true);
+    }
+    if (ok) {
+        ok = llama_decode(ctx, batch) == 0;
+    }
+    llama_batch_free(batch);
+
+    constexpr llama_state_seq_flags flags =
+        LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+    std::vector<common_prompt_checkpoint> checkpoints(n_seq);
+    for (llama_seq_id seq_id = 0; ok && seq_id < n_seq; ++seq_id) {
+        checkpoints[seq_id].update_tgt(ctx, seq_id, flags);
+        if (!checkpoints[seq_id].data_tgt_on_device) {
+            fprintf(stderr, "%s : sequence %d did not reuse its preallocated device checkpoint\n",
+                    __func__, seq_id);
+            ok = false;
+        }
+    }
+    for (llama_seq_id seq_id = 0; ok && seq_id < n_seq; ++seq_id) {
+        checkpoints[seq_id].load_tgt(ctx, seq_id, flags);
+    }
+
+    llama_free(ctx);
+    return ok;
+}
+
+static bool decode_range(
+        llama_context * ctx,
+        const std::vector<llama_token> & tokens,
+        uint32_t first,
+        uint32_t count) {
+    if (count == 0) {
+        return true;
+    }
+
+    llama_batch batch = llama_batch_init(count, 0, 1);
+    for (uint32_t i = 0; i < count; ++i) {
+        common_batch_add(batch, tokens[first + i], first + i, { 0 }, i + 1 == count);
+    }
+    const bool ok = llama_decode(ctx, batch) == 0;
+    llama_batch_free(batch);
+    return ok;
+}
+
+static bool compare_logits(
+        llama_context * ctx_full,
+        llama_context * ctx_depth,
+        int n_vocab,
+        uint32_t n_accepted,
+        const char * stage,
+        float eps) {
+    const float * logits_full  = llama_get_logits(ctx_full);
+    const float * logits_depth = llama_get_logits(ctx_depth);
+    if (logits_full == nullptr || logits_depth == nullptr) {
+        fprintf(stderr, "%s : missing %s logits after accepting %u draft tokens\n", __func__, stage, n_accepted);
+        return false;
+    }
+
+    int argmax_full  = 0;
+    int argmax_depth = 0;
+    float max_diff = 0.0f;
+    for (int token = 0; token < n_vocab; ++token) {
+        if (logits_full[token] > logits_full[argmax_full]) {
+            argmax_full = token;
+        }
+        if (logits_depth[token] > logits_depth[argmax_depth]) {
+            argmax_depth = token;
+        }
+
+        const float diff = std::fabs(logits_full[token] - logits_depth[token]);
+        max_diff = std::max(max_diff, diff);
+        if (eps >= 0.0f && diff > eps) {
+            fprintf(stderr, "%s : %s logits mismatch after accepting %u draft tokens, token %d (%g != %g)\n",
+                    __func__, stage, n_accepted, token, (double) logits_full[token], (double) logits_depth[token]);
+            return false;
+        }
+    }
+    if (eps < 0.0f && argmax_full != argmax_depth) {
+        fprintf(stderr, "%s : %s greedy token mismatch after accepting %u draft tokens (%d != %d, max logit diff %g)\n",
+                __func__, stage, n_accepted, argmax_full, argmax_depth, (double) max_diff);
+        return false;
+    }
+    if (eps < 0.0f) {
+        fprintf(stderr, "%s : %s greedy token %d preserved after accepting %u draft tokens (max logit diff %g)\n",
+                __func__, stage, argmax_full, n_accepted, (double) max_diff);
+    }
+    return true;
+}
+
+static bool test_recompute_fallback(
+        const common_params & params,
+        llama_model * model,
+        const std::vector<llama_token> & input_tokens,
+        int n_vocab) {
+    constexpr uint32_t n_prefix = 4;
+    constexpr uint32_t n_draft  = 5;
+    constexpr uint32_t n_verify = n_draft + 1;
+
+    std::vector<llama_token> tokens = input_tokens;
+    tokens.resize(n_prefix + n_verify + 2, tokens.back());
+
+    for (uint32_t n_accepted = 0; n_accepted < n_draft; ++n_accepted) {
+        llama_context * ctx_full  = make_ctx_rs(params, model, n_draft);
+        llama_context * ctx_depth = make_ctx_rs(params, model, 1);
+        if (ctx_full == nullptr || ctx_depth == nullptr) {
+            fprintf(stderr, "%s : failed to initialize fallback contexts\n", __func__);
+            llama_free(ctx_full);
+            llama_free(ctx_depth);
+            return false;
+        }
+
+        bool ok = decode_range(ctx_full, tokens, 0, n_prefix) &&
+                  decode_range(ctx_depth, tokens, 0, n_prefix);
+        if (ok) {
+            ok = compare_logits(ctx_full, ctx_depth, n_vocab, n_accepted, "prefix", 1e-5f);
+        }
+
+        constexpr llama_state_seq_flags partial_flags =
+            LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY | LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+        common_prompt_checkpoint ckpt;
+        if (ok) {
+            ckpt.update_tgt(ctx_depth, 0, partial_flags);
+            if (!ckpt.data_tgt_on_device) {
+                fprintf(stderr, "%s : contiguous recurrent checkpoint unexpectedly fell back to host storage\n", __func__);
+                ok = false;
+            }
+        }
+        if (ok) {
+            ok = decode_range(ctx_full, tokens, n_prefix, n_verify) &&
+                 decode_range(ctx_depth, tokens, n_prefix, n_verify);
+        }
+
+        const llama_pos rollback_pos = n_prefix + 1 + n_accepted;
+        const uint32_t n_rollback = n_draft - n_accepted;
+        if (ok) {
+            ok = llama_memory_seq_rm(llama_get_memory(ctx_full), 0, rollback_pos, -1);
+        }
+        if (ok && n_rollback <= 1) {
+            ok = llama_memory_seq_rm(llama_get_memory(ctx_depth), 0, rollback_pos, -1);
+        } else if (ok) {
+            ckpt.load_tgt(ctx_depth, 0, partial_flags);
+            ok = llama_memory_seq_rm(llama_get_memory(ctx_depth), 0, n_prefix, -1);
+        }
+
+        const llama_pos replacement_pos = rollback_pos;
+        if (ok && n_rollback <= 1) {
+            ok = decode_one(ctx_full, tokens[n_prefix + n_verify], replacement_pos) &&
+                 decode_one(ctx_depth, tokens[n_prefix + n_verify], replacement_pos);
+        } else if (ok) {
+            std::vector<llama_token> replay_tokens = tokens;
+            replay_tokens[replacement_pos] = tokens[n_prefix + n_verify];
+            ok = decode_one(ctx_full, tokens[n_prefix + n_verify], replacement_pos) &&
+                 decode_range(ctx_depth, replay_tokens, n_prefix, 2 + n_accepted);
+        }
+        if (ok) {
+            ok = compare_logits(ctx_full, ctx_depth, n_vocab, n_accepted, "replacement", -1.0f);
+        }
+        if (ok) {
+            ok = decode_one(ctx_full, tokens[n_prefix + n_verify + 1], replacement_pos + 1) &&
+                 decode_one(ctx_depth, tokens[n_prefix + n_verify + 1], replacement_pos + 1) &&
+                 compare_logits(ctx_full, ctx_depth, n_vocab, n_accepted, "continuation", -1.0f);
+        }
+
+        llama_free(ctx_full);
+        llama_free(ctx_depth);
+
+        if (!ok) {
+            fprintf(stderr, "%s : fallback validation failed after accepting %u draft tokens\n", __func__, n_accepted);
+            return false;
+        }
+    }
+
+    fprintf(stderr, "%s : depth-1 recomputation preserves depth-5 greedy decisions at every rejection position\n", __func__);
+    return true;
+}
+
 static int test_rollback(const common_params & params, llama_model * model, uint8_t fill) {
     const llama_vocab * vocab   = llama_model_get_vocab(model);
     const int           n_vocab = llama_vocab_n_tokens(vocab);
@@ -456,6 +712,29 @@ int main(int argc, char ** argv) {
         return 0;
     }
 
+    const llama_vocab * vocab   = llama_model_get_vocab(model);
+    const int           n_vocab = llama_vocab_n_tokens(vocab);
+
+    std::vector<llama_token> tokens;
+    if (llama_vocab_type(vocab) == LLAMA_VOCAB_TYPE_NONE) {
+        tokens = { 1, 2, 3, 4, 5, 6, 7, 8, 9 };
+    } else {
+        tokens = common_tokenize(vocab, "The quick brown fox jumps over the lazy dog", true);
+    }
+    if (tokens.empty()) {
+        fprintf(stderr, "%s : not enough prompt tokens\n", __func__);
+        return 1;
+    }
+    if (!test_fragmented_on_device_fallback(params, model, tokens)) {
+        return 1;
+    }
+    if (!test_preallocated_device_checkpoints(params, model, tokens)) {
+        return 1;
+    }
+    if (!test_recompute_fallback(params, model, tokens, n_vocab)) {
+        return 1;
+    }
+
     for (uint8_t fill : { 0, 0x3e }) {
         fprintf(stderr, "%s : testing with cache fill 0x%02x\n", __func__, fill);
         if (test_rollback(params, model, fill) != 0) {
diff --git a/tools/server/README.md b/tools/server/README.md
index ef903340..882a83a3 100644
--- a/tools/server/README.md
+++ b/tools/server/README.md
@@ -259,6 +259,8 @@ For the full list of features, please refer to [server's changelog](https://gith
 | `--spec-draft-cpu-moe, -cmoed, --cpu-moe-draft` | keep all Mixture of Experts (MoE) weights in the CPU for the draft model<br/>(env: LLAMA_ARG_SPEC_DRAFT_CPU_MOE) |
 | `--spec-draft-n-cpu-moe, --spec-draft-ncmoe, -ncmoed, --n-cpu-moe-draft N` | keep the Mixture of Experts (MoE) weights of the first N layers in the CPU for the draft model<br/>(env: LLAMA_ARG_SPEC_DRAFT_N_CPU_MOE) |
 | `--spec-draft-n-max N` | number of tokens to draft for speculative decoding (default: 3)<br/>(env: LLAMA_ARG_SPEC_DRAFT_N_MAX) |
+| `--spec-mtp-cr-depth N` | MTP Compact Rollback depth; lower values save memory but replay accepted tokens after deep rejection (default: `--spec-draft-n-max`)<br/>(env: LLAMA_ARG_SPEC_MTP_CR_DEPTH) |
+| `--spec-draft-adaptive` | size each draft from measured acceptance rather than always drafting `--spec-draft-n-max` (default: false)<br/>(env: LLAMA_ARG_SPEC_DRAFT_ADAPTIVE) |
 | `--spec-draft-n-min N` | minimum number of draft tokens to use for speculative decoding (default: 0)<br/>(env: LLAMA_ARG_SPEC_DRAFT_N_MIN) |
 | `--spec-synth-len L` | target mean synthetic acceptance length, including the target token (benchmarking only)<br/>(env: LLAMA_ARG_SPEC_SYNTH_LEN) |
 | `--spec-synth-rates P0,P1,...` | comma-separated unconditional per-position synthetic acceptance probabilities (benchmarking only)<br/>(env: LLAMA_ARG_SPEC_SYNTH_RATES) |
diff --git a/tools/server/server-common.cpp b/tools/server/server-common.cpp
index 7bf1138c..29f38301 100644
--- a/tools/server/server-common.cpp
+++ b/tools/server/server-common.cpp
@@ -100,6 +100,10 @@ json server_slot_stats::to_json() const {
         base["draft_n"]          = n_draft_tokens;
         base["draft_n_accepted"] = n_draft_accepted;
     }
+    if (n_draft_replay_count > 0) {
+        base["draft_replay_count"] = n_draft_replay_count;
+        base["draft_replay_n"]     = n_draft_replay_tokens;
+    }
 
     return base;
 }
diff --git a/tools/server/server-common.h b/tools/server/server-common.h
index 9894f5f0..37639f06 100644
--- a/tools/server/server-common.h
+++ b/tools/server/server-common.h
@@ -361,6 +361,8 @@ struct server_slot_stats {
     uint64_t n_draft_tokens      = 0;
     uint64_t n_draft_accepted    = 0;
     uint64_t n_draft_verif_steps = 0;
+    uint64_t n_draft_replay_count  = 0;
+    uint64_t n_draft_replay_tokens = 0;
 
     // these are absolute timestamps (in us)
     // note: must be signed - they are subtracted before the later ones are set
diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp
index b6835e43..e0050ff5 100644
--- a/tools/server/server-context.cpp
+++ b/tools/server/server-context.cpp
@@ -256,6 +256,7 @@ struct server_slot {
     std::vector<int32_t> spec_i_batch;
     common_prompt_checkpoint spec_ckpt;
     bool spec_is_replay = false;
+    int32_t spec_n_accepted_observed = -1;
     std::mt19937 spec_synth_rng;
 
     // TODO: move members that belong to the task (such as `generated_text`, `has_new_line`) to task_results_state
@@ -370,6 +371,7 @@ struct server_slot {
         SLT_DBG(*this, "%s", "\n");
 
         spec_is_replay = false;
+        spec_n_accepted_observed = -1;
 
         last_nl_pos    = 0;
         generated_text = "";
@@ -680,6 +682,12 @@ struct server_slot {
                     "     acc per pos = (%s)\n", acceptance_rates_per_pos.c_str());
         }
 
+        if (stats.n_draft_replay_count > 0) {
+            SLT_INF(*this,
+                    "     MTP replays = %10" PRIu64 " events / %5" PRIu64 " tokens\n",
+                    stats.n_draft_replay_count, stats.n_draft_replay_tokens);
+        }
+
         common_speculative_print_stats(spec);
     }
 
@@ -892,6 +900,7 @@ private:
 
     common_context_seq_rm_type ctx_tgt_seq_rm_type = COMMON_CONTEXT_SEQ_RM_TYPE_NO;
     common_context_seq_rm_type ctx_dft_seq_rm_type = COMMON_CONTEXT_SEQ_RM_TYPE_NO;
+    bool spec_mtp_device_checkpoint = false;
 
     common_speculative_ptr spec;
 
@@ -1088,6 +1097,90 @@ private:
             }
         }
 
+        // The on-device recurrent speculative checkpoint is allocated lazily, so reserve
+        // the additional recurrent-state group before the joint target / MTP fit.
+        if (params_base.fit_params && spec_mtp) {
+            const uint32_t n_rs_seq = params_base.speculative.need_n_rs_seq();
+            const int32_t draft_n_max = params_base.speculative.draft.n_max;
+            if (draft_n_max < 0) {
+                SRV_ERR("%s", "[spec] invalid negative MTP draft.n_max while fitting\n");
+                return false;
+            }
+            if (n_rs_seq < static_cast<uint32_t>(draft_n_max)) {
+                try {
+                    common_params params_rs = params_base;
+                    auto mparams_rs = common_model_params_to_llama(params_rs);
+                    auto cparams_rs = common_context_params_to_llama(params_rs);
+
+                    std::vector<ggml_backend_dev_t> devs_rs;
+                    uint32_t hp_ngl_rs = 0;
+                    uint32_t hp_nct_rs = 0;
+                    uint32_t hp_nex_rs = 0;
+                    cparams_rs.n_rs_seq = n_rs_seq;
+                    const auto dmd_rs = common_get_device_memory_data(
+                        params_base.model.path.c_str(), &mparams_rs, &cparams_rs,
+                        devs_rs, hp_ngl_rs, hp_nct_rs, hp_nex_rs, GGML_LOG_LEVEL_ERROR);
+
+                    std::vector<ggml_backend_dev_t> devs_rs_next;
+                    uint32_t hp_ngl_rs_next = 0;
+                    uint32_t hp_nct_rs_next = 0;
+                    uint32_t hp_nex_rs_next = 0;
+                    cparams_rs.n_rs_seq = n_rs_seq + 1;
+                    const auto dmd_rs_next = common_get_device_memory_data(
+                        params_base.model.path.c_str(), &mparams_rs, &cparams_rs,
+                        devs_rs_next, hp_ngl_rs_next, hp_nct_rs_next, hp_nex_rs_next, GGML_LOG_LEVEL_ERROR);
+
+                    if (devs_rs.size() != devs_rs_next.size() || dmd_rs.size() < devs_rs.size() ||
+                            dmd_rs_next.size() < devs_rs_next.size()) {
+                        throw std::runtime_error("target context device count changed during recurrent-state measurement");
+                    }
+
+                    // common_fit_params() consumes fit_params_target in the measured model-device
+                    // order. The configured device list can contain a nullptr sentinel, CPU/ACCEL
+                    // devices, or be reduced by split-mode none, so it is not a safe index map.
+                    if (params_base.fit_params_target.size() < devs_rs.size()) {
+                        throw std::runtime_error("fit_params_target has no entry for every target device");
+                    }
+
+                    size_t total = 0;
+                    for (size_t j = 0; j < devs_rs.size(); ++j) {
+                        const auto next_dev = std::find(devs_rs_next.begin(), devs_rs_next.end(), devs_rs[j]);
+                        if (next_dev == devs_rs_next.end()) {
+                            throw std::runtime_error("target context device mapping changed during recurrent-state measurement");
+                        }
+                        const size_t next_index = next_dev - devs_rs_next.begin();
+                        if (next_index != j) {
+                            throw std::runtime_error("target context device order changed during recurrent-state measurement");
+                        }
+                        if (dmd_rs_next[j].context < dmd_rs[j].context) {
+                            throw std::runtime_error("recurrent-state context measurement decreased");
+                        }
+                        const size_t delta = dmd_rs_next[j].context - dmd_rs[j].context;
+                        const size_t checkpoint = dmd_rs[j].checkpoint;
+                        if (checkpoint == 0 && delta != 0) {
+                            throw std::runtime_error("recurrent checkpoint sizing produced no device allocation");
+                        }
+                        params_base.fit_params_target[j] += checkpoint;
+                        total += checkpoint;
+                        SRV_INF("[spec] recurrent checkpoint fit reservation: device %s, %.2f MiB "
+                                "(exact backend allocation; recurrent-plane delta %.2f MiB, context %.2f -> %.2f MiB)\n",
+                                ggml_backend_dev_name(devs_rs[j]), checkpoint / (1024.0 * 1024.0),
+                                delta / (1024.0 * 1024.0),
+                                dmd_rs[j].context / (1024.0 * 1024.0),
+                                dmd_rs_next[j].context / (1024.0 * 1024.0));
+                    }
+                    if (total == 0) {
+                        throw std::runtime_error("recurrent-state measurement produced no device reservation");
+                    }
+                    SRV_INF("[spec] recurrent checkpoint fit reservation: %.2f MiB total\n",
+                            total / (1024.0 * 1024.0));
+                } catch (const std::exception & e) {
+                    SRV_ERR("[spec] failed to reserve recurrent-state memory before fitting: %s\n", e.what());
+                    return false;
+                }
+            }
+        }
+
         // note: the draft / MTP context is fitted together with the target model, see common_fit_extra_model
 
         // attach a progress callback
@@ -1279,6 +1372,16 @@ private:
             model_dft = nullptr;
         }
 
+        spec_mtp_device_checkpoint = false;
+        if (spec && ctx_tgt_seq_rm_type == COMMON_CONTEXT_SEQ_RM_TYPE_RS) {
+            spec_mtp_device_checkpoint = llama_state_seq_reserve_device_buffers(ctx_tgt);
+            if (spec_mtp_device_checkpoint) {
+                SRV_INF("%s", "[spec] reserved device recurrent checkpoints before evaluation\n");
+            } else {
+                SRV_WRN("%s", "[spec] device recurrent checkpoint reservation failed; using host checkpoints\n");
+            }
+        }
+
         if (!spec && params_base.speculative.has_synth()) {
             SRV_ERR("%s", "synthetic acceptance requires an initialized speculative decoding context\n");
             return false;
@@ -3076,9 +3179,23 @@ private:
                    (ctx_dft_seq_rm_type == COMMON_CONTEXT_SEQ_RM_TYPE_RS && draft.size() > llama_n_rs_seq(ctx_dft));
 
                 if (use_ckpt_tgt) {
+                    llama_state_seq_flags ckpt_flags = LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY;
+                    if (ctx_tgt_seq_rm_type == COMMON_CONTEXT_SEQ_RM_TYPE_RS && spec_mtp_device_checkpoint) {
+                        // Phase 1 proof of concept: avoid copying the recurrent fallback checkpoint
+                        // through host memory on every speculative round. The device buffer is
+                        // allocated during load_model(), before evaluation can consume the
+                        // fitted headroom reserved for the recurrent-state checkpoint.
+                        ckpt_flags |= LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+                    }
+
                     //const int64_t t_start = ggml_time_us();
 
-                    ckpt.update_tgt(ctx_tgt, slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+                    ckpt.update_tgt(ctx_tgt, slot.id, ckpt_flags);
+
+                    if ((ckpt_flags & LLAMA_STATE_SEQ_FLAGS_ON_DEVICE) && !ckpt.data_tgt_on_device) {
+                        spec_mtp_device_checkpoint = false;
+                        SRV_WRN("%s", "[spec] device recurrent checkpoint save failed; using host checkpoints for the rest of this server run\n");
+                    }
 
                     //const int64_t t_total = ggml_time_us() - t_start;
                     //printf("checkpoint total: %f ms\n", t_total / 1000.0);
@@ -3927,6 +4044,10 @@ private:
                     ctx_tgt_seq_rm_type == COMMON_CONTEXT_SEQ_RM_TYPE_FULL ||
                     (ctx_tgt_seq_rm_type == COMMON_CONTEXT_SEQ_RM_TYPE_RS && n_rollback > llama_n_rs_seq(ctx_tgt));
 
+                const bool use_rs_replay =
+                    ctx_tgt_seq_rm_type == COMMON_CONTEXT_SEQ_RM_TYPE_RS &&
+                    n_rollback > llama_n_rs_seq(ctx_tgt);
+
                 // check for partial draft acceptance
                 if (n_rollback > 0) {
                     if (use_ckpt_tgt) {
@@ -3934,15 +4055,32 @@ private:
                             SLT_INF(slot, "accepted %2zu/%2zu draft tokens (restore checkpoint)\n", accepted.size() - 1, slot.spec_draft.size());
                         }
 
+                        // Preserve the original observation for adaptive draft sizing. The
+                        // replay batch also contains the target replacement token, so its
+                        // physical acceptance count is one larger than the measurement that
+                        // selected the rollback path.
+                        if (!slot.spec_is_replay) {
+                            slot.spec_n_accepted_observed = static_cast<int32_t>(accepted.size() - 1);
+                        }
+
                         // partial acceptance is not supported by the context -> truncate the draft and restore the state
                         slot.spec_is_replay = true;
                         slot.spec_draft = std::move(accepted);
 
+                        if (use_rs_replay) {
+                            slot.stats.n_draft_replay_count  += 1;
+                            slot.stats.n_draft_replay_tokens += slot.spec_draft.size() + 1;
+                        }
+
                         const auto & ckpt = slot.spec_ckpt;
 
                         SLT_DBG(slot, "restoring speculative checkpoint (pos_min = %d, pos_max = %d, size = %zu)\n", ckpt.pos_min, ckpt.pos_max, ckpt.size());
 
-                        ckpt.load_tgt(slot.ctx_tgt, slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
+                        llama_state_seq_flags ckpt_flags = LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY;
+                        if (use_rs_replay) {
+                            ckpt_flags |= LLAMA_STATE_SEQ_FLAGS_ON_DEVICE;
+                        }
+                        ckpt.load_tgt(slot.ctx_tgt, slot.id, ckpt_flags);
 
                         if (slot.ctx_dft) {
                             ckpt.load_dft(slot.ctx_dft, slot.id, LLAMA_STATE_SEQ_FLAGS_PARTIAL_ONLY);
@@ -3961,7 +4099,12 @@ private:
                     SLT_INF(slot, "accepted %2zu/%2zu draft tokens\n", accepted.size() - 1, n_draft);
                 }
 
-                common_speculative_accept(spec.get(), slot.id, accepted.size() - 1);
+                common_speculative_accept(
+                        spec.get(),
+                        slot.id,
+                        accepted.size() - 1,
+                        slot.spec_is_replay ? slot.spec_n_accepted_observed : -1);
+                slot.spec_n_accepted_observed = -1;
 
                 slot.spec_draft = std::move(accepted);
             }
diff --git a/tools/server/tests/utils.py b/tools/server/tests/utils.py
index 826aef2d..268df76d 100644
--- a/tools/server/tests/utils.py
+++ b/tools/server/tests/utils.py
@@ -99,6 +99,7 @@ class ServerProcess:
     spec_type: str | None = None
     spec_draft_n_min: int | None = None
     spec_draft_n_max: int | None = None
+    spec_mtp_cr_depth: int | None = None
     spec_synth_len: float | None = None
     spec_synth_rates: List[float] | None = None
     no_ui: bool | None = None
@@ -247,6 +248,8 @@ class ServerProcess:
             server_args.extend(["--spec-draft-n-max", self.spec_draft_n_max])
         if self.spec_draft_n_min:
             server_args.extend(["--spec-draft-n-min", self.spec_draft_n_min])
+        if self.spec_mtp_cr_depth:
+            server_args.extend(["--spec-mtp-cr-depth", self.spec_mtp_cr_depth])
         if self.spec_synth_len is not None:
             server_args.extend(["--spec-synth-len", self.spec_synth_len])
         if self.spec_synth_rates is not None: