diff --git a/common/CMakeLists.txt b/common/CMakeLists.txt
index 36f1e0cd50f1..1e87312d2071 100644
--- a/common/CMakeLists.txt
+++ b/common/CMakeLists.txt
@@ -100,6 +100,8 @@ add_library(${TARGET}
     reasoning-budget.h
     sampling.cpp
     sampling.h
+    speculative-prefill.cpp
+    speculative-prefill.h
     speculative.cpp
     speculative.h
     subproc.cpp
diff --git a/common/arg.cpp b/common/arg.cpp
index 84e2800256bb..2b5e5f937273 100644
--- a/common/arg.cpp
+++ b/common/arg.cpp
@@ -361,6 +361,7 @@ static bool spec_types_is_default(const common_params & params) {
 common_models_handler common_models_handler_init(const common_params & params, llama_example curr_ex) {
     common_download_hf_plan plan;
     common_download_hf_plan plan_spec;
+    common_download_hf_plan plan_prefill;
     common_download_opts opts;
 
     const bool spec_type_draft_mtp = std::find(params.speculative.types.begin(),
@@ -413,7 +414,11 @@ common_models_handler common_models_handler_init(const common_params & params, l
         plan_spec = common_download_get_hf_plan(params.speculative.draft.mparams, opts_spec);
     }
 
-    return common_models_handler{plan, plan_spec, opts};
+    if (!params.speculative.prefill.model.hf_repo.empty()) {
+        plan_prefill = common_download_get_hf_plan(params.speculative.prefill.model, opts);
+    }
+
+    return common_models_handler{plan, plan_spec, plan_prefill, opts};
 }
 
 bool common_models_handler_is_preset_repo(const common_models_handler & handler) {
@@ -461,8 +466,9 @@ static std::vector<common_download_task> build_url_tasks(const common_params_mod
 void common_models_handler_apply(common_models_handler & handler, common_params & params, common_download_callback * callback) {
     std::vector<common_download_task> tasks;
 
-    auto & plan      = handler.plan;
-    auto & plan_spec = handler.plan_spec;
+    auto & plan         = handler.plan;
+    auto & plan_spec    = handler.plan_spec;
+    auto & plan_prefill = handler.plan_prefill;
 
     auto opts = handler.opts; // copy
     opts.callback = callback;
@@ -478,6 +484,7 @@ void common_models_handler_apply(common_models_handler & handler, common_params
     handle_url(params.model);
     handle_url(params.mmproj);
     handle_url(params.speculative.draft.mparams);
+    handle_url(params.speculative.prefill.model);
 
     // optionally, if docker repo is set, resolve it
     if (!params.model.docker_repo.empty()) {
@@ -513,6 +520,13 @@ void common_models_handler_apply(common_models_handler & handler, common_params
         tasks.push_back(task);
         had_spec_url = true;
     }
+    if (!params.speculative.prefill.model.url.empty()) {
+        common_download_task task;
+        task.url        = params.speculative.prefill.model.url;
+        task.local_path = params.speculative.prefill.model.path;
+        task.opts       = opts;
+        tasks.push_back(task);
+    }
 
     // handle hf_plan tasks
     auto add_tasks = [&opts, &tasks](const hf_cache::hf_files  & model_files,
@@ -626,6 +640,11 @@ void common_models_handler_apply(common_models_handler & handler, common_params
         had_spec_url = true;
     }
 
+    // handle plan_prefill (e.g. --spec-prefill-hf)
+    if (!plan_prefill.model_files.empty()) {
+        add_tasks(plan_prefill.model_files, plan_prefill.primary, params.speculative.prefill.model);
+    }
+
     if (!plan.model_files.empty()) {
         add_tasks(plan.model_files, plan.primary, params.model);
     }
@@ -4382,6 +4401,111 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
             params.speculative.draft.mparams.hf_file = value; // will be used if --spec-draft-hf is set
         }
     ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_DRAFT_MODEL"));
+    add_opt(common_arg(
+        {"--spec-prefill", "--speculative-prefill"},
+        "enable speculative prefill using draft model to filter prompt tokens",
+        [](common_params & params) {
+            params.speculative.prefill.enabled = true;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_PREFILL"));
+    add_opt(common_arg(
+        {"--spec-prefill-draft-model", "--spec-prefill-model", "-mpd", "--speculative-prefill-model", "--speculative-prefill-draft-model"}, "FNAME",
+        "draft model for speculative prefill (default: unused)",
+        [](common_params & params, const std::string & value) {
+            params.speculative.prefill.model.path    = value;
+            params.speculative.prefill.model.hf_file = value; // will be used if --spec-prefill-hf is set
+            params.speculative.prefill.enabled       = true;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_PREFILL_MODEL"));
+    add_opt(common_arg(
+        {"--spec-prefill-draft-hf", "--spec-prefill-hf", "-hfpd", "--speculative-prefill-hf", "--speculative-prefill-draft-hf"}, "<user>/<model>[:quant]",
+        "Hugging Face model repository for speculative prefill draft model (default: unused)",
+        [](common_params & params, const std::string & value) {
+            params.speculative.prefill.model.hf_repo = value;
+            params.speculative.prefill.enabled       = true;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_PREFILL_HF"));
+    add_opt(common_arg(
+        {"--spec-prefill-draft-ngl", "--spec-prefill-ngl", "-nglpd", "--speculative-prefill-ngl", "--speculative-prefill-draft-ngl"}, "N",
+        string_format("max. number of speculative prefill draft model layers to store in VRAM, either an exact number, 'auto', or 'all' (default: %s)",
+            params.speculative.prefill.n_gpu_layers == -1 ? "auto" : "all"),
+        [](common_params & params, const std::string & value) {
+            if (value == "auto") {
+                params.speculative.prefill.n_gpu_layers = -1;
+            } else if (value == "all") {
+                params.speculative.prefill.n_gpu_layers = -2;
+            } else {
+                params.speculative.prefill.n_gpu_layers = std::stoi(value);
+            }
+            if (!llama_supports_gpu_offload()) {
+                fprintf(stderr, "warning: no usable GPU found, --spec-prefill-ngl option will be ignored\n");
+                fprintf(stderr, "warning: one possible reason is that llama.cpp was compiled without GPU support\n");
+                fprintf(stderr, "warning: consult docs/build.md for compilation instructions\n");
+            }
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_PREFILL_N_GPU_LAYERS"));
+    add_opt(common_arg(
+        {"--spec-prefill-draft-device", "--spec-prefill-device", "-devpd", "--speculative-prefill-device", "--speculative-prefill-draft-device"}, "<dev1,dev2,..>",
+        "comma-separated list of devices to use for offloading the speculative prefill draft model (none = don't offload)\n"
+        "use --list-devices to see a list of available devices",
+        [](common_params & params, const std::string & value) {
+            params.speculative.prefill.devices = parse_device_list(value);
+            params.speculative.prefill.enabled = true;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_PREFILL_DEVICE"));
+    add_opt(common_arg(
+        {"--spec-prefill-draft-ctx", "--spec-prefill-ctx", "--spec-prefill-ctx-size", "--spec-prefill-max-ctx", "-cpd", "--speculative-prefill-ctx", "--speculative-prefill-max-ctx"}, "N",
+        string_format("context size for speculative prefill draft model (default: %d, 0 = draft training limit or main context)", params.speculative.prefill.n_ctx),
+        [](common_params & params, int value) {
+            if (value < 0) {
+                throw std::invalid_argument("spec-prefill context size must be >= 0");
+            }
+            params.speculative.prefill.n_ctx   = value;
+            params.speculative.prefill.enabled = true;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_PREFILL_CTX_SIZE"));
+    add_opt(common_arg(
+        {"--spec-prefill-p", "--spec-prefill-percentage"}, "P",
+        string_format("fraction of prompt tokens to retain during speculative prefill (default: %.2f)", (double) params.speculative.prefill.percentage),
+        [](common_params & params, const std::string & value) {
+            const float val = std::stof(value);
+            if (val <= 0.0f || val > 1.0f) {
+                throw std::invalid_argument("spec-prefill percentage must be between 0.0 (exclusive) and 1.0 (inclusive)");
+            }
+            params.speculative.prefill.enabled    = true;
+            params.speculative.prefill.percentage = val;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_SPEC_PREFILL_P"));
+    add_opt(common_arg(
+        {"--spec-prefill-chunk", "--spec-prefill-chunk-size"}, "N",
+        string_format("chunk grouping size for speculative prefill (default: %d, 0 to disable)", params.speculative.prefill.chunk_size),
+        [](common_params & params, int value) {
+            if (value < 0) {
+                throw std::invalid_argument("spec-prefill chunk size must be >= 0");
+            }
+            params.speculative.prefill.chunk_size = value;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}));
+    add_opt(common_arg(
+        {"--spec-prefill-lookahead", "--spec-prefill-lah"}, "N",
+        string_format("number of lookahead decode steps on draft model for attention estimation (default: %d)", params.speculative.prefill.look_ahead_cnt),
+        [](common_params & params, int value) {
+            if (value < 1) {
+                throw std::invalid_argument("spec-prefill lookahead count must be >= 1");
+            }
+            params.speculative.prefill.look_ahead_cnt = value;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}));
+    add_opt(common_arg(
+        {"--spec-prefill-pool-kernel"}, "N",
+        string_format("1D average pooling kernel size for attention smoothing (default: %d, 0 to disable)", params.speculative.prefill.pool_kernel_size),
+        [](common_params & params, int value) {
+            if (value < 0) {
+                throw std::invalid_argument("spec-prefill pool kernel size must be >= 0");
+            }
+            params.speculative.prefill.pool_kernel_size = value;
+        }
+    ).set_spec().set_examples({LLAMA_EXAMPLE_SPECULATIVE, LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}));
     add_opt(common_arg(
         {"--spec-type"}, common_speculative_all_types_str(),
         string_format("comma-separated list of types of speculative decoding to use (default: %s)\n",
@@ -4896,6 +5020,12 @@ void common_params_add_preset_options(std::vector<common_arg> & args) {
         [](common_params &, const std::string &) { /* unused */ }
     ).set_env(COMMON_ARG_PRESET_DEDUP_CACHE_MODELS).set_preset_only());
 
+    args.push_back(common_arg(
+        {"hidden"}, "0|1",
+        "in server router mode, omit this model from GET /models (it can still be loaded by name)",
+        [](common_params &, const std::string &) { /* unused */ }
+    ).set_env(COMMON_ARG_PRESET_HIDDEN).set_preset_only());
+
     // args.push_back(common_arg(
     //     {"pin"},
     //     "in server router mode, do not unload this model if models_max is exceeded",
diff --git a/common/arg.h b/common/arg.h
index 421bc295fc21..e86495a86971 100644
--- a/common/arg.h
+++ b/common/arg.h
@@ -14,6 +14,7 @@
 #define COMMON_ARG_PRESET_LOAD_ON_STARTUP    "__PRESET_LOAD_ON_STARTUP"
 #define COMMON_ARG_PRESET_STOP_TIMEOUT       "__PRESET_STOP_TIMEOUT"
 #define COMMON_ARG_PRESET_DEDUP_CACHE_MODELS "__PRESET_DEDUP_CACHE_MODELS"
+#define COMMON_ARG_PRESET_HIDDEN             "__PRESET_HIDDEN"
 
 //
 // CLI argument parsing
@@ -138,6 +139,7 @@ void common_params_add_preset_options(std::vector<common_arg> & args);
 struct common_models_handler {
     common_download_hf_plan plan;
     common_download_hf_plan plan_spec;
+    common_download_hf_plan plan_prefill;
     common_download_opts opts;
 };
 
diff --git a/common/common.h b/common/common.h
index 0162a5f1e638..431424304011 100644
--- a/common/common.h
+++ b/common/common.h
@@ -389,6 +389,20 @@ struct common_params_speculative_ngram_cache {
     std::string lookup_cache_dynamic; // path of dynamic ngram cache file for lookup decoding
 };
 
+struct common_params_speculative_prefill {
+    bool                enabled          = false; // enable speculative prefill
+    common_params_model model;                    // draft model for speculative prefill
+    int32_t             n_ctx            = 0;     // context size for draft model (0 = default/target context size)
+    int32_t             n_gpu_layers     = -1;    // max draft model layers to store in VRAM (-1 - use default)
+    std::vector<ggml_backend_dev_t> devices;      // devices to use for offloading the draft model
+    float               percentage       = 0.3f;  // fraction of prompt tokens to retain (0.0 < p <= 1.0)
+    int32_t             chunk_size       = 32;    // chunk grouping size (0 to disable chunking)
+    int32_t             look_ahead_cnt   = 8;     // lookahead decode steps on draft model
+    int32_t             pool_kernel_size = 13;    // 1D average pooling kernel size for smoothing
+    bool                keep_bos         = true;  // preserve first token (BOS)
+    bool                keep_last        = true;  // preserve last token / tail chunk
+};
+
 struct common_params_speculative {
     std::vector<enum common_speculative_type> types = { COMMON_SPECULATIVE_TYPE_NONE };
 
@@ -405,6 +419,8 @@ struct common_params_speculative {
 
     common_params_speculative_ngram_cache ngram_cache;
 
+    common_params_speculative_prefill prefill;
+
     bool has_dft() const {
         return !draft.mparams.empty();
     }
@@ -413,6 +429,9 @@ struct common_params_speculative {
         return synth_len != -1.0 || !synth_rates.empty();
     }
 
+    bool has_prefill() const {
+        return prefill.enabled && !prefill.model.empty();
+    }
     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;
diff --git a/common/speculative-prefill.cpp b/common/speculative-prefill.cpp
new file mode 100644
index 000000000000..ab20ad08f1e5
--- /dev/null
+++ b/common/speculative-prefill.cpp
@@ -0,0 +1,337 @@
+#include "speculative-prefill.h"
+
+#include "common.h"
+#include "ggml.h"
+#include "ggml-backend.h"
+#include "llama.h"
+#include "log.h"
+#include "sampling.h"
+
+#include <algorithm>
+#include <cmath>
+#include <cstring>
+#include <numeric>
+#include <vector>
+
+#define SPF_DBG(fmt, ...) LOG_DBG("spec-prefill: " fmt, __VA_ARGS__)
+#define SPF_INF(fmt, ...) LOG_INF("spec-prefill: " fmt, __VA_ARGS__)
+#define SPF_ERR(fmt, ...) LOG_ERR("spec-prefill: " fmt, __VA_ARGS__)
+
+std::vector<float> common_speculative_prefill_compute_importance(
+    const std::vector<std::vector<float>> & attn_layers,
+    int32_t n_prompt,
+    int32_t n_heads,
+    int32_t pool_kernel_size) {
+    if (attn_layers.empty() || n_prompt <= 0 || n_heads <= 0) {
+        return std::vector<float>(std::max(0, n_prompt), 1.0f);
+    }
+
+    std::vector<float> max_importance(n_prompt, 0.0f);
+    std::vector<float> smoothed(n_prompt, 0.0f);
+
+    const int32_t pad = pool_kernel_size > 0 ? (pool_kernel_size / 2) : 0;
+
+    for (const auto & layer_attn : attn_layers) {
+        if ((int32_t) layer_attn.size() < n_heads * n_prompt) {
+            continue;
+        }
+
+        for (int32_t h = 0; h < n_heads; ++h) {
+            const float * head_attn = layer_attn.data() + (size_t) h * n_prompt;
+
+            if (pool_kernel_size > 1) {
+                for (int32_t i = 0; i < n_prompt; ++i) {
+                    const int32_t start = std::max(0, i - pad);
+                    const int32_t end   = std::min(n_prompt - 1, i + pad);
+
+                    float sum = 0.0f;
+                    for (int32_t j = start; j <= end; ++j) {
+                        sum += head_attn[j];
+                    }
+                    smoothed[i] = sum / (float) (end - start + 1);
+                }
+            } else {
+                std::copy(head_attn, head_attn + n_prompt, smoothed.begin());
+            }
+
+            for (int32_t i = 0; i < n_prompt; ++i) {
+                if (smoothed[i] > max_importance[i]) {
+                    max_importance[i] = smoothed[i];
+                }
+            }
+        }
+    }
+
+    return max_importance;
+}
+
+std::vector<int32_t> common_speculative_prefill_select_indices(
+    const std::vector<float> & importance_scores,
+    const common_params_speculative_prefill & params) {
+    const int32_t n_prompt = (int32_t) importance_scores.size();
+    if (n_prompt <= 0) {
+        return {};
+    }
+
+    const float percentage = std::max(0.01f, std::min(1.0f, params.percentage));
+    std::vector<int32_t> selected;
+
+    if (params.chunk_size > 0 && n_prompt > params.chunk_size) {
+        const int32_t chunk_size = params.chunk_size;
+        const int32_t n_chunks = (n_prompt + chunk_size - 1) / chunk_size;
+
+        std::vector<std::pair<float, int32_t>> chunk_scores(n_chunks);
+
+        for (int32_t c = 0; c < n_chunks; ++c) {
+            const int32_t start = c * chunk_size;
+            const int32_t end   = std::min(n_prompt, (c + 1) * chunk_size);
+
+            float sum = 0.0f;
+            for (int32_t i = start; i < end; ++i) {
+                sum += importance_scores[i];
+            }
+            chunk_scores[c] = { sum / (float) (end - start), c };
+        }
+
+        const int32_t n_keep_chunks = std::max(1, (int32_t) std::ceil(n_chunks * percentage));
+
+        std::sort(chunk_scores.begin(), chunk_scores.end(),
+            [](const std::pair<float, int32_t> & a, const std::pair<float, int32_t> & b) {
+                return a.first > b.first;
+            });
+
+        for (int32_t k = 0; k < n_keep_chunks && k < n_chunks; ++k) {
+            const int32_t c = chunk_scores[k].second;
+            const int32_t start = c * chunk_size;
+            const int32_t end   = std::min(n_prompt, (c + 1) * chunk_size);
+
+            for (int32_t i = start; i < end; ++i) {
+                selected.push_back(i);
+            }
+        }
+    } else {
+        const int32_t n_keep = std::max(1, (int32_t) std::ceil(n_prompt * percentage));
+
+        std::vector<std::pair<float, int32_t>> token_scores(n_prompt);
+        for (int32_t i = 0; i < n_prompt; ++i) {
+            token_scores[i] = { importance_scores[i], i };
+        }
+
+        std::sort(token_scores.begin(), token_scores.end(),
+            [](const std::pair<float, int32_t> & a, const std::pair<float, int32_t> & b) {
+                return a.first > b.first;
+            });
+
+        for (int32_t k = 0; k < n_keep; ++k) {
+            selected.push_back(token_scores[k].second);
+        }
+    }
+
+    if (params.keep_bos) {
+        selected.push_back(0);
+    }
+
+    if (params.keep_last) {
+        if (params.chunk_size > 0 && n_prompt > params.chunk_size) {
+            const int32_t start = ((n_prompt - 1) / params.chunk_size) * params.chunk_size;
+            for (int32_t i = start; i < n_prompt; ++i) {
+                selected.push_back(i);
+            }
+        } else {
+            selected.push_back(n_prompt - 1);
+        }
+    }
+
+    std::sort(selected.begin(), selected.end());
+    selected.erase(std::unique(selected.begin(), selected.end()), selected.end());
+
+    return selected;
+}
+
+struct cb_attn_collector_data {
+    int32_t n_prompt = 0;
+    std::vector<std::vector<float>> layer_attns;
+};
+
+static bool cb_collect_attn(ggml_tensor * t, bool ask, void * user_data) {
+    if (ask) {
+        return strncmp(t->name, "kq_soft_max", 11) == 0;
+    }
+
+    if (strncmp(t->name, "kq_soft_max", 11) != 0) {
+        return true;
+    }
+
+    auto * data = (cb_attn_collector_data *) user_data;
+    if (data == nullptr || data->n_prompt <= 0) {
+        return true;
+    }
+
+    if (t->type != GGML_TYPE_F32) {
+        return true;
+    }
+
+    const int32_t n_past_total = (int32_t) t->ne[0];
+    const int32_t n_heads      = (int32_t) t->ne[2];
+    const int32_t n_prompt     = data->n_prompt;
+
+    if (t->nb[0] != sizeof(float) || n_heads <= 0 || n_past_total < n_prompt) {
+        return true;
+    }
+
+    std::vector<float> layer_data((size_t) n_heads * n_prompt);
+
+    for (int32_t h = 0; h < n_heads; ++h) {
+        float * dst = layer_data.data() + (size_t) h * n_prompt;
+        ggml_backend_tensor_get(t, dst, (size_t) h * t->nb[2], (size_t) n_prompt * sizeof(float));
+    }
+
+    data->layer_attns.push_back(std::move(layer_data));
+
+    return true;
+}
+
+common_speculative_prefill_result common_speculative_prefill_execute(
+    llama_context * ctx_dft,
+    common_sampler * smpl_dft,
+    const std::vector<llama_token> & prompt,
+    llama_seq_id seq_id,
+    const common_params_speculative_prefill & params) {
+    common_speculative_prefill_result res;
+    res.n_prompt_orig = (int32_t) prompt.size();
+
+    if (prompt.empty()) {
+        return res;
+    }
+
+    if (!params.enabled || params.percentage >= 1.0f || (int32_t) prompt.size() <= params.chunk_size) {
+        res.kept_indices.resize(prompt.size());
+        std::iota(res.kept_indices.begin(), res.kept_indices.end(), 0);
+        res.n_prompt_kept = (int32_t) res.kept_indices.size();
+        res.importance_scores.assign(prompt.size(), 1.0f);
+        return res;
+    }
+
+    const int32_t n_ctx_dft = llama_n_ctx(ctx_dft);
+    const int32_t lookahead = std::max(1, params.look_ahead_cnt);
+
+    if ((int32_t) prompt.size() + lookahead > n_ctx_dft) {
+        SPF_INF("prompt size (%d) + lookahead (%d) exceeds draft context (%d); skipping speculative prefill\n",
+                (int32_t) prompt.size(), lookahead, n_ctx_dft);
+        res.kept_indices.resize(prompt.size());
+        std::iota(res.kept_indices.begin(), res.kept_indices.end(), 0);
+        res.n_prompt_kept = (int32_t) res.kept_indices.size();
+        res.importance_scores.assign(prompt.size(), 1.0f);
+        return res;
+    }
+
+    const auto t_start = ggml_time_us();
+
+    const llama_model * model_dft = llama_get_model(ctx_dft);
+    const llama_vocab * vocab_dft = llama_model_get_vocab(model_dft);
+    const int32_t n_heads_dft     = llama_model_n_head(model_dft);
+
+    // 1. evaluate full prompt on draft model
+    {
+        const int32_t n_batch_dft = llama_n_batch(ctx_dft);
+        llama_batch batch_prompt = llama_batch_init(std::min((int32_t) prompt.size(), n_batch_dft), 0, 1);
+
+        for (int32_t i = 0; i < (int32_t) prompt.size(); i += n_batch_dft) {
+            const int32_t n_eval = std::min((int32_t) prompt.size() - i, n_batch_dft);
+            common_batch_clear(batch_prompt);
+
+            for (int32_t j = 0; j < n_eval; ++j) {
+                const int32_t idx = i + j;
+                const bool is_last = (idx == (int32_t) prompt.size() - 1);
+                common_batch_add(batch_prompt, prompt[idx], (llama_pos) idx, { seq_id }, is_last);
+            }
+
+            const int ret = llama_decode(ctx_dft, batch_prompt);
+            if (ret != 0) {
+                SPF_ERR("failed to decode prompt on draft model, ret = %d\n", ret);
+                llama_batch_free(batch_prompt);
+                res.kept_indices.resize(prompt.size());
+                std::iota(res.kept_indices.begin(), res.kept_indices.end(), 0);
+                res.n_prompt_kept = (int32_t) res.kept_indices.size();
+                return res;
+            }
+        }
+        llama_batch_free(batch_prompt);
+    }
+
+    const auto t_prefill_end = ggml_time_us();
+    res.t_draft_eval_us = t_prefill_end - t_start;
+
+    // 2. lookahead decode steps with attention extraction
+    std::vector<float> total_importance(prompt.size(), 0.0f);
+
+    cb_attn_collector_data cb_data;
+    cb_data.n_prompt = (int32_t) prompt.size();
+
+    // attach callback
+    llama_set_eval_callback(ctx_dft, cb_collect_attn, &cb_data);
+
+    llama_batch batch_decode = llama_batch_init(1, 0, 1);
+
+    int32_t actual_steps = 0;
+    int32_t cur_pos = (int32_t) prompt.size();
+
+    for (int32_t k = 0; k < lookahead; ++k) {
+        cb_data.layer_attns.clear();
+
+        const llama_token token_id = common_sampler_sample(smpl_dft, ctx_dft, -1);
+        common_sampler_accept(smpl_dft, token_id, true);
+
+        if (llama_vocab_is_eog(vocab_dft, token_id)) {
+            break;
+        }
+
+        common_batch_clear(batch_decode);
+        common_batch_add(batch_decode, token_id, cur_pos++, { seq_id }, true);
+
+        const int ret = llama_decode(ctx_dft, batch_decode);
+        if (ret != 0) {
+            SPF_ERR("failed lookahead decode step %d, ret = %d\n", k, ret);
+            break;
+        }
+
+        if (!cb_data.layer_attns.empty()) {
+            const auto step_importance = common_speculative_prefill_compute_importance(
+                cb_data.layer_attns,
+                (int32_t) prompt.size(),
+                n_heads_dft,
+                params.pool_kernel_size);
+
+            for (size_t i = 0; i < prompt.size(); ++i) {
+                total_importance[i] += step_importance[i];
+            }
+            actual_steps++;
+        }
+    }
+
+    llama_batch_free(batch_decode);
+
+    // detach callback
+    llama_set_eval_callback(ctx_dft, nullptr, nullptr);
+
+    if (actual_steps == 0) {
+        LOG_WRN("spec-prefill: attention was not captured (flash attn on?); keeping full prompt\n");
+        res.kept_indices.resize(prompt.size());
+        std::iota(res.kept_indices.begin(), res.kept_indices.end(), 0);
+        res.n_prompt_kept = (int32_t) res.kept_indices.size();
+        res.importance_scores.assign(prompt.size(), 1.0f);
+        return res;
+    }
+
+    for (size_t i = 0; i < prompt.size(); ++i) {
+        total_importance[i] /= (float) actual_steps;
+    }
+
+    const auto t_est_start = ggml_time_us();
+    res.importance_scores = total_importance;
+    res.kept_indices = common_speculative_prefill_select_indices(total_importance, params);
+    res.n_prompt_kept = (int32_t) res.kept_indices.size();
+    res.t_estimate_us = ggml_time_us() - t_est_start;
+
+    return res;
+}
diff --git a/common/speculative-prefill.h b/common/speculative-prefill.h
new file mode 100644
index 000000000000..0752449133df
--- /dev/null
+++ b/common/speculative-prefill.h
@@ -0,0 +1,37 @@
+#pragma once
+
+#include "llama.h"
+#include "common.h"
+
+#include <cstdint>
+#include <vector>
+
+
+struct common_speculative_prefill_result {
+    std::vector<int32_t> kept_indices;      // sorted original token positions kept for target model
+    std::vector<float>   importance_scores; // per-token aggregated importance scores (size = n_prompt)
+    int64_t              t_draft_eval_us = 0;
+    int64_t              t_estimate_us   = 0;
+    int32_t              n_prompt_orig   = 0;
+    int32_t              n_prompt_kept   = 0;
+};
+
+// calculate 1D average-pooled smoothed token importance from raw attention matrices
+std::vector<float> common_speculative_prefill_compute_importance(
+    const std::vector<std::vector<float>> & attn_layers,
+    int32_t n_prompt,
+    int32_t n_heads,
+    int32_t pool_kernel_size);
+
+// select kept token indices using chunk-based top-k or token-level ranking
+std::vector<int32_t> common_speculative_prefill_select_indices(
+    const std::vector<float> & importance_scores,
+    const common_params_speculative_prefill & params);
+
+// execute draft prefill, lookahead decoding, attention collection, and index selection
+common_speculative_prefill_result common_speculative_prefill_execute(
+    llama_context * ctx_dft,
+    common_sampler * smpl_dft,
+    const std::vector<llama_token> & prompt,
+    llama_seq_id seq_id,
+    const common_params_speculative_prefill & params);
diff --git a/docs/speculative-prefill.md b/docs/speculative-prefill.md
new file mode 100644
index 000000000000..e7eb9d571c15
--- /dev/null
+++ b/docs/speculative-prefill.md
@@ -0,0 +1,102 @@
+# Speculative prefill
+
+Port of upstream PR [ggml-org/llama.cpp#27692](https://github.com/ggml-org/llama.cpp/pull/27692),
+which implements *Speculative Prefill: Turbocharging TTFT with Lightweight and Training-Free
+Token Importance Estimation* ([arXiv:2502.02789](https://arxiv.org/abs/2502.02789)).
+
+A small draft model prefills the prompt and decodes a few lookahead tokens. The softmax
+attention of those tokens over the prompt is captured with a ggml eval callback, smoothed,
+max-reduced over heads and layers, and pooled into chunks. Only the top scoring chunks are
+fed to the target model. This is lossy: tokens that are dropped are gone.
+
+Ported: `common/speculative-prefill.{h,cpp}`, `llama_set_eval_callback`, the `--spec-prefill-*`
+options, `examples/speculative-prefill`, and the llama-server hookup. Not ported: the
+llama-bench changes and the python eval scripts from that PR.
+
+## Usage
+
+```bash
+llama-server -m target.gguf -mpd draft.gguf --spec-prefill-percentage 0.3
+```
+
+The estimator must be a separate small model with the same vocab as the target. MTP, DFlash2,
+DSpark and Eagle3 heads are target-dependent and the server refuses to reuse them for this.
+The draft context is forced to flash-attn off, because `kq_soft_max` does not exist on the
+flash-attn path. If no attention is captured the full prompt is kept.
+
+## Results
+
+Ryzen AI MAX+ 395 / Radeon 8060S, Vulkan, 119 GiB unified memory. Estimator is
+Qwen3.5-2B-UD-Q4_K_XL in every run. "eff PP" is original prompt tokens divided by TTFT.
+
+### TTFT vs context length
+
+Qwen3.8-27B-UD-Q4_K_XL, no decode speculation, median of 3 runs. A passcode was planted at
+mid-depth and retrieved in 3/3 runs in every cell.
+
+| N tokens | p = off | p = 0.30 | p = 0.15 |
+| ---: | ---: | ---: | ---: |
+|  1,498 |  6,225 ms | 2,647 ms (2.35x) |  1,516 ms (4.11x) |
+|  3,131 | 12,992 ms | 5,199 ms (2.50x) |  3,084 ms (4.21x) |
+|  6,212 | 26,494 ms | 9,906 ms (2.67x) |  6,113 ms (4.33x) |
+| 12,347 | 54,927 ms | 19,557 ms (2.81x) | 11,896 ms (4.62x) |
+
+Baseline eff PP is nearly flat over that range (241 -> 225 tok/s), so prefill here is close to
+linear in N rather than quadratic. The speedup still grows with N, but because the fixed draft
+cost amortizes, not because there is super-linear cost to remove.
+
+Decode is unaffected: 11.14 to 11.62 tok/s across all 12 cells.
+
+### With decode speculation
+
+Qwen3.8-27B-UD-Q4_K_XL, 1,869-token prompt, single run each.
+
+| decode spec | p | TTFT | eff PP | tg |
+| --- | ---: | ---: | ---: | ---: |
+| none        | off  | 5,604 ms |   334 tok/s | 11.56 tok/s |
+| none        | 0.15 | 1,587 ms | 1,180 tok/s | 11.66 tok/s |
+| MTP n=4     | off  | 5,755 ms |   325 tok/s | 19.47 tok/s |
+| MTP n=4     | 0.30 | 2,440 ms |   768 tok/s | 29.33 tok/s |
+| MTP n=4     | 0.15 | 1,636 ms | 1,145 tok/s | 23.61 tok/s |
+| DFlash2 n=7 | off  | 6,576 ms |   284 tok/s | 19.33 tok/s |
+| DFlash2 n=7 | 0.30 | 2,724 ms |   688 tok/s | 30.44 tok/s |
+| DFlash2 n=7 | 0.15 | 1,772 ms | 1,057 tok/s | 20.90 tok/s |
+
+The two mechanisms are independent. DFlash2 gains the most because its head also has to prefill
+the full prompt, and speculative prefill shrinks that too.
+
+### Qwen3.8-Flash-Next
+
+FP4 ple16 (87 GiB) with the Q8_0 MTP head, 1,869-token prompt, median of 3, needle 3/3.
+
+| p | TTFT | eff PP | tg | MTP accept |
+| ---: | ---: | ---: | ---: | ---: |
+| off  | 4,920 ms |   380 tok/s | 29.38 tok/s | 71% |
+| 0.30 | 2,441 ms |   767 tok/s | 28.53 tok/s | 66% |
+| 0.15 | 1,645 ms | 1,139 tok/s | 28.87 tok/s | 69% |
+
+Flash-Next gains less than the dense 27B because it is hybrid: most layers are linear-attention,
+so baseline prefill is already closer to linear in N. MTP acceptance and decode rate hold across
+keep ratios.
+
+The IQ4_XS build reaches 4,718 -> 1,599 ms over the same sweep when run with `--ngram-on-disk`.
+
+## Caveats
+
+- Greedy (temperature 0) output is not reproducible on this box even with no speculation and no
+  speculative prefill: 3 identical requests to the plain target gave 2 different completions.
+  Only TTFT is worth quoting from these runs. Cross-arm `tg` and draft-acceptance differences
+  are confounded, because each arm generates different text under `ignore_eos`.
+- `--spec-prefill-ctx`, `--spec-prefill-device` and `--spec-prefill-ngl` set `enabled = true` as
+  a side effect, so passing only a context size makes the server exit with "speculative prefill
+  enabled but no draft model was provided".
+- The example renumbers kept tokens to positions 0..K-1 instead of keeping their original
+  positions. The reference implementation keeps the original ones.
+- `t_draft_eval_us` covers the draft prompt prefill only, not the lookahead decode steps. About
+  150 ms of the reported TTFT is unattributed, because the eval callback forces node-by-node
+  execution.
+- Quality falls off when the token budget cannot hold all the non-redundant evidence. p = 0.30 is
+  a reasonable default; p = 0.15 still retrieved a needle and summed 3 facts spread across the
+  context; below p = 0.10 is risky for prompts with distributed evidence.
+
+`scripts/bench-spec-prefill.sh` reproduces the server measurements.
diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
index 39f802d250e1..f023b4bfb24a 100644
--- a/examples/CMakeLists.txt
+++ b/examples/CMakeLists.txt
@@ -30,6 +30,7 @@ else()
     add_subdirectory(simple)
     add_subdirectory(simple-chat)
     add_subdirectory(speculative)
+    add_subdirectory(speculative-prefill)
     add_subdirectory(speculative-simple)
     add_subdirectory(gen-docs)
     add_subdirectory(training)
diff --git a/examples/speculative-prefill/CMakeLists.txt b/examples/speculative-prefill/CMakeLists.txt
new file mode 100644
index 000000000000..9bfdec3148b9
--- /dev/null
+++ b/examples/speculative-prefill/CMakeLists.txt
@@ -0,0 +1,5 @@
+set(TARGET llama-speculative-prefill)
+add_executable(${TARGET} speculative-prefill.cpp)
+install(TARGETS ${TARGET} RUNTIME)
+target_link_libraries(${TARGET} PRIVATE llama-common llama ${CMAKE_THREAD_LIBS_INIT})
+target_compile_features(${TARGET} PRIVATE cxx_std_17)
diff --git a/examples/speculative-prefill/speculative-prefill.cpp b/examples/speculative-prefill/speculative-prefill.cpp
new file mode 100644
index 000000000000..1a7ae6560464
--- /dev/null
+++ b/examples/speculative-prefill/speculative-prefill.cpp
@@ -0,0 +1,270 @@
+#include "arg.h"
+#include "common.h"
+#include "log.h"
+#include "llama.h"
+#include "sampling.h"
+#include "speculative.h"
+#include "speculative-prefill.h"
+#include "../../src/llama-ext.h"
+
+#include <algorithm>
+#include <clocale>
+#include <cstdio>
+#include <cstring>
+#include <cinttypes>
+#include <string>
+#include <vector>
+
+#define SPEC_VOCAB_MAX_SIZE_DIFFERENCE 128
+
+int main(int argc, char ** argv) {
+    std::setlocale(LC_NUMERIC, "C");
+
+    common_params params;
+
+    common_init();
+
+    if (!common_params_parse(argc, argv, params, LLAMA_EXAMPLE_SPECULATIVE)) {
+        return 1;
+    }
+
+    common_params_model dft_model = params.speculative.prefill.model;
+    if (dft_model.empty()) {
+        dft_model = params.speculative.draft.mparams;
+    }
+
+    if (dft_model.empty()) {
+        LOG_ERR("%s: draft model is required for speculative prefill (specify with -mpd, --spec-prefill-model, or -md)\n", __func__);
+        return 1;
+    }
+
+    // force speculative prefill enabled
+    params.speculative.prefill.enabled = true;
+
+    llama_backend_init();
+    llama_numa_init(params.numa);
+
+    // load target model
+    LOG_INF("%s: loading target model...\n", __func__);
+    auto init_tgt = common_init_from_params(params);
+    if (!init_tgt) {
+        LOG_ERR("%s: failed to load target model\n", __func__);
+        return 1;
+    }
+
+    llama_model * model_tgt = init_tgt->model();
+    llama_context * ctx_tgt = init_tgt->context();
+
+    if (llama_model_is_recurrent(model_tgt)) {
+        LOG_ERR("%s: speculative prefill is not supported for recurrent models\n", __func__);
+        return 1;
+    }
+
+    // load draft model with standard attention to allow attention extraction
+    LOG_INF("%s: loading draft model...\n", __func__);
+    common_params params_dft = common_base_params_to_speculative(params);
+    params_dft.model = dft_model;
+    if (params.speculative.prefill.n_ctx > 0) {
+        params_dft.n_ctx = params.speculative.prefill.n_ctx;
+    }
+    if (params.speculative.prefill.n_gpu_layers != -1) {
+        params_dft.n_gpu_layers = params.speculative.prefill.n_gpu_layers;
+    }
+    if (!params.speculative.prefill.devices.empty()) {
+        params_dft.devices = params.speculative.prefill.devices;
+    } else if (!params.speculative.draft.devices.empty()) {
+        params_dft.devices = params.speculative.draft.devices;
+    }
+    params_dft.flash_attn_type = LLAMA_FLASH_ATTN_TYPE_DISABLED; // standard attention needed to capture kq_soft_max weights
+
+    auto init_dft = common_init_from_params(params_dft, /*model_only=*/true);
+    if (!init_dft) {
+        LOG_ERR("%s: failed to load draft model\n", __func__);
+        return 1;
+    }
+
+    llama_model * model_dft = init_dft->model();
+    if (!model_dft) {
+        LOG_ERR("%s: failed to load draft model\n", __func__);
+        return 1;
+    }
+
+    if (llama_model_is_recurrent(model_dft)) {
+        LOG_ERR("%s: draft model is recurrent and not supported for speculative prefill\n", __func__);
+        return 1;
+    }
+
+    if (llama_model_target_layer_ids_n(model_dft) > 0) {
+        LOG_ERR("%s: draft model '%s' is target-dependent and cannot be used for speculative prefill\n", __func__, dft_model.path.c_str());
+        return 1;
+    }
+
+    if (params.speculative.prefill.n_ctx <= 0 && params_dft.n_ctx > (int32_t) llama_model_n_ctx_train(model_dft)) {
+        params_dft.n_ctx = llama_model_n_ctx_train(model_dft);
+        LOG_INF("%s: capping speculative prefill draft context to training limit (%d tokens)\n", __func__, params_dft.n_ctx);
+    }
+
+    llama_context_params cparams_dft = common_context_params_to_llama(params_dft);
+    llama_context_ptr ctx_dft_own(llama_init_from_model(model_dft, cparams_dft));
+    llama_context * ctx_dft = ctx_dft_own.get();
+    if (!ctx_dft) {
+        LOG_ERR("%s: failed to create draft context\n", __func__);
+        return 1;
+    }
+
+    const llama_vocab * vocab_tgt = llama_model_get_vocab(model_tgt);
+    const llama_vocab * vocab_dft = llama_model_get_vocab(model_dft);
+
+    if (llama_vocab_get_add_bos(vocab_tgt) != llama_vocab_get_add_bos(vocab_dft) ||
+        (llama_vocab_get_add_bos(vocab_tgt) && llama_vocab_bos(vocab_tgt) != llama_vocab_bos(vocab_dft))) {
+        LOG_ERR("%s: draft model bos tokens must match target model. add: %d - %d, id: %d - %d\n",
+                __func__,
+                llama_vocab_get_add_bos(vocab_tgt), llama_vocab_get_add_bos(vocab_dft),
+                llama_vocab_bos(vocab_tgt), llama_vocab_bos(vocab_dft));
+        return 1;
+    }
+
+    {
+        const int n_vocab_tgt = llama_vocab_n_tokens(vocab_tgt);
+        const int n_vocab_dft = llama_vocab_n_tokens(vocab_dft);
+        const int vocab_diff  = n_vocab_tgt > n_vocab_dft ? n_vocab_tgt - n_vocab_dft : n_vocab_dft - n_vocab_tgt;
+        if (vocab_diff > SPEC_VOCAB_MAX_SIZE_DIFFERENCE) {
+            LOG_ERR("%s: target vocab size %d does not match draft vocab size %d - difference %d, max allowed %d\n",
+                    __func__, n_vocab_tgt, n_vocab_dft, vocab_diff, SPEC_VOCAB_MAX_SIZE_DIFFERENCE);
+            return 1;
+        }
+    }
+
+    // tokenize prompt
+    std::vector<llama_token> prompt_tokens = common_tokenize(ctx_tgt, params.prompt, true, true);
+    const int32_t n_prompt = (int32_t) prompt_tokens.size();
+
+    if (n_prompt == 0) {
+        LOG_ERR("%s: empty prompt provided\n", __func__);
+        return 1;
+    }
+
+    LOG_INF("%s: prompt tokens = %d\n", __func__, n_prompt);
+    LOG_INF("%s: speculative prefill config: percentage = %.2f, chunk_size = %d, lookahead = %d, pool_kernel = %d\n",
+            __func__,
+            (double) params.speculative.prefill.percentage,
+            params.speculative.prefill.chunk_size,
+            params.speculative.prefill.look_ahead_cnt,
+            params.speculative.prefill.pool_kernel_size);
+
+    llama_seq_id seq_id = 0;
+
+    // initialize draft sampler for lookahead steps
+    common_params_sampling sparams_dft = params.sampling;
+    sparams_dft.temp = 0.0f;
+    common_sampler_ptr smpl_dft(common_sampler_init(model_dft, sparams_dft));
+
+    // step 1: run speculative prefill on draft model
+    const auto t_spec_prefill_start = ggml_time_us();
+
+    common_speculative_prefill_result spec_res = common_speculative_prefill_execute(
+        ctx_dft,
+        smpl_dft.get(),
+        prompt_tokens,
+        seq_id,
+        params.speculative.prefill);
+
+    LOG_INF("%s: speculative prefill kept %d / %d tokens (%.1f%%)\n",
+            __func__,
+            spec_res.n_prompt_kept,
+            spec_res.n_prompt_orig,
+            100.0f * (float) spec_res.n_prompt_kept / (float) spec_res.n_prompt_orig);
+    LOG_INF("%s: draft eval time: %.2f ms, importance estimation time: %.2f ms\n",
+            __func__,
+            spec_res.t_draft_eval_us / 1000.0f,
+            spec_res.t_estimate_us / 1000.0f);
+
+    // step 2: sparse target model prefill
+    const auto t_tgt_prefill_start = ggml_time_us();
+
+    const int32_t n_batch_tgt = llama_n_batch(ctx_tgt);
+    const int32_t n_kept_total = (int32_t) spec_res.kept_indices.size();
+    llama_batch batch_tgt = llama_batch_init(std::min(n_kept_total, n_batch_tgt), 0, 1);
+
+    int ret = 0;
+    for (int32_t i = 0; i < n_kept_total; i += n_batch_tgt) {
+        const int32_t n_eval = std::min(n_kept_total - i, n_batch_tgt);
+        common_batch_clear(batch_tgt);
+
+        for (int32_t j = 0; j < n_eval; ++j) {
+            const int32_t k = i + j;
+            const int32_t orig_idx = spec_res.kept_indices[k];
+            const bool is_last = (k == n_kept_total - 1);
+            common_batch_add(batch_tgt, prompt_tokens[orig_idx], (llama_pos) k, { seq_id }, is_last);
+        }
+
+        ret = llama_decode(ctx_tgt, batch_tgt);
+        if (ret != 0) {
+            LOG_ERR("%s: failed to decode sparse prompt on target model, ret = %d\n", __func__, ret);
+            llama_batch_free(batch_tgt);
+            return 1;
+        }
+    }
+    llama_batch_free(batch_tgt);
+
+    llama_synchronize(ctx_tgt);
+    llama_synchronize(ctx_dft);
+
+    const auto t_tgt_prefill_end = ggml_time_us();
+    const double ttft_ms = (t_tgt_prefill_end - t_spec_prefill_start) / 1000.0;
+    const double tgt_prefill_ms = (t_tgt_prefill_end - t_tgt_prefill_start) / 1000.0;
+
+    LOG_INF("%s: sparse target prefill time = %.2f ms\n", __func__, tgt_prefill_ms);
+    LOG_INF("%s: total Time-To-First-Token (TTFT) = %.2f ms (%.2f effective prompt tokens/s)\n",
+            __func__, ttft_ms, ttft_ms > 0.0 ? (1000.0 * (double) n_prompt / ttft_ms) : 0.0);
+
+    // step 3: autoregressive generation
+    common_sampler_ptr smpl_tgt(common_sampler_init(model_tgt, params.sampling));
+
+    LOG("\n--- Generation Start ---\n");
+
+    int32_t n_predict = params.n_predict > 0 ? params.n_predict : 32;
+    int32_t cur_pos = n_kept_total;
+    int32_t n_generated = 0;
+
+    llama_batch batch_gen = llama_batch_init(1, 0, 1);
+
+    const auto t_gen_start = ggml_time_us();
+
+    for (int32_t i = 0; i < n_predict; ++i) {
+        const llama_token token_id = common_sampler_sample(smpl_tgt.get(), ctx_tgt, -1);
+        common_sampler_accept(smpl_tgt.get(), token_id, true);
+
+        if (llama_vocab_is_eog(vocab_tgt, token_id)) {
+            break;
+        }
+
+        const std::string piece = common_token_to_piece(ctx_tgt, token_id);
+        LOG("%s", piece.c_str());
+        fflush(stdout);
+
+        common_batch_clear(batch_gen);
+        common_batch_add(batch_gen, token_id, (llama_pos) cur_pos++, { seq_id }, true);
+
+        ret = llama_decode(ctx_tgt, batch_gen);
+        if (ret != 0) {
+            LOG_ERR("%s: failed to decode generated token %d, ret = %d\n", __func__, i, ret);
+            break;
+        }
+
+        n_generated++;
+    }
+
+    const auto t_gen_end = ggml_time_us();
+    const double gen_ms = (t_gen_end - t_gen_start) / 1000.0;
+
+    LOG("\n--- Generation End ---\n\n");
+
+    LOG_INF("generated %d tokens in %.2f ms (%.2f tokens/s)\n",
+            n_generated, gen_ms, n_generated > 0 ? (n_generated / (gen_ms / 1000.0)) : 0.0);
+
+    llama_batch_free(batch_gen);
+    llama_backend_free();
+
+    return 0;
+}
diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp
index fec7d7c92bf3..8f40fb2f7bc3 100644
--- a/ggml/src/ggml-backend.cpp
+++ b/ggml/src/ggml-backend.cpp
@@ -70,6 +70,7 @@ size_t ggml_backend_buft_get_alloc_size(ggml_backend_buffer_type_t buft, const s
         // if you hit this assert, update ggml_backend_op_alloc_size_may_expand() accordingly
         GGML_ASSERT(size <= ggml_nbytes(tensor) ||
                     ggml_op_is_empty(tensor->op) ||
+                    ggml_is_quantized(tensor->type) || // [TAG_ALLOC_SIZE_EXPAND]
                     ggml_backend_op_alloc_size_may_expand(tensor->op));
 
         return size;
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 7c2d3cad20a8..a6fc655c41c9 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -915,6 +915,7 @@ static size_t ggml_backend_cuda_buffer_type_get_alloc_size(ggml_backend_buffer_t
         : ggml_nbytes(tensor);
     int64_t ne0 = tensor->ne[0];
 
+    // [TAG_ALLOC_SIZE_EXPAND]
     if (ggml_is_quantized(tensor->type)) {
         if (ne0 % MATRIX_ROW_PADDING != 0) {
             GGML_ASSERT(tensor->nb[0] == ggml_element_size(tensor));
diff --git a/ggml/src/ggml-hexagon/htp/cpy-ops.c b/ggml/src/ggml-hexagon/htp/cpy-ops.c
index 15bc8dc244fb..c945425dab32 100644
--- a/ggml/src/ggml-hexagon/htp/cpy-ops.c
+++ b/ggml/src/ggml-hexagon/htp/cpy-ops.c
@@ -330,7 +330,7 @@ int op_cpy(struct htp_ops_context * octx) {
     }
 
     const struct htp_tensor *sync = octx->src[1];
-    if (sync) {
+    if (sync && (sync->flags & HTP_TENSOR_FENCE)) {
         if (!use_dma) {
             // htp_tensor_flush_all(octx->ctx, octx->dsts, 1);
             qurt_mem_cache_clean((qurt_addr_t) 0, 0, QURT_MEM_CACHE_FLUSH_INVALIDATE_ALL, QURT_MEM_DCACHE);
diff --git a/ggml/src/ggml-metal/ggml-metal-tuning.cpp b/ggml/src/ggml-metal/ggml-metal-tuning.cpp
index 9114ab7424de..4f26fd9c3733 100644
--- a/ggml/src/ggml-metal/ggml-metal-tuning.cpp
+++ b/ggml/src/ggml-metal/ggml-metal-tuning.cpp
@@ -516,6 +516,119 @@ constexpr fa_vec_entry_t fa_vec_tuned_table[] = {
     { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_0, 512, 512, 3, 3 }, { 2, 2 } },
     { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_0, 576, 512, -1, 0 }, { 1, 4 } },
     { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_0, 576, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 32, 32, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 32, 32, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 32, 32, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 32, 32, 2, 3 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 32, 32, 3, 2 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 32, 32, 3, 3 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 64, 64, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 64, 64, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 96, 96, 1, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 96, 96, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 96, 96, 2, 3 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 96, 96, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 96, 96, 3, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 128, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 128, 128, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 128, 128, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 128, 128, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 128, 128, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 192, 192, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 192, 192, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 192, 192, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 192, 192, 1, 4 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 192, 192, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 192, 192, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 192, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 192, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 256, 256, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 320, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 320, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 512, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 512, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 576, 512, 1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 576, 512, 1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 576, 512, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 576, 512, 1, 3 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q4_1, 576, 512, 1, 4 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 32, 32, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 32, 32, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 32, 32, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 32, 32, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 64, 64, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 64, 64, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 64, 64, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 64, 64, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 96, 96, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 96, 96, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 96, 96, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 96, 96, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 96, 96, 3, 4 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 128, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 128, 128, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 128, 128, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 128, 128, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 128, 128, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 192, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 192, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 192, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 192, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 128, -1, 1 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 128, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 128, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 128, 1, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 128, 1, 4 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 128, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 192, 128, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 256, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 256, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 320, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 320, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 512, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 512, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 576, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_0, 576, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 32, 32, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 32, 32, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 32, 32, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 32, 32, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 64, 64, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 64, 64, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 64, 64, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 64, 64, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 96, 96, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 96, 96, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 96, 96, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 96, 96, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 128, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 128, 128, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 128, 128, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 192, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 192, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 192, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 192, 1, 4 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 192, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 192, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 128, -1, 1 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 128, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 128, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 128, 1, 4 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 128, 2, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 192, 128, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 256, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 256, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 320, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 320, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 512, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 512, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 576, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q5_1, 576, 512, -1, 1 }, { 1, 4 } },
     { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 2, 4 } },
     { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q8_0, 32, 32, 1, 2 }, { 1, 4 } },
     { { GGML_METAL_DEVICE_M2, GGML_TYPE_Q8_0, 32, 32, 2, 2 }, { 1, 4 } },
diff --git a/ggml/src/ggml-rpc/ggml-rpc.cpp b/ggml/src/ggml-rpc/ggml-rpc.cpp
index 58a8a030cfa5..a97db24e624e 100644
--- a/ggml/src/ggml-rpc/ggml-rpc.cpp
+++ b/ggml/src/ggml-rpc/ggml-rpc.cpp
@@ -625,7 +625,7 @@ static bool ggml_backend_buffer_is_rpc(ggml_backend_buffer_t buffer) {
     return buffer->iface.free_buffer == ggml_backend_rpc_buffer_free_buffer;
 }
 
-static rpc_tensor serialize_tensor(const ggml_tensor * tensor) {
+static rpc_tensor serialize_tensor(const ggml_tensor * tensor, const std::shared_ptr<rpc_dispatcher> & dispatcher = nullptr) {
     rpc_tensor result;
     if (!tensor) {
         memset(&result, 0, sizeof(result));
@@ -637,8 +637,14 @@ static rpc_tensor serialize_tensor(const ggml_tensor * tensor) {
     if (tensor->buffer && ggml_backend_buffer_is_rpc(tensor->buffer)) {
         ggml_backend_buffer_t buffer = tensor->buffer;
         ggml_backend_rpc_buffer_context * ctx = (ggml_backend_rpc_buffer_context *)buffer->context;
-        result.buffer = ctx != nullptr ? ctx->remote_ptr : 0;
-        result.data = reinterpret_cast<uint64_t>(tensor->data);
+        // ref: https://github.com/ggml-org/llama.cpp/pull/26500
+        if (ctx != nullptr && (dispatcher == nullptr || ctx->dispatcher == dispatcher)) {
+            result.buffer = ctx->remote_ptr;
+            result.data = reinterpret_cast<uint64_t>(tensor->data);
+        } else {
+            result.buffer = 0;
+            result.data = 0;
+        }
     } else {
         result.buffer = 0;
         result.data   = 0;
@@ -958,7 +964,7 @@ static void ggml_backend_rpc_synchronize(ggml_backend_t backend) {
     rpc_ctx->dispatcher->synchronize();
 }
 
-static void add_tensor(ggml_tensor * tensor, const ggml_cgraph * cgraph, std::vector<rpc_tensor> & tensors, std::unordered_set<ggml_tensor*> & visited) {
+static void add_tensor(ggml_tensor * tensor, const ggml_cgraph * cgraph, const std::shared_ptr<rpc_dispatcher> & dispatcher, std::vector<rpc_tensor> & tensors, std::unordered_set<ggml_tensor*> & visited) {
     if (tensor == nullptr) {
         return;
     }
@@ -967,10 +973,10 @@ static void add_tensor(ggml_tensor * tensor, const ggml_cgraph * cgraph, std::ve
     }
     visited.insert(tensor);
     for (int i = 0; i < GGML_MAX_SRC; i++) {
-        add_tensor(tensor->src[i], cgraph, tensors, visited);
+        add_tensor(tensor->src[i], cgraph, dispatcher, tensors, visited);
     }
-    add_tensor(tensor->view_src, cgraph, tensors, visited);
-    rpc_tensor result = serialize_tensor(tensor);
+    add_tensor(tensor->view_src, cgraph, dispatcher, tensors, visited);
+    rpc_tensor result = serialize_tensor(tensor, dispatcher);
     const size_t hash_pos = ggml_hash_find(&cgraph->visited_hash_set, tensor);
     if (hash_pos != GGML_HASHSET_FULL && ggml_bitset_get(cgraph->visited_hash_set.used, hash_pos)) {
         result.use_count = cgraph->use_counts[hash_pos];
@@ -978,12 +984,12 @@ static void add_tensor(ggml_tensor * tensor, const ggml_cgraph * cgraph, std::ve
     tensors.push_back(result);
 }
 
-static uint8_t * serialize_graph(uint32_t device, const ggml_cgraph * cgraph, size_t * output_size) {
+static uint8_t * serialize_graph(uint32_t device, const ggml_cgraph * cgraph, const std::shared_ptr<rpc_dispatcher> & dispatcher, size_t * output_size) {
     uint32_t n_nodes = cgraph->n_nodes;
     std::vector<rpc_tensor> tensors;
     std::unordered_set<ggml_tensor*> visited;
     for (uint32_t i = 0; i < n_nodes; i++) {
-        add_tensor(cgraph->nodes[i], cgraph, tensors, visited);
+        add_tensor(cgraph->nodes[i], cgraph, dispatcher, tensors, visited);
     }
     // serialization format:
     // | device (4 bytes) | n_nodes (4 bytes) | nodes (n_nodes * sizeof(uint64_t) | n_tensors (4 bytes) | tensors (n_tensors * sizeof(rpc_tensor)) |
@@ -1020,7 +1026,7 @@ static enum ggml_status ggml_backend_rpc_graph_compute(ggml_backend_t backend, g
     } else {
         rpc_dev_ctx->last_graph_uid = cgraph->uid;
         size_t input_size = 0;
-        uint8_t * input = serialize_graph(rpc_ctx->device, cgraph, &input_size);
+        uint8_t * input = serialize_graph(rpc_ctx->device, cgraph, rpc_ctx->dispatcher, &input_size);
         std::shared_ptr<uint8_t> input_ptr(input, std::default_delete<uint8_t[]>());
         rpc_ctx->dispatcher->send_async(RPC_CMD_GRAPH_COMPUTE, input_ptr, input_size);
     }
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
index a4f442e8e736..de8e2673067b 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
@@ -1057,6 +1057,7 @@ struct vk_device_struct {
     vk_pipeline pipeline_argsort_f32[num_argsort_pipelines];
     vk_pipeline pipeline_argsort_large_f32[num_argsort_pipelines];
     vk_pipeline pipeline_topk_f32[num_topk_pipelines];
+    vk_pipeline pipeline_topk_radix_f32;
     vk_pipeline pipeline_sum_rows_f32;
     vk_pipeline pipeline_cross_entropy_loss_f32, pipeline_cross_entropy_loss_f32_wg512;
     vk_pipeline pipeline_cross_entropy_loss_back_f32, pipeline_cross_entropy_loss_back_f32_wg512;
@@ -3979,6 +3980,46 @@ static std::vector<uint32_t> get_fa_spec_constants(const vk_fa_pipeline_state& s
     };
 }
 
+// SHMEM_STRIDE_PAD for mul_mm.comp (constantID=12), in FLOAT_TYPEV2 units.
+//
+// mul_mm.comp stages the shared A/B tiles at SHMEM_STRIDE = BK/2 + pad. Each element is one
+// dword, so the stride maps 1:1 onto the 32 LDS banks and gcd(SHMEM_STRIDE, 32) is the conflict
+// factor. With BK=32 the shader default of 4 gives stride 20 (gcd 4, four-way conflict). Pad 2
+// gives stride 18 (gcd 2, 16 banks) and is a large win on the quantised path on gfx1151.
+//
+// But the coopmat path passes SHMEM_STRIDE to coopMatLoad as its Stride operand, and
+// VUID-RuntimeSpirv-OpCooperativeMatrixLoadKHR-08986 needs a 16 byte aligned stride for the
+// 16x16 f16 tiles. Stride bytes are (BK/2 + pad) * 4, so only pad % 4 == 0 is in contract.
+// RADV before 25.3 lowers coopMatLoad to ds_read_b128, which the contract entitles it to, and
+// the misaligned rows then pay runtime splits: pp512 drops by more than 2x. RADV 25.3 and later
+// lowers to ds_read_b64, which the 72 byte stride always suits, so the extra bank spread wins.
+// Use pad 2 only where it is measured to win.
+//
+// The float path keeps pad 4 on measurement, not theory: the f32/f16 shaders #define BK 32
+// whatever the host pushes, so they run the same stride, yet pad 2 measures worse on both. The
+// host warptile's BK labels the pipeline even where the shader overrides the value.
+//
+// GGML_VK_SHMEM_PAD=N overrides both paths, for probing.
+static uint32_t ggml_vk_coopmat_shmem_pad(const vk_device& device, uint32_t bk) {
+    if (device->vendor_id == VK_VENDOR_ID_INTEL && device->coopmat_support &&
+        device->driver_id == vk::DriverId::eIntelProprietaryWindows) {
+        return 0;
+    }
+    static const int env_pad = [] {
+        const char * e = getenv("GGML_VK_SHMEM_PAD");
+        return e ? atoi(e) : -1;
+    }();
+    if (env_pad >= 0) {
+        return (uint32_t) env_pad;
+    }
+    if (device->coopmat_support && bk >= 32 &&
+        device->driver_id == vk::DriverId::eMesaRadv &&
+        device->properties.driverVersion >= VK_MAKE_API_VERSION(0, 25, 3, 0)) {
+        return 2;
+    }
+    return 4;  // mul_mm.comp default
+}
+
 static bool ggml_vk_matmul_shmem_support(const vk_device& device, const std::vector<uint32_t>& warptile, bool mul_mat_id, ggml_type src0_type) {
 
     uint32_t lut_size = 0;
@@ -4019,9 +4060,9 @@ static bool ggml_vk_matmul_shmem_support(const vk_device& device, const std::vec
 
     // Needs to be kept up to date on shader changes
     // Needs to stay aligned with ggml_vk_mul_mm_spec.
-    const bool intel_shmem_stride_pad_zero = device->vendor_id == VK_VENDOR_ID_INTEL && device->coopmat_support &&
-                                              device->driver_id == vk::DriverId::eIntelProprietaryWindows;
-    const uint32_t bank_conflict_offset = intel_shmem_stride_pad_zero ? 0 : (device->coopmat_support ? 8 : 1);
+    // The shared tiles are (BK/2 + pad) FLOAT_TYPEV2 per row, i.e. (BK + 2*pad)
+    // elements, so the byte-space offset below is twice the shader's pad.
+    const uint32_t bank_conflict_offset = device->coopmat_support ? 2 * ggml_vk_coopmat_shmem_pad(device, warptile[3]) : 1;
     const uint32_t type_size = device->fp16 ? sizeof(ggml_fp16_t) : sizeof(float);
     const uint32_t warps = warptile[0] / warptile[10];
 
@@ -4646,10 +4687,12 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
 
     auto const &ggml_vk_mul_mm_spec = [&device](std::vector<uint32_t> spec, bool aligned) {
         spec.push_back(aligned ? 1u : 0u);  // constantID=11: ALIGNED
-        if (device->vendor_id == VK_VENDOR_ID_INTEL && device->coopmat_support &&
-            device->driver_id == vk::DriverId::eIntelProprietaryWindows) {
-            spec.push_back(0u);  // constantID=12: SHMEM_STRIDE_PAD = 0
-            spec.push_back(1u);  // constantID=13: APPLY_SLM_A_RESHAPE = true
+        const uint32_t pad = ggml_vk_coopmat_shmem_pad(device, spec[3]);  // spec[3] is the warptile BK
+        const bool intel_slm = device->vendor_id == VK_VENDOR_ID_INTEL && device->coopmat_support &&
+                               device->driver_id == vk::DriverId::eIntelProprietaryWindows;
+        if (intel_slm || pad != 4) {
+            spec.push_back(pad);                  // constantID=12: SHMEM_STRIDE_PAD
+            spec.push_back(intel_slm ? 1u : 0u);  // constantID=13: APPLY_SLM_A_RESHAPE
         }
         return spec;
     };
@@ -5814,6 +5857,13 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
         }
     }
 
+    // radix selection for k beyond the shared-memory shaders' single-workgroup reach:
+    // one workgroup per row, single dispatch, no scratch buffers
+    {
+        const uint32_t RADIX_WG = std::min<uint32_t>(1024, 1u << device->max_workgroup_size_log2);
+        ggml_vk_create_pipeline2(device, device->pipeline_topk_radix_f32, "topk_radix_f32", topk_radix_f32_len, topk_radix_f32_data, "main", 2, sizeof(vk_op_topk_push_constants), {RADIX_WG, 1, 1}, {RADIX_WG}, 1, true);
+    }
+
     ggml_vk_create_pipeline(device, device->pipeline_argmax_f32, "argmax_f32", argmax_f32_len, argmax_f32_data, "main", 2, sizeof(vk_op_push_constants), {1, 1, 1}, { device->subgroup_size }, 1);
 
     ggml_vk_create_pipeline(device, device->pipeline_sum_rows_f32, "sum_rows_f32", sum_rows_f32_len, sum_rows_f32_data, "main", 2, sizeof(vk_op_sum_rows_push_constants), {1, 1, 1}, { device->subgroup_size }, 1);
@@ -8677,7 +8727,13 @@ static void ggml_vk_buffer_read_2d(vk_buffer& src, size_t offset, void * dst, si
     // If the device is not an UMA device the memory is host-accessible through rebar. While writing
     // through PCIe is sufficient fast reading back data from PCIe is slower than going through
     // the HW device to host copy path.
-    if(src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible && src->device->uma) {
+    // On UMA a direct CPU read is only fast from a host-cached mapping. A write-combined mapping
+    // (host-visible without HOST_CACHED, which is what amdgpu gives for GTT) reads back at
+    // uncached speed, so bulk reads must use the device copy path. Small reads stay direct to
+    // avoid the fence round trip.
+    const bool host_cached = bool(src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostCached);
+    if((src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible) && src->device->uma &&
+       (host_cached || width * height <= 64 * 1024)) {
         GGML_ASSERT(src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostCoherent);
 
         std::lock_guard<std::recursive_mutex> guard(src->device->mutex);
@@ -13971,7 +14027,7 @@ static void ggml_vk_argsort(ggml_backend_vk_context * ctx, vk_context& subctx, c
     std::array<uint32_t, 3> elements;
 
     elements[0] = ncolsp2;
-    elements[1] = std::min((uint32_t)ggml_nrows(src0), ctx->device->properties.limits.maxComputeWorkGroupCount[1]);
+    elements[1] = std::min(nrows, ctx->device->properties.limits.maxComputeWorkGroupCount[1]);
     elements[2] = 1;
 
     // First dispatch initializes tmp_idx and does the first N passes where
@@ -14015,7 +14071,10 @@ static void ggml_vk_argsort(ggml_backend_vk_context * ctx, vk_context& subctx, c
     }
 }
 
-static void ggml_vk_topk(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
+// Tournament-reduction top_k, only correct while k fits a single workgroup's
+// tournament pipeline (see ggml_vk_topk's dispatcher and the large-k path below
+// for why this can't just be given a bigger workgroup).
+static void ggml_vk_topk_small(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
     uint32_t ncols = src0->ne[0];
     uint32_t nrows = ggml_nrows(src0);
     uint32_t k = dst->ne[0];
@@ -14123,6 +14182,48 @@ static void ggml_vk_topk(ggml_backend_vk_context * ctx, vk_context& subctx, cons
     ctx->prealloc_x_need_sync = true;
 }
 
+// Large-k top_k: the tournament reduction above can only discard elements a
+// workgroup can see all of at once, so it can never make progress once k is
+// >= the max deployable workgroup size (~1024 on real hardware) - any element
+// in a smaller chunk could belong to the true global top-k. Radix selection
+// discards on a *global* bucket boundary instead, so it converges regardless
+// of the workgroup-size-vs-k relationship: one workgroup per row, a single
+// dispatch, no scratch buffers.
+static void ggml_vk_topk_large(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
+    const uint32_t ncols = (uint32_t) src0->ne[0];
+    const uint32_t nrows = (uint32_t) ggml_nrows(src0);
+    const uint32_t k     = (uint32_t) dst->ne[0];
+
+    vk_pipeline pipeline = ctx->device->pipeline_topk_radix_f32;
+    GGML_ASSERT(pipeline != nullptr);
+
+    vk_op_topk_push_constants pc { ncols, ncols, ncols, k, nrows, 0, 0 };
+
+    vk_subbuffer src0_buf = ggml_vk_tensor_subbuffer(ctx, src0);
+    vk_subbuffer dst_buf  = ggml_vk_tensor_subbuffer(ctx, dst);
+
+    // one workgroup per row
+    std::array<uint32_t, 3> elements = { nrows*pipeline->wg_denoms[0], 1, 1 };
+
+    ggml_pipeline_request_descriptor_sets(ctx, pipeline, 1);
+    ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { src0_buf, dst_buf }, pc, elements);
+}
+
+// true when k fits a single-workgroup tournament pipeline; the index is only
+// valid to probe after the range check, hence the explicit ordering here
+static bool ggml_vk_topk_k_fits_pipeline(const vk_device_struct * device, uint32_t k) {
+    const uint32_t min_pipeline = (uint32_t) log2f(float(k)) + 1;
+    return min_pipeline < num_topk_pipelines && device->pipeline_topk_f32[min_pipeline] != nullptr;
+}
+
+static void ggml_vk_topk(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
+    if (ggml_vk_topk_k_fits_pipeline(ctx->device.get(), (uint32_t) dst->ne[0])) {
+        ggml_vk_topk_small(ctx, subctx, src0, dst);
+    } else {
+        ggml_vk_topk_large(ctx, subctx, src0, dst);
+    }
+}
+
 static void ggml_vk_sum(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, ggml_tensor * dst) {
     vk_op_sum_rows_push_constants p = vk_op_sum_rows_push_constants_init(src0, dst, ggml_nelements(src0));
     ggml_vk_op_f32(ctx, subctx, src0, nullptr, nullptr, nullptr, dst, GGML_OP_SUM, p);
@@ -18803,15 +18904,13 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                 if (!ggml_is_contiguous(op) || !ggml_is_contiguous(op->src[0])) {
                     return false;
                 }
-                // We could potentially support larger, using argsort to sort the
-                // whole thing. Not clear if this is needed.
-                uint32_t min_pipeline = (uint32_t)log2f(float(op->ne[0])) + 1;
-                if (min_pipeline >= num_topk_pipelines ||
-                    !device->pipeline_topk_f32[min_pipeline]) {
-                    return false;
-                }
+                // op is the dst tensor here, so op->ne[0] == k (ggml_top_k builds dst
+                // as ggml_new_tensor_4d(..., k, a->ne[1], ...)), not the candidate pool
+                // width - ggml_vk_topk_small's multi-pass loop already handles arbitrarily
+                // large ncols correctly as long as k fits a single tournament pipeline.
+                // Beyond that, the radix path takes over, which has no extra requirements.
+                return true;
             }
-            return true;
         case GGML_OP_UPSCALE:
             if (op->op_params[0] & GGML_SCALE_FLAG_ANTIALIAS) {
                 if ((op->op_params[0] & 0xFF) != GGML_SCALE_MODE_BILINEAR) {
diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_iq3_s.comp b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_iq3_s.comp
index 5cdf2a89d0fd..42f52b4a1273 100644
--- a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_iq3_s.comp
+++ b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mat_vec_iq3_s.comp
@@ -7,7 +7,14 @@ layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
 
 FLOAT_TYPE temp[NUM_COLS][NUM_ROWS];
 
-void calc_superblock(const uint a_offset, const uint b_offset, const uint ib32, const uint i, const uint num_blocks_per_row, const uint first_row, const uint num_rows) {
+// invocations per superblock. with many columns, 8 invocations need too many
+// registers and spill, so use 16 to halve the per-invocation B working set
+const uint TPB = NUM_COLS <= 4 ? 8 : 16;
+const uint NL  = 32 / TPB; // l steps per invocation
+
+void calc_superblock(const uint a_offset, const uint b_offset, const uint itid, const uint i, const uint num_blocks_per_row, const uint first_row, const uint num_rows) {
+    const uint ib32 = itid / (TPB / 8);
+    const uint l0   = (itid % (TPB / 8)) * NL;
     const uint y_idx = i * QUANT_K + 32 * ib32;
 
     uint ibi = a_offset + first_row * num_blocks_per_row + i;
@@ -16,11 +23,8 @@ void calc_superblock(const uint a_offset, const uint b_offset, const uint ib32,
         const uint scale = (data_a[ibi].scales[ib32/2] >> (4 * (ib32 & 1))) & 0xF;
         const float dscale = d * (1 + 2 * scale);
         const uint qh = data_a[ibi].qh[ib32];
-        FLOAT_TYPE sum[NUM_COLS];
-        [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
-            sum[j] = 0.0;
-        }
-        [[unroll]] for (uint l = 0; l < 4; ++l) {
+        [[unroll]] for (uint ll = 0; ll < NL; ++ll) {
+            const uint l = l0 + ll;
             const u8vec2 qs = unpack8(uint32_t(data_a_packed16[ibi].qs[4 * ib32 + l])).xy; // vec4 used due to #12147
             const uint sign = data_a[ibi].signs[4 * ib32 + l];
             const vec4 grid0 = vec4(unpack8(iq3s_grid[qs.x | ((qh << (8 - 2*l)) & 0x100)]));
@@ -30,7 +34,7 @@ void calc_superblock(const uint a_offset, const uint b_offset, const uint ib32,
                 const vec4 b0 = vec4(data_b_v4[(j*p.batch_stride_b + b_offset + y_idx) / 4 + 2*l + 0]);
                 const vec4 b4 = vec4(data_b_v4[(j*p.batch_stride_b + b_offset + y_idx) / 4 + 2*l + 1]);
 
-                sum[j] =
+                const FLOAT_TYPE sum =
                       fma(FLOAT_TYPE(b0.x), FLOAT_TYPE((sign &   1) != 0 ? -grid0.x : grid0.x),
                       fma(FLOAT_TYPE(b0.y), FLOAT_TYPE((sign &   2) != 0 ? -grid0.y : grid0.y),
                       fma(FLOAT_TYPE(b0.z), FLOAT_TYPE((sign &   4) != 0 ? -grid0.z : grid0.z),
@@ -39,12 +43,11 @@ void calc_superblock(const uint a_offset, const uint b_offset, const uint ib32,
                       fma(FLOAT_TYPE(b4.y), FLOAT_TYPE((sign &  32) != 0 ? -grid1.y : grid1.y),
                       fma(FLOAT_TYPE(b4.z), FLOAT_TYPE((sign &  64) != 0 ? -grid1.z : grid1.z),
                       fma(FLOAT_TYPE(b4.w), FLOAT_TYPE((sign & 128) != 0 ? -grid1.w : grid1.w),
-                      sum[j]))))))));
+                      FLOAT_TYPE(0.0)))))))));
+
+                temp[j][n] = fma(dscale, sum, temp[j][n]);
             }
         }
-        [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
-            temp[j][n] = fma(dscale, sum[j], temp[j][n]);
-        }
         ibi += num_blocks_per_row;
     }
 }
@@ -55,11 +58,11 @@ void compute_outputs(const uint32_t first_row, const uint32_t num_rows) {
 
     const uint num_blocks_per_row = p.ncols / QUANT_K;
 
-    // 8 threads are used to process each block
-    const uint blocks_per_wg = gl_WorkGroupSize.x/8;
+    // TPB invocations are used to process each block
+    const uint blocks_per_wg = gl_WorkGroupSize.x/TPB;
     const uint tid = gl_LocalInvocationID.x;
-    const uint itid = tid % 8;  // 0...7
-    const uint ix = tid / 8;
+    const uint itid = tid % TPB;
+    const uint ix = tid / TPB;
 
     [[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
         [[unroll]] for (uint i = 0; i < NUM_ROWS; ++i) {
diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/topk_radix.comp b/ggml/src/ggml-vulkan/vulkan-shaders/topk_radix.comp
new file mode 100644
index 000000000000..abb88fe4e1e5
--- /dev/null
+++ b/ggml/src/ggml-vulkan/vulkan-shaders/topk_radix.comp
@@ -0,0 +1,120 @@
+#version 450
+
+#extension GL_EXT_control_flow_attributes : enable
+
+#include "types.glsl"
+
+// Exact top-k selection for rows too wide (or k too large) for the shared-memory
+// nary-search shader: iterative 8-bit radix selection on the order-preserving
+// unsigned transform of the keys, one workgroup per row, no temporary buffers.
+// The resulting indices are in no particular order (the ggml_top_k contract).
+// Ported to Vulkan from the HIP kernel in EngramHalo.cpp (Aristo94).
+
+layout(constant_id = 0) const int BLOCK_SIZE = 1024;
+
+layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
+
+layout (binding = 0) readonly  buffer A {float data_a[];};
+layout (binding = 1) writeonly buffer D {int data_d[];};
+
+layout (push_constant) uniform parameter {
+    uint orig_ncols;
+    uint ncols_input;
+    uint ncols_output;
+    uint k;
+    uint nrows;
+    uint first_pass;
+    uint last_pass;
+} p;
+
+shared uint hist[256];
+shared uint sh_prefix;
+shared uint sh_rank;
+shared uint sh_greater;
+shared uint sh_equal;
+
+// map float bits to uint such that unsigned comparison matches float comparison
+uint f2ui(float x) {
+    uint y = floatBitsToUint(x);
+    if ((y & 0x80000000u) != 0u) {
+        y ^= ~0u;
+    } else {
+        y |= 0x80000000u;
+    }
+    return y;
+}
+
+void main() {
+    const uint row  = gl_WorkGroupID.x;
+    const uint tid  = gl_LocalInvocationID.x;
+    const uint wgs  = gl_WorkGroupSize.x;
+    const uint base = row*p.ncols_input;
+
+    if (row >= p.nrows) {
+        return;
+    }
+
+    // find the exact key of the k-th largest element, 8 bits per pass
+    uint prefix = 0u;
+    uint pmask  = 0u;
+    uint rank   = p.k;
+
+    for (int shift = 24; shift >= 0; shift -= 8) {
+        // strided: the workgroup can be narrower than the 256 bins (Vulkan only
+        // guarantees maxComputeWorkGroupInvocations >= 128)
+        for (uint bin = tid; bin < 256u; bin += wgs) {
+            hist[bin] = 0u;
+        }
+        barrier();
+
+        // count only the elements that still match the prefix found so far
+        for (uint col = tid; col < p.ncols_input; col += wgs) {
+            const uint key = f2ui(data_a[base + col]);
+            if ((key & pmask) == prefix) {
+                atomicAdd(hist[(key >> shift) & 255u], 1u);
+            }
+        }
+        barrier();
+
+        // walk the bins from the largest keys down to the one holding the k-th largest
+        if (tid == 0u) {
+            uint r   = rank;
+            int  bin = 255;
+            while (bin > 0 && hist[bin] < r) {
+                r -= hist[bin];
+                bin--;
+            }
+            sh_prefix = prefix | (uint(bin) << shift);
+            sh_rank   = r;
+        }
+        barrier();
+
+        prefix = sh_prefix;
+        rank   = sh_rank;
+        pmask |= 255u << shift;
+    }
+
+    // gather: all elements strictly greater than the threshold key are in the top k,
+    // ties on the exact threshold fill the remaining `rank` slots
+    if (tid == 0u) {
+        sh_greater = 0u;
+        sh_equal   = 0u;
+    }
+    barrier();
+
+    const uint threshold = prefix;
+    const uint out_base  = row*p.k;
+
+    for (uint col = tid; col < p.ncols_input; col += wgs) {
+        const uint key = f2ui(data_a[base + col]);
+        if (key > threshold) {
+            const uint pos = atomicAdd(sh_greater, 1u);
+            data_d[out_base + pos] = int(col);
+        } else if (key == threshold) {
+            const uint pos = atomicAdd(sh_equal, 1u);
+            if (pos < rank) {
+                data_d[out_base + (p.k - rank) + pos] = int(col);
+            }
+        }
+    }
+}
diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
index f0610f4cd82d..3624fb51a43f 100644
--- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
+++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
@@ -1029,6 +1029,8 @@ void process_shaders() {
     string_to_spv("topk_argsort_f32", "topk_argsort.comp", {{"A_TYPE", "float"}});
     string_to_spv("topk_nary_search_f32", "topk_nary_search.comp", {{"A_TYPE", "float"}});
 
+    string_to_spv("topk_radix_f32", "topk_radix.comp", {{"A_TYPE", "float"}});
+
     string_to_spv("argmax_f32", "argmax.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "int"}}));
     string_to_spv("sum_rows_f32", "sum_rows.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));
     string_to_spv("cross_entropy_loss_f32", "cross_entropy_loss.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"B_TYPE", "float"}, {"D_TYPE", "float"}}));
diff --git a/include/llama.h b/include/llama.h
index c2ef101f6c44..d0dda88ffd2e 100644
--- a/include/llama.h
+++ b/include/llama.h
@@ -1028,6 +1028,9 @@ extern "C" {
     // Set abort callback
     LLAMA_API void llama_set_abort_callback(struct llama_context * ctx, ggml_abort_callback abort_callback, void * abort_callback_data);
 
+    // Set evaluation callback
+    LLAMA_API void llama_set_eval_callback(struct llama_context * ctx, ggml_backend_sched_eval_callback cb_eval, void * cb_eval_user_data);
+
     // Wait until all computations are finished
     // This is automatically done when using one of the functions below to obtain the computation results
     // and is not necessary to call it explicitly in most cases
diff --git a/scripts/bench-spec-prefill.sh b/scripts/bench-spec-prefill.sh
new file mode 100755
index 000000000000..faad2c62b179
--- /dev/null
+++ b/scripts/bench-spec-prefill.sh
@@ -0,0 +1,63 @@
+#!/bin/bash
+# Measure TTFT and decode rate for speculative prefill against a llama-server arm.
+#
+#   scripts/bench-spec-prefill.sh <label> <prompt-file> -m target.gguf [server args...]
+#
+# Runs one warmup plus 3 timed completions with prompt caching off, then reports the median
+# TTFT. Pass -mpd/--spec-prefill-percentage in the server args to measure a speculative
+# prefill arm, or leave them out for the baseline.
+
+set -u
+
+LABEL="$1"; shift
+PROMPT_FILE="$1"; shift
+
+PORT="${PORT:-8899}"
+NEEDLE="${NEEDLE:-}"
+N_PREDICT="${N_PREDICT:-128}"
+BIN="${BIN:-./build/bin/llama-server}"
+LOG="/tmp/bench-spf-$LABEL.log"
+
+"$BIN" --port "$PORT" --no-webui "$@" > "$LOG" 2>&1 &
+SRV=$!
+trap 'kill $SRV 2>/dev/null' EXIT
+
+for _ in $(seq 1 600); do
+    curl -sf "http://127.0.0.1:$PORT/health" > /dev/null 2>&1 && break
+    kill -0 $SRV 2>/dev/null || { echo "$LABEL: server died"; tail -8 "$LOG"; exit 1; }
+    sleep 2
+done
+
+PROMPT=$(python3 -c "import json,sys;print(json.dumps(open(sys.argv[1]).read()))" "$PROMPT_FILE")
+REQ="{\"prompt\":$PROMPT,\"n_predict\":$N_PREDICT,\"temperature\":0,\"cache_prompt\":false,\"ignore_eos\":true}"
+
+curl -s "http://127.0.0.1:$PORT/completion" -H 'Content-Type: application/json' -d "$REQ" > /dev/null
+
+for _ in 1 2 3; do
+    curl -s "http://127.0.0.1:$PORT/completion" -H 'Content-Type: application/json' -d "$REQ"
+    echo
+done | LABEL="$LABEL" NEEDLE="$NEEDLE" python3 -c "
+import sys, json, os
+label  = os.environ['LABEL']
+needle = os.environ.get('NEEDLE', '')
+ttft = []; tg = []; n = 0; hits = 0; runs = 0; acc = 0; drafted = 0
+for line in sys.stdin:
+    line = line.strip()
+    if not line:
+        continue
+    d = json.loads(line); t = d['timings']
+    ttft.append(t['prompt_ms']); tg.append(t['predicted_per_second']); n = t['prompt_n']
+    acc += t.get('draft_n_accepted', 0); drafted += t.get('draft_n', 0)
+    runs += 1
+    hits += bool(needle) and needle in d['content']
+ttft.sort()
+med = ttft[len(ttft)//2]
+out = f'{label}  N={n}  ttft_med={med:.0f}ms  eff_pp={1000*n/med:.1f} tok/s  tg_avg={sum(tg)/len(tg):.2f} tok/s'
+if drafted:
+    out += f'  accept={100*acc/drafted:.0f}%'
+if needle:
+    out += f'  needle={hits}/{runs}'
+print(out)
+"
+
+grep -oE "kept [0-9]+ / [0-9]+ tokens \([0-9.]+%\)" "$LOG" | sort -u
diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index ca387021d221..835c98c8128d 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -1158,6 +1158,17 @@ void llama_context::set_abort_callback(bool (*abort_callback)(void * data), void
     }
 }
 
+void llama_context::set_eval_callback(ggml_backend_sched_eval_callback cb_eval, void * cb_eval_user_data) {
+    LLAMA_LOG_DEBUG("%s: call\n", __func__);
+
+    cparams.cb_eval           = cb_eval;
+    cparams.cb_eval_user_data = cb_eval_user_data;
+
+    if (sched) {
+        ggml_backend_sched_set_eval_callback(sched.get(), cb_eval, cb_eval_user_data);
+    }
+}
+
 void llama_context::set_embeddings(bool value) {
     LLAMA_LOG_DEBUG("%s: value = %d\n", __func__, value);
 
@@ -3811,6 +3822,14 @@ void llama_set_abort_callback(llama_context * ctx, bool (*abort_callback)(void *
     ctx->set_abort_callback(abort_callback, abort_callback_data);
 }
 
+void llama_set_eval_callback(llama_context * ctx, ggml_backend_sched_eval_callback cb_eval, void * cb_eval_user_data) {
+    if (!ctx) {
+        return;
+    }
+
+    ctx->set_eval_callback(cb_eval, cb_eval_user_data);
+}
+
 void llama_set_embeddings(llama_context * ctx, bool embeddings) {
     ctx->set_embeddings(embeddings);
 }
diff --git a/src/llama-context.h b/src/llama-context.h
index bf91daa8b562..b089e1267aaf 100644
--- a/src/llama-context.h
+++ b/src/llama-context.h
@@ -111,6 +111,7 @@ struct llama_context {
     void set_n_threads(int32_t n_threads, int32_t n_threads_batch);
 
     void set_abort_callback(bool (*abort_callback)(void * data), void * abort_callback_data);
+    void set_eval_callback(ggml_backend_sched_eval_callback cb_eval, void * cb_eval_user_data);
 
     void set_embeddings (bool value);
     void set_embeddings_nextn(bool value, bool masked);
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 996b88db296e..516c8094739e 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -6172,16 +6172,17 @@ struct test_top_k : public test_case {
     const std::array<int64_t, 4> ne;
     const int k;
     const bool ties;
+    const bool neg_inf;
     ggml_tensor * input {};
 
     std::string vars() override {
-        return VARS_TO_STR4(type, ne, k, ties);
+        return VARS_TO_STR5(type, ne, k, ties, neg_inf);
     }
 
     test_top_k(ggml_type type = GGML_TYPE_F32,
             std::array<int64_t, 4> ne = {16, 10, 10, 10},
-            int k = 4, bool ties = false)
-        : type(type), ne(ne), k(k), ties(ties) {}
+            int k = 4, bool ties = false, bool neg_inf = false)
+        : type(type), ne(ne), k(k), ties(ties), neg_inf(neg_inf) {}
 
     double max_err() override {
         return 0.0;
@@ -6281,6 +6282,14 @@ struct test_top_k : public test_case {
                     }
                 }
                 std::shuffle(data.begin(), data.end(), rng);
+                if (neg_inf) {
+                    // keep exactly k finite candidates; the rest are masked (-INFINITY),
+                    // mirroring qwen4exp's indexer score tensor where invisible cells are
+                    // -INFINITY and must never be selected regardless of magnitude
+                    for (int64_t i = k; i < t->ne[0]; i++) {
+                        data[i] = -INFINITY;
+                    }
+                }
                 ggml_backend_tensor_set(t, data.data(), r * t->nb[1], t->ne[0] * sizeof(float));
             }
         }
@@ -9241,6 +9250,10 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
             //test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 18,  i, 32*256, { 1,  1}, {8, 1}));
             //test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 19,  i, 33*256, { 1,  1}, {1, 1}));
         }
+        // mat-vec shaders split k across lanes and loop over the blocks in strides. k must be
+        // long enough that the loop wraps, else the stride is never exercised
+        test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 16,  1, 16*256, { 1,  1}, {1, 1}));
+        test_cases.emplace_back(new test_mul_mat(type_a,    GGML_TYPE_F32, 16,  8, 16*256, { 1,  1}, {1, 1}));
     }
 
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 2880, 32, 2880, {1, 1}, {1, 1}));
@@ -9787,6 +9800,12 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
                 test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {(1<<i), 1, 1, 1}, k));
                 test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {(1<<i) + 11, 1, 2, 1}, k));
                 test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {(1<<i) + 11, 1, 2, 1}, k, true));
+                // -INFINITY sentinels (masked cells) must sort to the bottom and never be
+                // selected - exercises the large-k argsort-then-slice path specifically,
+                // mirroring the DeepSeek-V4-style sparse indexer's masked score tensor
+                if (k == 9999) {
+                    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {(1<<i), 1, 1, 1}, k, false, true));
+                }
             }
         }
     }
@@ -9800,6 +9819,8 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
             test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {202048, nrows, 1, 1}, k, true));
         }
     }
+    // sparse-attention indexer shape: ncols tracks a growing KV cache, k is far past one workgroup
+    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {32768, 1, 1, 1}, 2051, false, true));
 
     for (int k : {1, 2, 3, 7, 15}) {
         test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {16, 10, 10, 10}, k));
diff --git a/tests/test-rpc-multi-server.cpp b/tests/test-rpc-multi-server.cpp
new file mode 100644
index 000000000000..4502e2ce71fa
--- /dev/null
+++ b/tests/test-rpc-multi-server.cpp
@@ -0,0 +1,47 @@
+#include "ggml-alloc.h"
+#include "ggml-backend.h"
+#include "ggml-impl.h"
+#include "ggml-rpc.h"
+#include "ggml.h"
+
+int main(int argc, char ** argv) {
+    GGML_ASSERT(argc == 3);
+    ggml_backend_load_all();
+
+    const char * endpoint_a = argv[1];
+    const char * endpoint_b = argv[2];
+
+    ggml_backend_t backend_a = ggml_backend_rpc_init(endpoint_a, 0);
+    ggml_backend_t backend_b = ggml_backend_rpc_init(endpoint_b, 0);
+    GGML_ASSERT(backend_a != nullptr);
+    GGML_ASSERT(backend_b != nullptr);
+
+    ggml_init_params params = {
+        /* .mem_size   = */ ggml_tensor_overhead() + ggml_graph_overhead_custom(1, false),
+        /* .mem_buffer = */ nullptr,
+        /* .no_alloc   = */ true,
+    };
+    ggml_context * ctx = ggml_init(params);
+    GGML_ASSERT(ctx != nullptr);
+
+    ggml_tensor * tensor = ggml_new_tensor_1d(ctx, GGML_TYPE_F32, 1);
+    ggml_backend_buffer_t buffer = ggml_backend_alloc_ctx_tensors(ctx, backend_a);
+    GGML_ASSERT(buffer != nullptr);
+
+    // A remote pointer allocated by server A is not meaningful to server B.
+    ggml_cgraph * graph = ggml_new_graph_custom(ctx, 1, false);
+    graph->nodes[0] = tensor;
+    graph->n_nodes = 1;
+
+    GGML_ASSERT(ggml_backend_graph_compute(backend_b, graph) == GGML_STATUS_SUCCESS);
+    // Wait for server B to finish the graph before the script checks its log.
+    size_t free_mem;
+    size_t total_mem;
+    ggml_backend_rpc_get_device_memory(endpoint_b, 0, &free_mem, &total_mem);
+    GGML_ASSERT(total_mem > 0);
+    ggml_backend_buffer_free(buffer);
+    ggml_free(ctx);
+    ggml_backend_free(backend_b);
+    ggml_backend_free(backend_a);
+    return 0;
+}
diff --git a/tests/test-rpc-multi-server.sh b/tests/test-rpc-multi-server.sh
new file mode 100755
index 000000000000..a8c72316202a
--- /dev/null
+++ b/tests/test-rpc-multi-server.sh
@@ -0,0 +1,43 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+server=$1
+client=$2
+port_a=$((40000 + $$ % 10000))
+port_b=$((port_a + 1))
+endpoint_a="127.0.0.1:${port_a}"
+endpoint_b="127.0.0.1:${port_b}"
+test_dir=$(mktemp -d)
+
+cleanup() {
+    kill "${pid_a:-}" "${pid_b:-}" 2>/dev/null || true
+    rm -rf "$test_dir"
+}
+trap cleanup EXIT
+
+wait_for_port() {
+    local port=$1
+    for _ in {1..600}; do
+        if (exec 3<>"/dev/tcp/127.0.0.1/$port") 2>/dev/null; then
+            exec 3>&-
+            exec 3<&-
+            return 0
+        fi
+        sleep 0.05
+    done
+    return 1
+}
+
+"$server" --device CPU --host 127.0.0.1 --port "$port_a" >"$test_dir/server-a.log" 2>&1 &
+pid_a=$!
+"$server" --device CPU --host 127.0.0.1 --port "$port_b" >"$test_dir/server-b.log" 2>&1 &
+pid_b=$!
+wait_for_port "$port_a"
+wait_for_port "$port_b"
+
+"$client" "$endpoint_a" "$endpoint_b"
+
+if grep -q "invalid data ptr" "$test_dir/server-b.log"; then
+    cat "$test_dir/server-b.log"
+    exit 1
+fi
diff --git a/tools/rpc/CMakeLists.txt b/tools/rpc/CMakeLists.txt
index 0eee9a922e77..2891c7d034cf 100644
--- a/tools/rpc/CMakeLists.txt
+++ b/tools/rpc/CMakeLists.txt
@@ -3,6 +3,18 @@ add_executable(${TARGET} rpc-server.cpp)
 target_link_libraries(${TARGET} PRIVATE ggml)
 target_compile_features(${TARGET} PRIVATE cxx_std_17)
 
+if (LLAMA_BUILD_TESTS AND UNIX AND NOT GGML_BACKEND_DL)
+    add_executable(test-rpc-multi-server ${PROJECT_SOURCE_DIR}/tests/test-rpc-multi-server.cpp)
+    target_link_libraries(test-rpc-multi-server PRIVATE ggml ggml-rpc)
+    target_include_directories(test-rpc-multi-server PRIVATE ${PROJECT_SOURCE_DIR}/ggml/src)
+    add_test(
+        NAME test-rpc-multi-server
+        COMMAND bash ${PROJECT_SOURCE_DIR}/tests/test-rpc-multi-server.sh
+            $<TARGET_FILE:ggml-rpc-server>
+            $<TARGET_FILE:test-rpc-multi-server>)
+    set_property(TEST test-rpc-multi-server PROPERTY LABELS main)
+endif()
+
 if(LLAMA_TOOLS_INSTALL)
     install(TARGETS ${TARGET} RUNTIME)
 endif()
diff --git a/tools/server/README.md b/tools/server/README.md
index 10c6a0e796c4..a35ab5974e36 100644
--- a/tools/server/README.md
+++ b/tools/server/README.md
@@ -1787,6 +1787,15 @@ We also offer additional options that are exclusive to presets (these aren't tre
 - `load-on-startup` (boolean): Controls whether the model loads automatically when the server starts. Only applies at startup: if the model list is reloaded later (for example after editing the preset file), a newly added model is listed but not loaded
 - `stop-timeout` (int, seconds): After requested unload, wait for this many seconds before forcing termination (default: 10)
 - `dedup-cache-models` (boolean): When the preset uses `hf-repo` pointing to a model that is already downloaded, hide the corresponding cached model entry from `GET /models` (the preset entry remains visible). Set it in the `[*]` section to apply to all presets.
+- `hidden` (boolean): Omit this model from `GET /models` and `GET /v1/models`. The model can still be loaded and used by requesting it by name, so this only affects listings such as the model picker in the web UI. Useful for entries that are not meant to be picked directly, e.g. a draft/MTP GGUF that is discovered from the HF cache:
+
+  ```ini
+  ; hides the auto-discovered cache entry, which is only useful as a draft model
+  [ggml-org/MY-MODEL-MTP-GGUF:Q4_K_M]
+  hidden = 1
+  ```
+
+  A section whose name matches an existing model merges into that model's entry rather than creating a new one, so a stub section with just `hidden = 1` is enough to hide a model that came from the HF cache or from `--models-dir`.
 
 ### Routing requests
 
diff --git a/tools/server/server-common.cpp b/tools/server/server-common.cpp
index 19c0e94ebe95..4650d0a10490 100644
--- a/tools/server/server-common.cpp
+++ b/tools/server/server-common.cpp
@@ -546,7 +546,7 @@ void server_tokens::insert(const llama_tokens & inp_tokens) {
 }
 
 const llama_tokens & server_tokens::get_tokens() const {
-    GGML_ASSERT(!has_mtmd);
+    GGML_ASSERT(!has_media());
     return tokens;
 }
 
@@ -629,7 +629,7 @@ llama_tokens server_tokens::get_text_tokens() const {
 }
 
 void server_tokens::set_token(llama_pos pos, llama_token id) {
-    GGML_ASSERT(!has_mtmd); // only allow this if mtmd is disabled
+    GGML_ASSERT(!has_media());
     tokens[pos] = id;
 }
 
diff --git a/tools/server/server-common.h b/tools/server/server-common.h
index ddfc544cb9a0..8dce0976381b 100644
--- a/tools/server/server-common.h
+++ b/tools/server/server-common.h
@@ -220,6 +220,8 @@ struct server_tokens {
 
     bool empty() const { return tokens.empty(); }
 
+    bool has_media() const { return !map_idx_to_media.empty(); }
+
     void clear() {
         map_idx_to_media.clear();
         tokens.clear();
diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp
index 90501bcb7f73..81228c18cc79 100644
--- a/tools/server/server-context.cpp
+++ b/tools/server/server-context.cpp
@@ -14,6 +14,8 @@
 #include "log.h"
 #include "sampling.h"
 #include "speculative.h"
+#include "speculative-prefill.h"
+#include "src/llama-ext.h"
 #include "mtmd.h"
 #include "mtmd-helper.h"
 
@@ -241,6 +243,7 @@ struct server_slot {
 
     llama_context * ctx_tgt = nullptr;
     llama_context * ctx_dft = nullptr;
+    llama_context * ctx_spf = nullptr;
 
     common_memory mem;
 
@@ -258,6 +261,9 @@ struct server_slot {
     bool spec_is_replay = false;
     std::mt19937 spec_synth_rng;
 
+    bool spec_prefill_active = false;
+    server_tokens spec_prefill_tokens;
+
     // TODO: move members that belong to the task (such as `generated_text`, `has_new_line`) to task_results_state
     //       see https://github.com/ggml-org/llama.cpp/pull/18283#issuecomment-3710175837
     std::unique_ptr<const server_task> task;
@@ -335,6 +341,9 @@ struct server_slot {
         SLT_TRC(*this, "clearing prompt with %zu tokens\n", prompt.tokens.size());
 
         mem.seq_rm(id, -1, -1);
+        if (ctx_spf) {
+            llama_memory_seq_rm(llama_get_memory(ctx_spf), id, -1, -1);
+        }
 
         prompt.clear();
     }
@@ -404,6 +413,9 @@ struct server_slot {
 
         // clear multimodal state
         mbatch.reset();
+
+        spec_prefill_active = false;
+        spec_prefill_tokens.clear();
     }
 
     void init_sampler() const {
@@ -471,6 +483,14 @@ struct server_slot {
         return !!spec;
     }
 
+    const server_tokens & prompt_src() const {
+        return spec_prefill_active ? spec_prefill_tokens : task->tokens;
+    }
+
+    int32_t n_prompt_src() const {
+        return spec_prefill_active ? (int32_t) spec_prefill_tokens.size() : task->n_tokens();
+    }
+
     void add_token(const completion_token_output & token) {
         if (!is_processing()) {
             SLT_WRN(*this, "%s", "slot is not processing\n");
@@ -508,10 +528,12 @@ struct server_slot {
             // no speculative decoding
             i_batch = batch.size();
 
+            llama_pos pos = prompt.tokens.pos_next();
+
             if (!inp_embd.empty()) {
-                add_ok &= batch.add(id, inp_embd, prompt.tokens.pos_next(), true, false);
+                add_ok &= batch.add(id, inp_embd, pos, true, false);
             } else {
-                add_ok &= batch.add(id, sampled, prompt.tokens.pos_next(), true, false);
+                add_ok &= batch.add(id, sampled, pos, true, false);
             }
 
             SLT_DBG(*this, "slot decode token, id=%d, n_ctx = %d, n_tokens = %d, truncated = %d\n",
@@ -527,7 +549,7 @@ struct server_slot {
                 spec_i_batch.push_back(batch.size() + i + 1);
             }
 
-            auto pos0 = prompt.tokens.pos_next();
+            llama_pos pos0 = prompt.tokens.pos_next();
 
             add_ok &= batch.add(id, sampled, pos0++, true, false);
             for (auto token : spec_draft) {
@@ -551,6 +573,10 @@ struct server_slot {
 
             state = SLOT_STATE_IDLE;
 
+            if (ctx_spf) {
+                llama_memory_seq_rm(llama_get_memory(ctx_spf), id, -1, -1);
+            }
+
             // do not keep context of the child slots - the parent's context is enough
             if (task->is_child()) {
                 prompt_clear();
@@ -622,11 +648,18 @@ struct server_slot {
             return;
         }
 
-        const double n_prompt_second = stats.n_prompt_tps();
-        const double f_progress = task->n_tokens() > 0 ? (double) prompt.n_tokens() / task->n_tokens() : 0.0;
+        const double f_progress = n_prompt_src() > 0 ? (double) prompt.n_tokens() / n_prompt_src() : 0.0;
+        int n_tokens_disp = (int) stats.n_prompt_processed;
+        double n_prompt_second = stats.n_prompt_tps();
+
+        if (spec_prefill_active && task && n_prompt_src() > 0) {
+            const int n_uncached = (int) (task->n_tokens() - stats.n_prompt_cached);
+            n_tokens_disp = (int) (f_progress * n_uncached);
+            n_prompt_second = t_prompt_total > 0.0 ? (1e3 / t_prompt_total) * n_tokens_disp : 0.0;
+        }
 
         SLT_INF(*this, "prompt processing, n_tokens = %6d, progress = %.2f, t = %6.2f s / %.2f tokens per second\n",
-                (int) stats.n_prompt_processed, f_progress, t_prompt_total / 1e3, n_prompt_second);
+                n_tokens_disp, f_progress, t_prompt_total / 1e3, n_prompt_second);
     }
 
     void print_timings() const {
@@ -697,7 +730,7 @@ struct server_slot {
 
         if (ptask) {
             res["id_task"] = ptask->id;
-            res["n_prompt_tokens"]           = (int32_t) prompt.tokens.size();
+            res["n_prompt_tokens"]           = ptask->n_tokens();
             res["n_prompt_tokens_processed"] = stats.n_prompt_processed;
             res["n_prompt_tokens_cache"]     = stats.n_prompt_cached;
             res["params"] = ptask->params.to_json(only_metrics);
@@ -725,6 +758,10 @@ struct server_slot {
         mem.seq_rm(other.id,     -1, -1);
         mem.seq_cp(id, other.id, -1, -1);
 
+        if (ctx_spf) {
+            llama_memory_seq_rm(llama_get_memory(ctx_spf), other.id, -1, -1);
+        }
+
         other.i_batch = i_batch;
 
         other.stats = stats;
@@ -890,6 +927,12 @@ struct server_context_impl {
 
     common_speculative_init_result_ptr spec_init;
 
+    common_init_result_ptr llama_init_spf;
+    llama_context_ptr ctx_spf_own;
+    llama_model *   model_spf = nullptr;
+    llama_context * ctx_spf   = nullptr;
+    common_sampler_ptr smpl_spf;
+
     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;
 
@@ -937,6 +980,13 @@ struct server_context_impl {
 
     void destroy() {
         spec.reset();
+
+        smpl_spf.reset();
+        ctx_spf_own.reset();
+        llama_init_spf.reset();
+        ctx_spf   = nullptr;
+        model_spf = nullptr;
+
         spec_init.reset();
 
         ctx_dft   = nullptr;
@@ -1001,6 +1051,53 @@ struct server_context_impl {
         return true;
     }
 
+    void apply_spec_prefill(server_slot & slot) {
+        slot.spec_prefill_active = false;
+        slot.spec_prefill_tokens.clear();
+
+        if (!ctx_spf || !smpl_spf || !params_base.speculative.prefill.enabled) {
+            return;
+        }
+        if (!slot.task || !slot.task->need_logits() || slot.task->tokens.has_media()) {
+            return;
+        }
+
+        const llama_tokens & prompt = slot.task->tokens.get_tokens();
+        if (prompt.empty()) {
+            return;
+        }
+
+        llama_memory_seq_rm(llama_get_memory(ctx_spf), slot.id, -1, -1);
+        common_sampler_reset(smpl_spf.get());
+
+        const common_speculative_prefill_result res = common_speculative_prefill_execute(
+            ctx_spf, smpl_spf.get(), prompt, slot.id, params_base.speculative.prefill);
+
+        llama_memory_seq_rm(llama_get_memory(ctx_spf), slot.id, -1, -1);
+
+        if (res.kept_indices.empty() || (int32_t) res.kept_indices.size() >= (int32_t) prompt.size()) {
+            return;
+        }
+
+        llama_tokens kept;
+        kept.reserve(res.kept_indices.size());
+        for (const int32_t idx : res.kept_indices) {
+            if (idx >= 0 && idx < (int32_t) prompt.size()) {
+                kept.push_back(prompt[idx]);
+            }
+        }
+        if (kept.size() < 2) {
+            return;
+        }
+
+        slot.spec_prefill_tokens = server_tokens(kept, slot.task->tokens.has_mtmd);
+        slot.spec_prefill_active = true;
+
+        SLT_INF(slot, "speculative prefill kept %d / %d tokens (%.1f%%)\n",
+                (int) kept.size(), (int) prompt.size(),
+                100.0f * (float) kept.size() / (float) prompt.size());
+    }
+
     // load the model and initialize llama_context
     // this may also be called to resume from sleeping state
     bool load_model(common_params & params) {
@@ -1150,6 +1247,107 @@ struct server_context_impl {
             load_progress_callback(1.0f, &load_progress_spec);
         }
 
+        if (params_base.speculative.prefill.enabled) {
+            if (llama_model_is_recurrent(model_tgt)) {
+                SRV_ERR("%s", "speculative prefill is not supported for recurrent models\n");
+                return false;
+            }
+
+            const bool has_target_dependent_dft = std::any_of(
+                params_base.speculative.types.begin(),
+                params_base.speculative.types.end(),
+                [](common_speculative_type t) {
+                    return t == COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH ||
+                           t == COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK ||
+                           t == COMMON_SPECULATIVE_TYPE_DRAFT_EAGLE3 ||
+                           t == COMMON_SPECULATIVE_TYPE_DRAFT_MTP;
+                });
+
+            common_params_model spf_model = params_base.speculative.prefill.model;
+            if (spf_model.empty() && !has_target_dependent_dft) {
+                spf_model = params_base.speculative.draft.mparams;
+            }
+
+            if (spf_model.empty()) {
+                SRV_ERR("%s", "speculative prefill enabled but no draft model was provided\n");
+                return false;
+            }
+
+            common_params params_spf = common_base_params_to_speculative(params_base);
+            params_spf.model = spf_model;
+            if (params_base.speculative.prefill.n_ctx > 0) {
+                params_spf.n_ctx = params_base.speculative.prefill.n_ctx;
+            }
+            if (params_base.speculative.prefill.n_gpu_layers != -1) {
+                params_spf.n_gpu_layers = params_base.speculative.prefill.n_gpu_layers;
+            }
+            if (!params_base.speculative.prefill.devices.empty()) {
+                params_spf.devices = params_base.speculative.prefill.devices;
+            }
+            params_spf.flash_attn_type = LLAMA_FLASH_ATTN_TYPE_DISABLED;
+            params_spf.fit_params = false;
+            params_spf.n_parallel = params_base.n_parallel;
+            params_spf.speculative.prefill.enabled = false;
+            params_spf.speculative.draft.mparams = {};
+
+            const bool reuse_dft = model_dft && !has_target_dependent_dft && (
+                params_base.speculative.prefill.model.empty() ||
+                spf_model.path == params_base.speculative.draft.mparams.path);
+
+            if (reuse_dft) {
+                SRV_INF("reusing draft model for speculative prefill '%s'\n", params_spf.model.get_name().c_str());
+                model_spf = model_dft;
+            } else {
+                SRV_INF("loading speculative prefill draft model '%s'\n", params_spf.model.get_name().c_str());
+                llama_init_spf = common_init_from_params(params_spf, /*model_only=*/true);
+                model_spf = llama_init_spf ? llama_init_spf->model() : nullptr;
+            }
+            if (model_spf == nullptr) {
+                SRV_ERR("failed to load speculative prefill draft model, '%s'\n", params_spf.model.path.c_str());
+                return false;
+            }
+
+            if (params_base.speculative.prefill.n_ctx <= 0 && params_spf.n_ctx > (int32_t) llama_model_n_ctx_train(model_spf)) {
+                params_spf.n_ctx = llama_model_n_ctx_train(model_spf);
+                SRV_INF("capping speculative prefill draft context to training limit (%d tokens)\n", params_spf.n_ctx);
+            }
+
+            llama_context_params cparams = common_context_params_to_llama(params_spf);
+            ctx_spf_own.reset(llama_init_from_model(model_spf, cparams));
+            ctx_spf   = ctx_spf_own.get();
+            if (ctx_spf == nullptr) {
+                SRV_ERR("failed to create speculative prefill context for '%s'\n", params_spf.model.path.c_str());
+                return false;
+            }
+
+            if (llama_model_is_recurrent(model_spf)) {
+                SRV_ERR("speculative prefill draft model '%s' is recurrent and not supported\n", params_spf.model.get_name().c_str());
+                return false;
+            }
+
+            if (llama_model_target_layer_ids_n(model_spf) > 0) {
+                SRV_ERR("speculative prefill draft model '%s' is target-dependent and cannot be used for speculative prefill\n", params_spf.model.get_name().c_str());
+                return false;
+            }
+
+            const llama_vocab * vocab_spf = llama_model_get_vocab(model_spf);
+            const int n_vocab_tgt = llama_vocab_n_tokens(vocab);
+            const int n_vocab_spf = llama_vocab_n_tokens(vocab_spf);
+            const int vocab_diff  = n_vocab_tgt > n_vocab_spf ? n_vocab_tgt - n_vocab_spf : n_vocab_spf - n_vocab_tgt;
+            if (vocab_diff > 128) {
+                SRV_ERR("speculative prefill vocab size difference %d exceeds 128\n", vocab_diff);
+                return false;
+            }
+
+            common_params_sampling sparams_spf;
+            sparams_spf.temp = 0.0f;
+            smpl_spf.reset(common_sampler_init(model_spf, sparams_spf));
+            if (!smpl_spf) {
+                SRV_ERR("%s", "failed to init speculative prefill sampler\n");
+                return false;
+            }
+        }
+
         if (has_mmproj) {
             if (callback_state) {
                 callback_state(SERVER_STATE_LOADING, {{"stage", "mmproj_model"}});
@@ -1290,6 +1488,7 @@ struct server_context_impl {
             slot.id      = i;
             slot.ctx_tgt = ctx_tgt;
             slot.ctx_dft = ctx_dft;
+            slot.ctx_spf = ctx_spf;
             slot.mem.init(ctx_tgt, ctx_dft);
             slot.spec    = spec.get();
             slot.n_ctx   = n_ctx_slot();
@@ -3134,19 +3333,23 @@ struct server_context_impl {
 
                 // this slot still has a prompt to be processed
                 if (slot.state == SLOT_STATE_PROCESSING_PROMPT || slot.state == SLOT_STATE_STARTED) {
-                    const auto & input_tokens = slot.task->tokens;
+                    if (slot.state == SLOT_STATE_STARTED) {
+                        slot.stats.update_prompt_start();
+                        apply_spec_prefill(slot);
+                    }
+
+                    const auto & input_tokens = slot.prompt_src();
+                    const int32_t n_prompt_in = slot.n_prompt_src();
 
                     // used to determine the number of tokens added to the batch for the current slot
                     const auto n_tokens_prev = batch.size();
 
                     // TODO: maybe move branch to outside of this loop in the future
                     if (slot.state == SLOT_STATE_STARTED) {
-                        slot.stats.update_prompt_start();
-
                         slot.state = SLOT_STATE_PROCESSING_PROMPT;
 
                         SLT_TRC(slot, "new prompt, n_ctx_slot = %d, n_keep = %d, task.n_tokens = %d\n",
-                                slot.n_ctx, slot.task->params.n_keep, slot.task->n_tokens());
+                                slot.n_ctx, slot.task->params.n_keep, n_prompt_in);
 
                         // print prompt tokens (for debugging)
                         /*if (1) {
@@ -3183,39 +3386,39 @@ struct server_context_impl {
                         }
 
                         if (!slot.can_split()) {
-                            if (slot.task->n_tokens() > n_ubatch) {
+                            if (n_prompt_in > n_ubatch) {
                                 send_error(slot,
                                            string_format(
                                                "input (%d tokens) is too large to process. increase the physical batch "
                                                "size (current batch size: %d)",
-                                               slot.task->n_tokens(), n_ubatch),
+                                               n_prompt_in, n_ubatch),
                                            ERROR_TYPE_SERVER);
                                 slot.release();
                                 return;
                             }
 
-                            if (slot.task->n_tokens() > slot.n_ctx) {
+                            if (n_prompt_in > slot.n_ctx) {
                                 send_error(
                                     slot,
                                     string_format(
                                         "input (%d tokens) is larger than the max context size (%d tokens). skipping",
-                                        slot.task->n_tokens(), slot.n_ctx),
+                                        n_prompt_in, slot.n_ctx),
                                     ERROR_TYPE_EXCEED_CONTEXT_SIZE);
                                 slot.release();
                                 return;
                             }
                         } else {
-                            if (slot.task->n_tokens() >= slot.n_ctx) {
+                            if (n_prompt_in >= slot.n_ctx) {
                                 send_error(slot,
                                            string_format("request (%d tokens) exceeds the available context size (%d "
                                                          "tokens), try increasing it",
-                                                         slot.task->n_tokens(), slot.n_ctx),
+                                                         n_prompt_in, slot.n_ctx),
                                            ERROR_TYPE_EXCEED_CONTEXT_SIZE);
                                 slot.release();
                                 return;
                             }
 
-                            if (slot.task->params.cache_prompt) {
+                            if (slot.task->params.cache_prompt && !slot.spec_prefill_active) {
                                 // reuse any previously computed tokens that are common with the new prompt
                                 n_past = slot.prompt.tokens.get_common_prefix(input_tokens);
 
@@ -3292,7 +3495,7 @@ struct server_context_impl {
                             llama_pos pos_next = slot.prompt.tokens.pos_next(n_past);
 
                             // ref: https://github.com/ggml-org/llama.cpp/pull/24110
-                            const bool has_new_tokens = (n_past < slot.task->n_tokens());
+                            const bool has_new_tokens = (n_past < n_prompt_in);
 
                             // the largest pos_min required for a checkpoint to be useful
                             const auto pos_min_thold = std::max(0, pos_next - n_swa - (has_new_tokens ? 0 : 1));
@@ -3308,7 +3511,7 @@ struct server_context_impl {
                                 // this is useful for debugging prompt caching
                                 if (slots_debug) {
                                     const int np0 = std::max<int>(n_past - slots_n_diff, 0);
-                                    const int np1 = std::min<int>(n_past + slots_n_diff + 2, std::min(slot.prompt.tokens.size(), slot.task->tokens.size()));
+                                    const int np1 = std::min<int>(n_past + slots_n_diff + 2, std::min(slot.prompt.tokens.size(), input_tokens.size()));
 
                                     std::stringstream ss0;
                                     std::stringstream ss1;
@@ -3333,7 +3536,7 @@ struct server_context_impl {
                                         }
 
                                         {
-                                            const auto token = slot.task->tokens[i];
+                                            const auto token = input_tokens[i];
                                             const auto piece = token != LLAMA_TOKEN_NULL ? common_token_to_piece(ctx_tgt, token) : "[mtmd]";
                                             ss1 << piece;
                                             st1 << std::setw(8) << token;
@@ -3401,8 +3604,8 @@ struct server_context_impl {
                         }
 
                         // [TAG_PROMPT_LOGITS]
-                        if (n_past == slot.task->n_tokens() && n_past > 0) {
-                            SLT_WRN(slot, "need to evaluate at least 1 token for each active slot (n_past = %d, task.n_tokens() = %d)\n", n_past, slot.task->n_tokens());
+                        if (n_past == n_prompt_in && n_past > 0) {
+                            SLT_WRN(slot, "need to evaluate at least 1 token for each active slot (n_past = %d, task.n_tokens() = %d)\n", n_past, n_prompt_in);
                             n_past--;
                             SLT_WRN(slot, "n_past was set to %d\n", n_past);
                         }
@@ -3428,7 +3631,7 @@ struct server_context_impl {
 
                     if (!slot.can_split()) {
                         // cannot fit the prompt in the current batch - will try next iter
-                        if (batch.size() + slot.task->n_tokens() > n_batch) {
+                        if (batch.size() + n_prompt_in > n_batch) {
                             return;
                         }
                     }
@@ -3479,7 +3682,7 @@ struct server_context_impl {
                     while (true) {
                         auto cur_token_idx = slot.prompt.n_tokens();
                         if (
-                            cur_token_idx >= slot.task->n_tokens() ||
+                            cur_token_idx >= n_prompt_in ||
                             input_tokens[cur_token_idx] != LLAMA_TOKEN_NULL // encountered a text token
                         ) {
                             break;
@@ -3522,7 +3725,7 @@ struct server_context_impl {
                     const auto last_user_pos = spans.last_user_message_pos();
 
                     // add prompt tokens for processing in the current batch
-                    while (slot.prompt.n_tokens() < slot.task->n_tokens() && batch.size() < n_batch) {
+                    while (slot.prompt.n_tokens() < n_prompt_in && batch.size() < n_batch) {
                         // get next token to process
                         llama_token cur_tok = input_tokens[slot.prompt.n_tokens()];
                         if (cur_tok == LLAMA_TOKEN_NULL) {
@@ -3540,9 +3743,10 @@ struct server_context_impl {
                         // embedding requires all tokens in the batch to be output;
                         // MTP also wants logits at every prompt position so the
                         // streaming hook can mirror t_h_nextn into ctx_dft.
+                        llama_pos pos = slot.prompt.tokens.pos_next();
                         add_ok &= batch.add(slot.id,
                             cur_tok,
-                            /* pos       = */ slot.prompt.tokens.pos_next(),
+                            /* pos       = */ pos,
                             /* output    = */ slot.need_embd(),
                             /* is_prompt = */ true);
                         slot.prompt.tokens.push_back(cur_tok);
@@ -3568,7 +3772,7 @@ struct server_context_impl {
                             bool should_break = false;
                             for (int offset : checkpoint_offsets) {
                                 const int n_last = std::min(n_batch, offset);
-                                if (slot.task->n_tokens() == slot.prompt.n_tokens() + n_last) {
+                                if (n_prompt_in == slot.prompt.n_tokens() + n_last) {
                                     should_break = true;
                                     break;
                                 }
@@ -3584,15 +3788,19 @@ struct server_context_impl {
 
                     const auto n_tokens_start = slot.prompt.n_tokens() - n_tokens_cur;
 
-                    const bool near_prompt_end = slot.task->n_tokens() < slot.prompt.n_tokens() + n_ubatch;
+                    const bool near_prompt_end = n_prompt_in < slot.prompt.n_tokens() + n_ubatch;
 
                     const bool is_user_start = spans.is_user_start(n_tokens_start);
                     const bool is_last_user_message = n_tokens_start == last_user_pos;
 
                     // entire prompt has been processed
-                    if (slot.prompt.n_tokens() == slot.task->n_tokens()) {
+                    if (slot.prompt.n_tokens() == n_prompt_in) {
                         slot.state = SLOT_STATE_DONE_PROMPT;
 
+                        if (slot.spec_prefill_active && slot.task) {
+                            slot.stats.n_prompt_processed = slot.task->n_tokens() - slot.stats.n_prompt_cached;
+                        }
+
                         GGML_ASSERT(batch.size() > 0);
 
                         // extract the logits only for the last token
diff --git a/tools/server/server-models.cpp b/tools/server/server-models.cpp
index db0fac99527b..3ef010b5d979 100644
--- a/tools/server/server-models.cpp
+++ b/tools/server/server-models.cpp
@@ -626,7 +626,10 @@ void server_models::load_models() {
     };
     auto apply_hidden = [&]() {
         for (auto & [name, inst] : mapping) {
-            inst.meta.hidden = hidden_models.count(name) > 0;
+            std::string val;
+            bool by_preset = inst.meta.preset.get_option(COMMON_ARG_PRESET_HIDDEN, val)
+                             && common_arg_utils::is_truthy(val);
+            inst.meta.hidden = by_preset || hidden_models.count(name) > 0;
         }
     };
     // update_args() injects HOST/PORT/ALIAS, so strip them before comparing presets
@@ -1971,7 +1974,7 @@ void server_models_routes::init_routes() {
         std::time_t t = std::time(0);
         for (const auto & meta : all_models) {
             if (meta.hidden) {
-                continue; // cache model deduplicated by a preset
+                continue; // preset asked to hide it, or a cache model deduplicated by a preset
             }
             json status {
                 {"value",  server_model_status_to_string(meta.status)},
diff --git a/tools/ui/CMakeLists.txt b/tools/ui/CMakeLists.txt
index 208b46a5c15a..6c00c19621c4 100644
--- a/tools/ui/CMakeLists.txt
+++ b/tools/ui/CMakeLists.txt
@@ -49,10 +49,26 @@ if(CMAKE_CROSSCOMPILING)
         set(LLAMA_UI_EMBED_EXE "${CMAKE_CURRENT_BINARY_DIR}/llama-ui-embed-host")
     endif()
 
+    # The host compiler is not necessarily the one the toolchain file picked. A
+    # Windows toolchain file sets CMAKE_SYSTEM_NAME, which makes CMake treat the
+    # build as cross-compiling even on a Windows host, and HOST_CXX_COMPILER then
+    # resolves to MSVC. MSVC does not take GNU-style flags: it ignores -std=c++17
+    # with only a D9002 warning, drops to C++14, and embed.cpp then fails because
+    # <filesystem> is C++17-only. Pick the spelling the host compiler understands.
+    get_filename_component(HOST_CXX_NAME "${HOST_CXX_COMPILER}" NAME_WE)
+    if(HOST_CXX_NAME STREQUAL "cl")
+        set(LLAMA_UI_EMBED_FLAGS
+            /nologo /O2 /EHsc /std:c++17
+            "/Fo:${CMAKE_CURRENT_BINARY_DIR}/llama-ui-embed-host.obj"
+            "/Fe:${LLAMA_UI_EMBED_EXE}")
+    else()
+        set(LLAMA_UI_EMBED_FLAGS -O2 -std=c++17 -o "${LLAMA_UI_EMBED_EXE}")
+    endif()
+
     add_custom_command(
         OUTPUT  "${LLAMA_UI_EMBED_EXE}"
-        COMMAND "${HOST_CXX_COMPILER}" -O2 -std=c++17
-                -o "${LLAMA_UI_EMBED_EXE}" "${CMAKE_CURRENT_SOURCE_DIR}/embed.cpp"
+        COMMAND "${HOST_CXX_COMPILER}" ${LLAMA_UI_EMBED_FLAGS}
+                "${CMAKE_CURRENT_SOURCE_DIR}/embed.cpp"
         DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/embed.cpp"
         COMMENT "Building llama-ui-embed (host)"
         VERBATIM