imatrix : fix --nextn crash (#31522)

The note under 'Run the inference' still said that the CANN backend only
supports FP16/Q4_0/Q8_0 models, while the DataType table in the same file,
and ggml-cann.cpp (BF16 is enabled for MUL_MAT, GET_ROWS, CPY and CONT for
every non-310P device), say otherwise. Spell out the per-device support and
mark BF16 as unsupported on 310P instead of leaving the cell empty.

GGML_CANN_NO_PINNED is read by ggml-cann.cpp when allocating host buffers
and when advertising the host_buffer device capability, but it was not
documented, so the escape hatch for a failing pinned host allocation was
not discoverable.

Signed-off-by: li-lizhe <147392333@qq.com>

README.md

@@ -13,7 +13,7 @@
 [![Docker](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/docker.yml?label=Docker)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml)
 [![Winget](https://img.shields.io/github/actions/workflow/status/ggml-org/llama.cpp/winget.yml?label=Winget)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml)
 
-[ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md) / [maintainer PRs](https://github.com/ggml-org/llama.cpp/issues?q=is%3Apr%20is%3Aopen%20draft%3AFalse%20(author%3Asanmai%20OR%20author%3Argerganov%20OR%20author%3AKitaitiMakoto%20OR%20author%3Adanbev%20OR%20author%3Aaldehir%20OR%20author%3Amax-krasnyansky%20OR%20author%3ACISC%20OR%20author%3Aggerganov%20OR%20author%3Aam17an%20OR%20author%3Ajhen0409%20OR%20author%3Abartowski1182%20OR%20author%3Anikwen%20OR%20author%3Ahipudding%20OR%20author%3Aravi9%20OR%20author%3AServeurpersoCom%20OR%20author%3Athelittlefireman%20OR%20author%3Apwilkin%20OR%20author%3Areeselevine%20OR%20author%3Angxson%20OR%20author%3Ajeffbolznv%20OR%20author%3Amarty1885%20OR%20author%3A0cc4m%20OR%20author%3ATitaniumtown%20OR%20author%3Aangt%20OR%20author%3AIMbackK%20OR%20author%3Aarthw%20OR%20author%3AJohannesGaessler%20OR%20author%3AORippler%20OR%20author%3Aruixiang63%20OR%20author%3Axctan%20OR%20author%3AmasterFoad%20OR%20author%3Aallozaur%20OR%20author%3AAndreas-Krebbel%20OR%20author%3Ayomaytk%20OR%20author%3Aaendk%20OR%20author%3Awine99%20OR%20author%3Agaugarg-nv%20OR%20author%3Ataronaeo%20OR%20author%3Aforforever73%20OR%20author%3Achaxu01%20OR%20author%3Alhez%20OR%20author%3AAlekseiNikiforovIBM%20OR%20author%3Anetrunnereve%20OR%20author%3Afairydreaming%20OR%20author%3Aynankani)%20sort%3Aupdated-desc) / [dev stats](https://github.com/ggml-org/llama.cpp-dev) / [lib llama API](https://github.com/ggml-org/llama.cpp/issues/9289) / [llama-server REST API](https://github.com/ggml-org/llama.cpp/issues/9291)
+[ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md) / [maintainer PRs](https://github.com/ggml-org/llama.cpp/pulls?q=is%3Apr%20is%3Aopen%20draft%3AFalse%20(author%3Asanmai%20OR%20author%3Argerganov%20OR%20author%3AKitaitiMakoto%20OR%20author%3Adanbev%20OR%20author%3Aaldehir%20OR%20author%3Amax-krasnyansky%20OR%20author%3ACISC%20OR%20author%3Aggerganov%20OR%20author%3Aam17an%20OR%20author%3Ajhen0409%20OR%20author%3Abartowski1182%20OR%20author%3Anikwen%20OR%20author%3Ahipudding%20OR%20author%3Aravi9%20OR%20author%3AServeurpersoCom%20OR%20author%3Athelittlefireman%20OR%20author%3Apwilkin%20OR%20author%3Areeselevine%20OR%20author%3Angxson%20OR%20author%3Ajeffbolznv%20OR%20author%3Amarty1885%20OR%20author%3A0cc4m%20OR%20author%3ATitaniumtown%20OR%20author%3Aangt%20OR%20author%3AIMbackK%20OR%20author%3Aarthw%20OR%20author%3AJohannesGaessler%20OR%20author%3AORippler%20OR%20author%3Aruixiang63%20OR%20author%3Axctan%20OR%20author%3AmasterFoad%20OR%20author%3Aallozaur%20OR%20author%3AAndreas-Krebbel%20OR%20author%3Ayomaytk%20OR%20author%3Aaendk%20OR%20author%3Awine99%20OR%20author%3Agaugarg-nv%20OR%20author%3Ataronaeo%20OR%20author%3Aforforever73%20OR%20author%3Achaxu01%20OR%20author%3Alhez%20OR%20author%3AAlekseiNikiforovIBM%20OR%20author%3Anetrunnereve%20OR%20author%3Afairydreaming%20OR%20author%3Aynankani)%20sort%3Aupdated-desc) / [dev stats](https://github.com/ggml-org/llama.cpp-dev) / [lib llama API](https://github.com/ggml-org/llama.cpp/issues/9289) / [llama-server REST API](https://github.com/ggml-org/llama.cpp/issues/9291)
 
 </div>
 
@@ -79,7 +79,7 @@ The `llama.cpp` project is build on top of the [ggml](https://github.com/ggml-or
 | --- | --- |
 | [BLAS](docs/build.md#blas-build) | All |
 | [BLIS](docs/backend/BLIS.md) | All |
-| [CANN](docs/build.md#cann) | Ascend NPU |
+| [CANN](docs/build.md#cann) | Huawei Ascend NPU |
 | [CUDA](docs/build.md#cuda) | Nvidia GPU |
 | [HIP](docs/build.md#hip) | AMD GPU |
 | [Hexagon](docs/backend/snapdragon/README.md) | Snapdragon |

common/jinja/runtime.cpp

@@ -307,58 +307,13 @@ value binary_expression::execute_impl(context & ctx) const {
     throw std::runtime_error("Unknown operator \"" + op.value + "\" between " + left_val->type() + " and " + right_val->type());
 }
 
-static value try_builtin_func(context & ctx, const std::string & name, value & input, bool undef_on_missing = false) {
-    JJ_DEBUG("Trying built-in function '%s' for type %s", name.c_str(), input->type().c_str());
-    if (ctx.is_get_stats) {
-        value_t::stats_t::mark_used(input);
-        input->stats.ops.insert(name);
-    }
-    auto builtins = input->get_builtins();
-    auto it = builtins.find(name);
-    if (it != builtins.end()) {
-        JJ_DEBUG("Binding built-in '%s'", name.c_str());
-        return mk_val<value_func>(name, it->second, input);
-    }
-    if (undef_on_missing) {
-        return mk_val<value_undefined>(name);
-    }
-    throw std::runtime_error("Unknown (built-in) filter '" + name + "' for type " + input->type());
-}
-
 static value apply_filter(context & ctx, const statement_ptr & filter, value input) {
     JJ_DEBUG("Applying filter to %s", input->type().c_str());
 
-    auto set_filter_alias = [](auto & filter_id) {
-        if (filter_id == "count") {
-            filter_id = "length";
-        } else if (filter_id == "d") {
-            filter_id = "default";
-        } else if (filter_id == "e") {
-            filter_id = "escape";
-        } else if (filter_id == "trim") {
-            filter_id = "strip";
-        }
-    };
-
     if (is_stmt<identifier>(filter)) {
         auto filter_id = cast_stmt<identifier>(filter)->val;
-
-        set_filter_alias(filter_id);
         JJ_DEBUG("Applying filter '%s' to %s", filter_id.c_str(), input->type().c_str());
-        // TODO: Refactor filters so this coercion can be done automatically
-        if (!input->is_undefined() && !is_val<value_string>(input) && (
-            filter_id == "capitalize" ||
-            filter_id == "lower" ||
-            filter_id == "replace" ||
-            filter_id == "strip" ||
-            filter_id == "title" ||
-            filter_id == "upper" ||
-            filter_id == "wordcount"
-        )) {
-            JJ_DEBUG("Coercing %s to String for '%s' filter", input->type().c_str(), filter_id.c_str());
-            input = mk_val<value_string>(input->as_string());
-        }
-        return try_builtin_func(ctx, filter_id, input)->invoke(func_args(ctx));
+        return apply_filter_by_name(ctx, filter_id, input, func_args(ctx));
 
     } else if (is_stmt<call_expression>(filter)) {
         auto call = cast_stmt<call_expression>(filter);
@@ -367,14 +322,13 @@ static value apply_filter(context & ctx, const statement_ptr & filter, value inp
         }
         auto filter_id = cast_stmt<identifier>(call->callee)->val;
 
-        set_filter_alias(filter_id);
         JJ_DEBUG("Applying filter '%s' with arguments to %s", filter_id.c_str(), input->type().c_str());
         func_args args(ctx);
         for (const auto & arg_expr : call->args) {
             args.push_back(arg_expr->execute(ctx));
         }
 
-        return try_builtin_func(ctx, filter_id, input)->invoke(args);
+        return apply_filter_by_name(ctx, filter_id, input, args);
 
     } else {
         throw std::runtime_error("Invalid filter expression");

common/jinja/value.cpp

@@ -72,6 +72,17 @@ const std::vector<value> & func_args::get_args() const {
     return args;
 }
 
+void func_args::ensure_kwargs(std::initializer_list<const char *> names) const {
+    for (const auto & arg : args) {
+        if (is_val<value_kwarg>(arg)) {
+            auto * kwarg = cast_val<value_kwarg>(arg);
+            if (std::find(names.begin(), names.end(), kwarg->key) == names.end()) {
+                throw raised_exception("Function '" + func_name + "' got an unexpected keyword argument '" + kwarg->key + "'");
+            }
+        }
+    }
+}
+
 /**
  * Function that mimics Python's array slicing.
  */
@@ -134,6 +145,14 @@ static value empty_value_fn(const func_args &) {
         return mk_val<T>();
     }
 }
+
+template<typename T>
+static value empty_value_no_args_fn(const func_args & args) {
+    args.ensure_kwargs();
+    args.ensure_count(1, 1);
+    return empty_value_fn<T>(args);
+}
+
 template<typename T>
 static value test_type_fn(const func_args & args) {
     args.ensure_count(1);
@@ -594,6 +613,52 @@ const func_builtins & global_builtins() {
     return builtins;
 }
 
+value try_builtin_func(context & ctx, const std::string & name, value & input, bool undef_on_missing) {
+    JJ_DEBUG("Trying built-in function '%s' for type %s", name.c_str(), input->type().c_str());
+    if (ctx.is_get_stats) {
+        value_t::stats_t::mark_used(input);
+        input->stats.ops.insert(name);
+    }
+    auto builtins = input->get_builtins();
+    auto it = builtins.find(name);
+    if (it != builtins.end()) {
+        JJ_DEBUG("Binding built-in '%s'", name.c_str());
+        return mk_val<value_func>(name, it->second, input);
+    }
+    if (undef_on_missing) {
+        return mk_val<value_undefined>(name);
+    }
+    throw std::runtime_error("Unknown (built-in) filter '" + name + "' for type " + input->type());
+}
+
+value apply_filter_by_name(context & ctx, std::string filter_id, value input, const func_args & args) {
+    if (filter_id == "count") {
+        filter_id = "length";
+    } else if (filter_id == "d") {
+        filter_id = "default";
+    } else if (filter_id == "e") {
+        filter_id = "escape";
+    } else if (filter_id == "trim") {
+        filter_id = "strip";
+    }
+
+    // TODO: Refactor filters so this coercion can be done automatically
+    if (!input->is_undefined() && !is_val<value_string>(input) && (
+        filter_id == "capitalize" ||
+        filter_id == "lower" ||
+        filter_id == "replace" ||
+        filter_id == "strip" ||
+        filter_id == "title" ||
+        filter_id == "upper" ||
+        filter_id == "wordcount"
+    )) {
+        JJ_DEBUG("Coercing %s to String for '%s' filter", input->type().c_str(), filter_id.c_str());
+        input = mk_val<value_string>(input->as_string());
+    }
+
+    return try_builtin_func(ctx, filter_id, input)->invoke(args);
+}
+
 
 const func_builtins & value_int_t::get_builtins() const {
     static const func_builtins builtins = {
@@ -667,16 +732,22 @@ const func_builtins & value_string_t::get_builtins() const {
     static const func_builtins builtins = {
         {"default", default_value},
         {"upper", [](const func_args & args) -> value {
+            args.ensure_kwargs();
+            args.ensure_count(1, 1);
             args.ensure_vals<value_string>();
             jinja::string str = args.get_pos(0)->as_string().uppercase();
             return mk_val<value_string>(str);
         }},
         {"lower", [](const func_args & args) -> value {
+            args.ensure_kwargs();
+            args.ensure_count(1, 1);
             args.ensure_vals<value_string>();
             jinja::string str = args.get_pos(0)->as_string().lowercase();
             return mk_val<value_string>(str);
         }},
         {"strip", [](const func_args & args) -> value {
+            args.ensure_kwargs({"chars"});
+            args.ensure_count(1, 2);
             value val_input = args.get_pos(0);
             if (!is_val<value_string>(val_input)) {
                 throw raised_exception("strip() first argument must be a string");
@@ -707,11 +778,15 @@ const func_builtins & value_string_t::get_builtins() const {
             }
         }},
         {"title", [](const func_args & args) -> value {
+            args.ensure_kwargs();
+            args.ensure_count(1, 1);
             args.ensure_vals<value_string>();
             jinja::string str = args.get_pos(0)->as_string().titlecase();
             return mk_val<value_string>(str);
         }},
         {"capitalize", [](const func_args & args) -> value {
+            args.ensure_kwargs();
+            args.ensure_count(1, 1);
             args.ensure_vals<value_string>();
             jinja::string str = args.get_pos(0)->as_string().capitalize();
             return mk_val<value_string>(str);
@@ -1013,6 +1088,10 @@ const func_builtins & value_bool_t::get_builtins() const {
     return builtins;
 }
 
+[[noreturn]] static value array_sum_not_implemented(const func_args &) {
+    throw not_implemented_exception("Array sum builtin not implemented");
+}
+
 const func_builtins & value_array_t::get_builtins() const {
     static const func_builtins builtins = {
         {"default", default_value},
@@ -1079,7 +1158,9 @@ const func_builtins & value_array_t::get_builtins() const {
         {"select", select_or_reject<false, false>},
         {"rejectattr", select_or_reject<true, true>},
         {"reject", select_or_reject<true, false>},
+        {"sum", array_sum_not_implemented},
         {"join", [](const func_args & args) -> value {
+            args.ensure_kwargs({"d", "attribute"});
             args.ensure_count(1, 3);
             if (!is_val<value_array>(args.get_pos(0))) {
                 throw raised_exception("join() first argument must be an array");
@@ -1117,13 +1198,28 @@ const func_builtins & value_array_t::get_builtins() const {
         {"tojson", tojson},
         {"map", [](const func_args & args) -> value {
             args.ensure_count(2);
-            if (!is_val<value_array>(args.get_pos(0))) {
+            value val = args.get_pos(0);
+            if (!is_val<value_array>(val)) {
                 throw raised_exception("map: first argument must be an array");
             }
             if (!is_val<value_kwarg>(args.get_args().at(1))) {
-                throw not_implemented_exception("map: filter-mapping not implemented");
+                value filter = args.get_pos(1);
+                if (!is_val<value_string>(filter)) {
+                    throw raised_exception("map: filter name must be a string");
+                }
+                auto out = mk_val<value_array>();
+                func_args filter_args(args.ctx);
+                const auto & all_args = args.get_args();
+                for (size_t i = 2; i < all_args.size(); ++i) {
+                    filter_args.push_back(all_args[i]);
+                }
+                const std::string filter_id = filter->as_string().str();
+                auto arr = val->as_array();
+                for (const auto & item : arr) {
+                    out->push_back(apply_filter_by_name(args.ctx, filter_id, item, filter_args));
+                }
+                return is_val<value_tuple>(val) ? mk_val<value_tuple>(out->as_array()) : out;
             }
-            value val       = args.get_pos(0);
             value attribute = args.get_kwarg_or_pos("attribute", 1);
             value default_val = args.get_kwarg("default", mk_val<value_undefined>());
             auto out = mk_val<value_array>();
@@ -1132,7 +1228,7 @@ const func_builtins & value_array_t::get_builtins() const {
                 value attr_val = get_attribute(item, attribute, default_val);
                 out->push_back(attr_val);
             }
-            return is_val<value_tuple>(val) ? mk_val<value_tuple>(std::move(out->as_array())) : out;
+            return is_val<value_tuple>(val) ? mk_val<value_tuple>(out->as_array()) : out;
         }},
         {"append", [](const func_args & args) -> value {
             args.ensure_count(2);
@@ -1405,14 +1501,18 @@ const func_builtins & value_none_t::get_builtins() const {
 const func_builtins & value_undefined_t::get_builtins() const {
     static const func_builtins builtins = {
         {"default", default_value},
-        {"capitalize", empty_value_fn<value_string>},
+        {"capitalize", empty_value_no_args_fn<value_string>},
         {"first", empty_value_fn<value_undefined>},
         {"items", empty_value_fn<value_array>},
-        {"join", empty_value_fn<value_string>},
+        {"join", [](const func_args & args) -> value {
+            args.ensure_kwargs({"d", "attribute"});
+            args.ensure_count(1, 3);
+            return empty_value_fn<value_string>(args);
+        }},
         {"last", empty_value_fn<value_undefined>},
         {"length", empty_value_fn<value_int>},
         {"list", empty_value_fn<value_array>},
-        {"lower", empty_value_fn<value_string>},
+        {"lower", empty_value_no_args_fn<value_string>},
         {"map", empty_value_fn<value_array>},
         {"max", empty_value_fn<value_undefined>},
         {"min", empty_value_fn<value_undefined>},
@@ -1425,12 +1525,16 @@ const func_builtins & value_undefined_t::get_builtins() const {
         {"selectattr", empty_value_fn<value_array>},
         {"sort", empty_value_fn<value_array>},
         {"string", empty_value_fn<value_string>},
-        {"strip", empty_value_fn<value_string>},
+        {"strip", [](const func_args & args) -> value {
+            args.ensure_kwargs({"chars"});
+            args.ensure_count(1, 2);
+            return empty_value_fn<value_string>(args);
+        }},
         {"sum", empty_value_fn<value_int>},
-        {"title", empty_value_fn<value_string>},
+        {"title", empty_value_no_args_fn<value_string>},
         {"truncate", empty_value_fn<value_string>},
         {"unique", empty_value_fn<value_array>},
-        {"upper", empty_value_fn<value_string>},
+        {"upper", empty_value_no_args_fn<value_string>},
         {"wordcount", empty_value_fn<value_int>},
     };
     return builtins;

common/jinja/value.h

@@ -7,6 +7,7 @@
 #include <cmath>
 #include <cstdint>
 #include <functional>
+#include <initializer_list>
 #include <map>
 #include <memory>
 #include <set>
@@ -176,6 +177,8 @@ protected:
 //
 
 const func_builtins & global_builtins();
+value try_builtin_func(context & ctx, const std::string & name, value & input, bool undef_on_missing = false);
+value apply_filter_by_name(context & ctx, std::string filter_id, value input, const func_args & args);
 
 std::string value_to_json(const value & val, int indent = -1, const std::string_view item_sep = ", ", const std::string_view key_sep = ": ");
 
@@ -658,6 +661,7 @@ public:
     size_t count() const { return args.size(); }
     void push_back(const value & val);
     void push_front(const value & val);
+    void ensure_kwargs(std::initializer_list<const char *> names = {}) const;
     void ensure_count(size_t min, size_t max = 999) const {
         size_t n = args.size();
         if (n < min || n > max) {

docs/backend/CANN.md

@@ -161,7 +161,7 @@ lspci -n | grep -Eo '19e5:d[0-9a-f]{3}' | cut -d: -f2
 | FP16                   | Support | Support |
 | Q8_0                   | Support | Partial |
 | Q4_0                   | Support | Partial |
-| BF16                   | Support |         |
+| BF16                   | Support |    No   |
 
 > **310P note**
 > - `Q8_0`: data transform / buffer path is implemented, and `GET_ROWS` is supported, but quantized `MUL_MAT` / `MUL_MAT_ID` are not supported.
@@ -274,7 +274,9 @@ cmake --build build --config release
 
     **Notes**:
 
-      - CANN backend only supports FP16/Q4_0/Q8_0 models currently.
+      - On 910B the CANN backend supports FP16, BF16 and Q4_0/Q8_0 models. On 310P, BF16 is not
+        supported and quantized `MUL_MAT` / `MUL_MAT_ID` are not implemented, so only FP16 models
+        are supported (see [DataType Supports](#datatype-supports)).
 
 2. **Launch inference**
 
@@ -355,3 +357,7 @@ Enable ACL graph execution during the prefill stage, default is false. This opti
 ### GGML_CANN_OPERATOR_FUSION
 
 Enable operator fusion during computation, default is false. This option fuses compatible operators (e.g., ADD + RMS_NORM) to reduce overhead and improve performance.
+
+### GGML_CANN_NO_PINNED
+
+Disable pinned (host-locked) host buffers. When it is set (any value), `aclrtMallocHost` is not called: the CANN host buffer type allocates regular CPU memory instead and the backend reports no `host_buffer` capability, so host-side staging uses pageable memory. Use it when pinned host allocation fails, or when the pinned memory is needed by other processes on the host.

docs/build.md

@@ -481,7 +481,7 @@ cmake --build build --config Release
 Now you can load the model in conversation mode using `Vulkan`
 
 ```sh
-build/bin/Release/llama-cli -m "[PATH TO MODEL]" -ngl 100 -c 16384 -t 10 -n -2 -cnv
+build/bin/Release/llama-cli -m "[PATH TO MODEL]" -ngl 100 -c 16384 -t 10 -n -2
 ```
 
 **MSYS2**

ggml/src/ggml-cpu/ggml-cpu-impl.h

@@ -30,7 +30,7 @@ struct ggml_compute_params {
 };
 
 
-#if defined(_MSC_VER)
+#if defined(_MSC_VER) && !defined(__clang__)
 
 #define m512bh(p) p
 #define m512i(p) p

ggml/src/ggml-metal/ggml-metal-device.m

@@ -2102,19 +2102,11 @@ static void ggml_metal_buffer_rset_free(ggml_metal_buffer_t buf) {
 static void * ggml_metal_host_malloc(size_t n) {
     void * data = NULL;
 
-#if TARGET_OS_OSX
     kern_return_t err = vm_allocate((vm_map_t) mach_task_self(), (void *) &data, n, VM_FLAGS_ANYWHERE);
     if (err != KERN_SUCCESS) {
         GGML_LOG_ERROR("%s: error: vm_allocate failed\n", __func__);
         return NULL;
     }
-#else
-    const int result = posix_memalign((void **) &data, sysconf(_SC_PAGESIZE), n);
-    if (result != 0) {
-        GGML_LOG_ERROR("%s: error: posix_memalign failed\n", __func__);
-        return NULL;
-    }
-#endif
 
     return data;
 }
@@ -2296,11 +2288,7 @@ void ggml_metal_buffer_free(ggml_metal_buffer_t buf) {
     }
 
     if (buf->is_shared && buf->owned) {
-#if TARGET_OS_OSX
         vm_deallocate((vm_map_t)mach_task_self(), (vm_address_t)buf->all_data, buf->all_size);
-#else
-        free(buf->all_data);
-#endif
     }
 
     free(buf);

ggml/src/ggml-opencl/CMakeLists.txt

@@ -67,6 +67,7 @@ set(GGML_OPENCL_KERNELS
     argmax
     argsort
     top_k
+    top_k_mask
     lightning_indexer
     tri
     fill

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -873,6 +873,7 @@ struct ggml_backend_opencl_context {
     ggml_cl_adreno_xmem_attn_state adreno_xmem_attn;
 #endif
     cl_kernel kernel_top_k_f32, kernel_top_k_merge_f32;
+    cl_kernel kernel_top_k_mask[2], kernel_top_k_mask_fill[2];
     cl_kernel kernel_lightning_indexer;
     cl_kernel kernel_get_rows_f32, kernel_get_rows_f16, kernel_get_rows_i32, kernel_get_rows_q4_0;
     cl_kernel kernel_set_rows_f32_i64, kernel_set_rows_f32_i32, kernel_set_rows_f16_i64, kernel_set_rows_f16_i32;
@@ -2298,6 +2299,27 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
         GGML_LOG_CONT(".");
     }
 
+    // top_k_mask
+    {
+#ifdef GGML_OPENCL_EMBED_KERNELS
+        const std::string kernel_src {
+            #include "top_k_mask.cl.h"
+        };
+#else
+        const std::string kernel_src = read_file("top_k_mask.cl");
+#endif
+        // Scores can contain NaN and -inf, so do not use the default fast-math options.
+        const std::string opts = "-cl-std=" + opencl_c_std +
+            (backend_ctx->adreno_use_large_buffer ? " -qcom-enable-large-buffer" : "");
+        for (int i = 0; i < 2; ++i) {
+            cl_program prog = build_program_from_source(backend_ctx, kernel_src.c_str(), opts + " -DMASK_F16=" + std::to_string(i == 0));
+            CL_CHECK((backend_ctx->kernel_top_k_mask[i] = clCreateKernel(prog, "kernel_top_k_mask", &err), err));
+            CL_CHECK((backend_ctx->kernel_top_k_mask_fill[i] = clCreateKernel(prog, "kernel_top_k_mask_fill", &err), err));
+            CL_CHECK(clReleaseProgram(prog));
+        }
+        GGML_LOG_CONT(".");
+    }
+
     // get_rows
     {
 #ifdef GGML_OPENCL_EMBED_KERNELS
@@ -11800,6 +11822,14 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
                    op->src[2]->nb[0] == sizeof(float) && op->src[3]->nb[0] == sizeof(ggml_fp16_t);
         case GGML_OP_TOP_K:
             return op->src[0]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_I32;
+        case GGML_OP_TOP_K_MASK:
+            return (op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_F32) &&
+                op->src[0]->type == op->type &&
+                op->src[0]->nb[0] == ggml_type_size(op->type) && op->nb[0] == ggml_type_size(op->type) &&
+                op->src[1]->type == GGML_TYPE_I32 &&
+                (!op->src[2] || op->src[2]->type == GGML_TYPE_F32) &&
+                (!op->src[3] || op->src[3]->type == GGML_TYPE_I32) &&
+                (!op->src[4] || op->src[4]->type == GGML_TYPE_I32);
         case GGML_OP_ARGSORT: {
             load_cl_kernels_argsort(backend_ctx);
 
@@ -36968,6 +36998,119 @@ static void ggml_cl_top_k(ggml_backend_t backend, const ggml_tensor * src0, cons
     }
 }
 
+static void ggml_cl_top_k_mask(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    ggml_backend_opencl_context * backend_ctx = (ggml_backend_opencl_context *) backend->context;
+    const ggml_tensor * score  = dst->src[2];
+    const ggml_tensor * blocks = dst->src[3];
+    const ggml_tensor * extra  = dst->src[4];
+    const int type = dst->type == GGML_TYPE_F16 ? 0 : 1;
+    cl_kernel kernel = backend_ctx->kernel_top_k_mask[type];
+
+    const ggml_tensor_extra_cl * extram = (const ggml_tensor_extra_cl *) src0->extra;
+    const ggml_tensor_extra_cl * extrak = (const ggml_tensor_extra_cl *) src1->extra;
+    const ggml_tensor_extra_cl * extras = score  ? (const ggml_tensor_extra_cl *) score->extra  : nullptr;
+    const ggml_tensor_extra_cl * extrab = blocks ? (const ggml_tensor_extra_cl *) blocks->extra : nullptr;
+    const ggml_tensor_extra_cl * extrae = extra  ? (const ggml_tensor_extra_cl *) extra->extra  : nullptr;
+    const ggml_tensor_extra_cl * extrad = (const ggml_tensor_extra_cl *) dst->extra;
+
+    const cl_mem buffer_s = extras ? extras->data_device : nullptr;
+    const cl_mem buffer_b = extrab ? extrab->data_device : nullptr;
+    const cl_mem buffer_e = extrae ? extrae->data_device : nullptr;
+
+    const cl_ulong offsetm = extram->offset + src0->view_offs;
+    const cl_ulong offsetk = extrak->offset + src1->view_offs;
+    const cl_ulong offsets = extras ? extras->offset + score->view_offs  : 0;
+    const cl_ulong offsetb = extrab ? extrab->offset + blocks->view_offs : 0;
+    const cl_ulong offsete = extrae ? extrae->offset + extra->view_offs  : 0;
+    const cl_ulong offsetd = extrad->offset + dst->view_offs;
+
+    // single launch: fill == true, two launches: fill = false
+    // X2 benefits from a parallel fill for long single rows; mobile Adreno favors one launch.
+    const cl_int fill = !(backend_ctx->adreno_gen == ADRENO_GPU_GEN::X2E &&
+        ggml_nrows(dst) == 1 && dst->ne[0] >= 65536 && ggml_is_contiguous(dst));
+
+    if (!fill) {
+        GGML_ASSERT(ggml_is_contiguous(dst));
+
+        cl_kernel kernel_fill = backend_ctx->kernel_top_k_mask_fill[type];
+
+        const cl_ulong n = ggml_nelements(dst);
+        const size_t nv = 16/ggml_type_size(dst->type);
+        const size_t nth = MIN((size_t) 256, backend_ctx->get_kernel_workgroup_size(kernel_fill));
+
+        CL_CHECK(clSetKernelArg(kernel_fill, 0, sizeof(cl_mem), &extrad->data_device));
+        CL_CHECK(clSetKernelArg(kernel_fill, 1, sizeof(cl_ulong), &offsetd));
+        CL_CHECK(clSetKernelArg(kernel_fill, 2, sizeof(cl_ulong), &n));
+
+        size_t global[] = {((n + nv - 1)/nv + nth - 1)/nth*nth};
+        size_t local[] = {nth};
+        backend_ctx->enqueue_ndrange_kernel(kernel_fill, 1, global, local, dst);
+    }
+
+    GGML_TENSOR_LOCALS(cl_ulong, nem, src0,   ne);
+    GGML_TENSOR_LOCALS(cl_ulong, nbm, src0,   nb);
+    GGML_TENSOR_LOCALS(cl_ulong, nek, src1,   ne);
+    GGML_TENSOR_LOCALS(cl_ulong, nbk, src1,   nb);
+    GGML_TENSOR_LOCALS(cl_ulong, nes, score,  ne);
+    GGML_TENSOR_LOCALS(cl_ulong, nbs, score,  nb);
+    GGML_TENSOR_LOCALS(cl_ulong, neb, blocks, ne);
+    GGML_TENSOR_LOCALS(cl_ulong, nbb, blocks, nb);
+    GGML_TENSOR_LOCALS(cl_ulong, nee, extra,  ne);
+    GGML_TENSOR_LOCALS(cl_ulong, nbe, extra,  nb);
+    GGML_TENSOR_LOCALS(cl_ulong, ned, dst,    ne);
+    GGML_TENSOR_LOCALS(cl_ulong, nbd, dst,    nb);
+
+    CL_CHECK(clSetKernelArg(kernel,  0, sizeof(cl_mem),   &extram->data_device));
+    CL_CHECK(clSetKernelArg(kernel,  1, sizeof(cl_ulong), &offsetm));
+    CL_CHECK(clSetKernelArg(kernel,  2, sizeof(cl_mem),   &extrak->data_device));
+    CL_CHECK(clSetKernelArg(kernel,  3, sizeof(cl_ulong), &offsetk));
+    CL_CHECK(clSetKernelArg(kernel,  4, sizeof(cl_mem),   &buffer_s));
+    CL_CHECK(clSetKernelArg(kernel,  5, sizeof(cl_ulong), &offsets));
+    CL_CHECK(clSetKernelArg(kernel,  6, sizeof(cl_mem),   &buffer_b));
+    CL_CHECK(clSetKernelArg(kernel,  7, sizeof(cl_ulong), &offsetb));
+    CL_CHECK(clSetKernelArg(kernel,  8, sizeof(cl_mem),   &buffer_e));
+    CL_CHECK(clSetKernelArg(kernel,  9, sizeof(cl_ulong), &offsete));
+    CL_CHECK(clSetKernelArg(kernel, 10, sizeof(cl_mem),   &extrad->data_device));
+    CL_CHECK(clSetKernelArg(kernel, 11, sizeof(cl_ulong), &offsetd));
+    CL_CHECK(clSetKernelArg(kernel, 12, sizeof(cl_ulong), &nem0));
+    CL_CHECK(clSetKernelArg(kernel, 13, sizeof(cl_ulong), &nem1));
+    CL_CHECK(clSetKernelArg(kernel, 14, sizeof(cl_ulong), &nem2));
+    CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_ulong), &nem3));
+    CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_ulong), &nbm1));
+    CL_CHECK(clSetKernelArg(kernel, 17, sizeof(cl_ulong), &nbm2));
+    CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_ulong), &nbm3));
+    CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_ulong), &nek0));
+    CL_CHECK(clSetKernelArg(kernel, 20, sizeof(cl_ulong), &nbk0));
+    CL_CHECK(clSetKernelArg(kernel, 21, sizeof(cl_ulong), &nbk1));
+    CL_CHECK(clSetKernelArg(kernel, 22, sizeof(cl_ulong), &nbk2));
+    CL_CHECK(clSetKernelArg(kernel, 23, sizeof(cl_ulong), &nbk3));
+    CL_CHECK(clSetKernelArg(kernel, 24, sizeof(cl_ulong), &nes0));
+    CL_CHECK(clSetKernelArg(kernel, 25, sizeof(cl_ulong), &nbs0));
+    CL_CHECK(clSetKernelArg(kernel, 26, sizeof(cl_ulong), &nbs1));
+    CL_CHECK(clSetKernelArg(kernel, 27, sizeof(cl_ulong), &nbs2));
+    CL_CHECK(clSetKernelArg(kernel, 28, sizeof(cl_ulong), &nbs3));
+    CL_CHECK(clSetKernelArg(kernel, 29, sizeof(cl_ulong), &neb0));
+    CL_CHECK(clSetKernelArg(kernel, 30, sizeof(cl_ulong), &neb1));
+    CL_CHECK(clSetKernelArg(kernel, 31, sizeof(cl_ulong), &nbb0));
+    CL_CHECK(clSetKernelArg(kernel, 32, sizeof(cl_ulong), &nbb1));
+    CL_CHECK(clSetKernelArg(kernel, 33, sizeof(cl_ulong), &nee0));
+    CL_CHECK(clSetKernelArg(kernel, 34, sizeof(cl_ulong), &nbe0));
+    CL_CHECK(clSetKernelArg(kernel, 35, sizeof(cl_ulong), &nbe1));
+    CL_CHECK(clSetKernelArg(kernel, 36, sizeof(cl_ulong), &nbe2));
+    CL_CHECK(clSetKernelArg(kernel, 37, sizeof(cl_ulong), &nbe3));
+    CL_CHECK(clSetKernelArg(kernel, 38, sizeof(cl_ulong), &ned1));
+    CL_CHECK(clSetKernelArg(kernel, 39, sizeof(cl_ulong), &ned2));
+    CL_CHECK(clSetKernelArg(kernel, 40, sizeof(cl_ulong), &nbd1));
+    CL_CHECK(clSetKernelArg(kernel, 41, sizeof(cl_ulong), &nbd2));
+    CL_CHECK(clSetKernelArg(kernel, 42, sizeof(cl_ulong), &nbd3));
+    CL_CHECK(clSetKernelArg(kernel, 43, sizeof(cl_int),   &fill));
+
+    const size_t nth = MIN((size_t) 512, backend_ctx->get_kernel_workgroup_size(kernel));
+    size_t global[] = {nth, (size_t) ggml_nrows(dst)};
+    size_t local[] = {nth, 1};
+    backend_ctx->enqueue_ndrange_kernel(kernel, 2, global, local, dst);
+}
+
 static void ggml_cl_argsort(ggml_backend_t backend, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
     GGML_ASSERT(src0);
     GGML_ASSERT(src0->extra);
@@ -37893,6 +38036,9 @@ bool ggml_cl_compute_forward(ggml_backend_t backend, struct ggml_tensor * tensor
         case GGML_OP_TOP_K:
             func = ggml_cl_top_k;
             break;
+        case GGML_OP_TOP_K_MASK:
+            func = ggml_cl_top_k_mask;
+            break;
         case GGML_OP_ARGSORT:
             if (!any_on_device) {
                 return false;

ggml/src/ggml-opencl/kernels/top_k_mask.cl

@@ -0,0 +1,129 @@
+// Integer storage preserves the mask bits without FP16 conversion.
+#if MASK_F16
+#define MASK_T ushort
+#define MASK_V ushort8
+#define MASK_NV 8
+#define MASK_NINF 0xfc00
+#define MASK_STORE vstore8
+#else
+#define MASK_T uint
+#define MASK_V uint4
+#define MASK_NV 4
+#define MASK_NINF 0xff800000u
+#define MASK_STORE vstore4
+#endif
+
+kernel void kernel_top_k_mask_fill(
+        global char * dst,
+        ulong         offsetd,
+        ulong         n
+) {
+    global MASK_T * out = (global MASK_T *) (dst + offsetd);
+    const ulong c = get_global_id(0)*MASK_NV;
+    if (c + MASK_NV <= n) {
+        MASK_STORE((MASK_V) (MASK_NINF), 0, out + c);
+    } else {
+        for (ulong j = c; j < n; ++j) {
+            out[j] = MASK_NINF;
+        }
+    }
+}
+
+kernel void kernel_top_k_mask(
+        global const char * mask,
+        ulong               offsetm,
+        global const char * top_k,
+        ulong               offsetk,
+        global const char * score,
+        ulong               offsets,
+        global const char * blocks,
+        ulong               offsetb,
+        global const char * extra,
+        ulong               offsete,
+        global       char * dst,
+        ulong               offsetd,
+        ulong nem0,
+        ulong nem1,
+        ulong nem2,
+        ulong nem3,
+        ulong nbm1,
+        ulong nbm2,
+        ulong nbm3,
+        ulong nek0,
+        ulong nbk0,
+        ulong nbk1,
+        ulong nbk2,
+        ulong nbk3,
+        ulong nes0,
+        ulong nbs0,
+        ulong nbs1,
+        ulong nbs2,
+        ulong nbs3,
+        ulong neb0,
+        ulong neb1,
+        ulong nbb0,
+        ulong nbb1,
+        ulong nee0,
+        ulong nbe0,
+        ulong nbe1,
+        ulong nbe2,
+        ulong nbe3,
+        ulong ned1,
+        ulong ned2,
+        ulong nbd1,
+        ulong nbd2,
+        ulong nbd3,
+        int   fill
+) {
+    const ulong row = get_group_id(1);
+    const ulong i1 = row % ned1;
+    const ulong i2 = (row / ned1) % ned2;
+    const ulong i3 = row / (ned1*ned2);
+    const ulong tid = get_local_id(0);
+    const ulong nth = get_local_size(0);
+    const ulong n_pick = blocks ? neb1 : score ? nes0 : nem0;
+
+    global const MASK_T * mask_row = (global const MASK_T *) (mask + offsetm
+        + (i1 % nem1)*nbm1 + (i2 % nem2)*nbm2 + (i3 % nem3)*nbm3);
+    global MASK_T * dst_row = (global MASK_T *) (dst + offsetd + i1*nbd1 + i2*nbd2 + i3*nbd3);
+    global const char * top_k_row = top_k + offsetk + i1*nbk1 + i2*nbk2 + i3*nbk3;
+
+    if (fill) {
+        for (ulong c = tid; c < nem0/MASK_NV; c += nth) {
+            MASK_STORE((MASK_V) (MASK_NINF), c, dst_row);
+        }
+        for (ulong c = nem0/MASK_NV*MASK_NV + tid; c < nem0; c += nth) {
+            dst_row[c] = MASK_NINF;
+        }
+        barrier(CLK_GLOBAL_MEM_FENCE);
+    }
+
+    // Repeated cells copy the same mask value, as in the CUDA and Metal kernels.
+    for (ulong j = tid; j < nek0 + nee0; j += nth) {
+        if (j >= nek0) {
+            const int c = *(global const int *) (extra + offsete + (j - nek0)*nbe0 + i1*nbe1 + i2*nbe2 + i3*nbe3);
+            if (c >= 0 && (ulong) c < nem0) {
+                dst_row[c] = mask_row[c];
+            }
+            continue;
+        }
+
+        const int p = *(global const int *) (top_k_row + j*nbk0);
+        if (p < 0 || (ulong) p >= n_pick) {
+            continue;
+        }
+        if (score && !(*(global const float *) (score + offsets + p*nbs0 + i1*nbs1 + i2*nbs2 + i3*nbs3) > -INFINITY)) {
+            continue;
+        }
+        if (blocks) {
+            for (ulong k = 0; k < neb0; ++k) {
+                const int c = *(global const int *) (blocks + offsetb + k*nbb0 + p*nbb1);
+                if (c >= 0 && (ulong) c < nem0) {
+                    dst_row[c] = mask_row[c];
+                }
+            }
+        } else if ((ulong) p < nem0) {
+            dst_row[p] = mask_row[p];
+        }
+    }
+}

ggml/src/ggml-sycl/dmmv.cpp

@@ -112,11 +112,12 @@ static void dequantize_mul_mat_vec(const void * __restrict__ vx, const dfloat *
     }
 }
 
-template <int qk, int qr, dequantize_kernel_t_reorder dequantize_kernel_reorder>
+template <int qk, int qr, int qbytes, dequantize_kernel_t_reorder dequantize_kernel_reorder>
 static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const dfloat * __restrict__ y, float * __restrict__ dst, const int ncols, const int nrows,
                                    const sycl::nd_item<3> &item_ct1) {
     // qk = quantized weights per x block
     // qr = number of quantized weights per data value in x block
+    // qbytes = bytes of the quants region per block (not qk/qr for bit-packed types)
     const int row = item_ct1.get_group(2) * item_ct1.get_local_range(1) +
                     item_ct1.get_local_id(1);
 
@@ -127,7 +128,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
     const int tid = item_ct1.get_local_id(2);
 
 
-    const int ncols_left = ncols % (QK4_0*WARP_SIZE);
+    const int ncols_left = ncols % (qk*WARP_SIZE);
     const int ncols_align = ncols - ncols_left;
     const int iter_stride = 8*2*GGML_SYCL_DMMV_X;
     const int vals_per_iter = iter_stride / WARP_SIZE; // num quantized vals per thread and i iter //64/16=4, 512/16/2= 16
@@ -139,7 +140,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
 #else
     float tmp = 0.0f;
 #endif // GGML_SYCL_F16
-    const char *d_ptr = (const char*)vx+ncols*nrows/2;
+    const char *d_ptr = (const char*)vx + (size_t) ncols * nrows / qk * qbytes;
     int i=0;
     for (i = 0; i < ncols_align; i += iter_stride) {
         const int col = i + vals_per_iter*tid;
@@ -155,7 +156,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
             // dequantize
             // for qr = 2 the iqs needs to increase by 1 per j iter because 2 weights per data val
             dfloat2 v;
-            dequantize_kernel_reorder((const void *)d_ptr, ib, (const void *)vx, ib * QK4_0 / 2 +iqs+j/qr, v);
+            dequantize_kernel_reorder((const void *)d_ptr, ib, (const void *)vx, ib * (qk/qr) +iqs+j/qr, v);
 
             // matrix multiplication
             // for qr = 2 the y index needs to increase by 1 per j iter because of y_offset = qk/2
@@ -172,7 +173,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
     }
 
     for (; i < ncols; i += iter_stride) {
-        if (tid>=ncols_left/QK4_0) continue;
+        if (tid>=ncols_left/qk) continue;
         const int col = i + vals_per_iter*tid;
         const int ib = (row*ncols + col)/qk; // x block index
         const int iqs = (col%qk)/qr; // x quant index
@@ -186,7 +187,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
             // dequantize
             // for qr = 2 the iqs needs to increase by 1 per j iter because 2 weights per data val
             dfloat2 v;
-            dequantize_kernel_reorder((const void *)d_ptr, ib, (const void *)vx, ib * QK4_0 / 2 +iqs+j/qr, v);
+            dequantize_kernel_reorder((const void *)d_ptr, ib, (const void *)vx, ib * (qk/qr) +iqs+j/qr, v);
 
             // matrix multiplication
             // for qr = 2 the y index needs to increase by 1 per j iter because of y_offset = qk/2
@@ -1487,10 +1488,10 @@ static void dequantize_mul_mat_vec_q6_k_reorder(const void * __restrict__ vx, co
     }
 }
 
-static void dequantize_mul_mat_vec_q4_0_sycl_reorder(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
+template <int qk, int qr, dequantize_kernel_t dequantize_kernel>
+static void launch_dequantize_mul_mat_vec(const void *vx, const dfloat *y, float *dst,
+                                          const int ncols, const int nrows,
+                                          dpct::queue_ptr stream) {
     GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
     const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
     // the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
@@ -1503,17 +1504,16 @@ static void dequantize_mul_mat_vec_q4_0_sycl_reorder(const void *vx, const dfloa
         stream->parallel_for(
             sycl::nd_range<3>(block_nums * block_dims, block_dims),
             [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                dequantize_mul_mat_vec_reorder<QK4_0, QR4_0, dequantize_q4_0_reorder>(
+                dequantize_mul_mat_vec<qk, qr, dequantize_kernel>(
                     vx, y, dst, ncols, nrows, item_ct1);
             });
     }
 }
 
-
-static void dequantize_mul_mat_vec_q4_0_sycl(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
+template <int qk, int qr, int qbytes, dequantize_kernel_t_reorder dequantize_kernel>
+static void launch_dequantize_mul_mat_vec_reorder(const void *vx, const dfloat *y, float *dst,
+                                                  const int ncols, const int nrows,
+                                                  dpct::queue_ptr stream) {
     GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
     const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
     // the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
@@ -1526,346 +1526,24 @@ static void dequantize_mul_mat_vec_q4_0_sycl(const void *vx, const dfloat *y,
         stream->parallel_for(
             sycl::nd_range<3>(block_nums * block_dims, block_dims),
             [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                dequantize_mul_mat_vec<QK4_0, QR4_0, dequantize_q4_0>(
+                dequantize_mul_mat_vec_reorder<qk, qr, qbytes, dequantize_kernel>(
                     vx, y, dst, ncols, nrows, item_ct1);
             });
     }
 }
 
-static void dequantize_mul_mat_vec_q1_0_sycl_reorder(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
-    const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
-    // the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
-    {
-        dpct::has_capability_or_fail(stream->get_device(),
-                                     {sycl::aspect::fp16});
-
-        stream->parallel_for(
-            sycl::nd_range<3>(block_nums * block_dims, block_dims),
-            [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                dequantize_mul_mat_vec_reorder<QK1_0, QR1_0, dequantize_q1_0_reorder>(
-                    vx, y, dst, ncols, nrows, item_ct1);
-            });
-    }
-}
-
-static void dequantize_mul_mat_vec_q1_0_sycl(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
-    const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
-    // the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
-    {
-        dpct::has_capability_or_fail(stream->get_device(),
-                                     {sycl::aspect::fp16});
-
-        stream->parallel_for(
-            sycl::nd_range<3>(block_nums * block_dims, block_dims),
-            [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                dequantize_mul_mat_vec<QK1_0, QR1_0, dequantize_q1_0>(
-                    vx, y, dst, ncols, nrows, item_ct1);
-            });
-    }
-}
-
-static void dequantize_mul_mat_vec_q4_1_sycl(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
-    const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
-    {
-        dpct::has_capability_or_fail(stream->get_device(),
-                                     {sycl::aspect::fp16});
-
-        stream->parallel_for(
-            sycl::nd_range<3>(block_nums * block_dims, block_dims),
-            [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                dequantize_mul_mat_vec<QK4_1, QR4_1, dequantize_q4_1>(
-                    vx, y, dst, ncols, nrows, item_ct1);
-            });
-    }
-}
-
-static void dequantize_mul_mat_vec_q5_0_sycl(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
-    const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
-    {
-        dpct::has_capability_or_fail(stream->get_device(),
-                                     {sycl::aspect::fp16});
-
-        stream->parallel_for(
-            sycl::nd_range<3>(block_nums * block_dims, block_dims),
-            [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                dequantize_mul_mat_vec<QK5_0, QR5_0, dequantize_q5_0>(
-                    vx, y, dst, ncols, nrows, item_ct1);
-            });
-    }
-}
-
-static void dequantize_mul_mat_vec_q5_1_sycl(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
-    const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
-    {
-        dpct::has_capability_or_fail(stream->get_device(),
-                                     {sycl::aspect::fp16});
-
-        stream->parallel_for(
-            sycl::nd_range<3>(block_nums * block_dims, block_dims),
-            [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                dequantize_mul_mat_vec<QK5_1, QR5_1, dequantize_q5_1>(
-                    vx, y, dst, ncols, nrows, item_ct1);
-            });
-    }
-}
-
-static void dequantize_mul_mat_vec_q8_0_sycl_reorder(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
-    const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
-    {
-        dpct::has_capability_or_fail(stream->get_device(),
-                                     {sycl::aspect::fp16});
-
-        stream->parallel_for(
-            sycl::nd_range<3>(block_nums * block_dims, block_dims),
-            [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                // Q8_0 reorder layout: [all qs (ncols*nrows bytes)][all d values]
-                // Cannot reuse dequantize_mul_mat_vec_reorder template because it has
-                // Q4_0-specific constants hardcoded (d_ptr offset and qs stride).
-                const int row = item_ct1.get_group(2) * item_ct1.get_local_range(1) +
-                    item_ct1.get_local_id(1);
-                if (row >= nrows) return;
-
-                const int tid = item_ct1.get_local_id(2);
-                const int iter_stride = 8*2*GGML_SYCL_DMMV_X;
-                const int vals_per_iter = iter_stride / WARP_SIZE;
-                const int ncols_left = ncols % (QK8_0*WARP_SIZE);
-                const int ncols_align = ncols - ncols_left;
-
-#ifdef GGML_SYCL_F16
-                sycl::half2 tmp = {0.0f, 0.0f};
-#else
-                float tmp = 0.0f;
-#endif
-                const char *d_ptr = (const char*)vx + ncols*nrows;  // d after all qs
-
-                int i = 0;
-                for (i = 0; i < ncols_align; i += iter_stride) {
-                    const int col = i + vals_per_iter*tid;
-                    const int ib = (row*ncols + col)/QK8_0;
-                    const int iqs = col % QK8_0;
-
-#pragma unroll
-                    for (int j = 0; j < vals_per_iter; j += 2) {
-                        dfloat2 v;
-                        dequantize_q8_0_reorder((const void *)d_ptr, ib, (const void *)vx,
-                                                ib * QK8_0 + iqs + j, v);
-
-#ifdef GGML_SYCL_F16
-                        dfloat2 t1{y[col + j + 0], y[col + j + 1]};
-                        tmp += v * t1;
-#else
-                        tmp += v.x() * y[col + j + 0];
-                        tmp += v.y() * y[col + j + 1];
-#endif
-                    }
-                }
-
-                // handle remaining columns
-                for (; i < ncols; i += iter_stride) {
-                    if (tid >= ncols_left/QK8_0) continue;
-                    const int col = i + vals_per_iter*tid;
-                    const int ib = (row*ncols + col)/QK8_0;
-                    const int iqs = col % QK8_0;
-
-#pragma unroll
-                    for (int j = 0; j < vals_per_iter; j += 2) {
-                        dfloat2 v;
-                        dequantize_q8_0_reorder((const void *)d_ptr, ib, (const void *)vx,
-                                                ib * QK8_0 + iqs + j, v);
-
-#ifdef GGML_SYCL_F16
-                        dfloat2 t1{y[col + j + 0], y[col + j + 1]};
-                        tmp += v * t1;
-#else
-                        tmp += v.x() * y[col + j + 0];
-                        tmp += v.y() * y[col + j + 1];
-#endif
-                    }
-                }
-
-                // reduce
-                const int mask_start = ncols > GGML_SYCL_DMMV_X ? WARP_SIZE >> 1 : WARP_SIZE >> 2;
-                for (int mask = mask_start; mask > 0; mask >>= 1) {
-                    tmp += dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
-                }
-
-                if (tid == 0) {
-#ifdef GGML_SYCL_F16
-                    dst[row] = tmp.x() + tmp.y();
-#else
-                    dst[row] = tmp;
-#endif
-                }
-            });
-    }
-}
-
-static void dequantize_mul_mat_vec_q8_0_sycl(const void *vx, const dfloat *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
-    const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
-    {
-        dpct::has_capability_or_fail(stream->get_device(),
-                                     {sycl::aspect::fp16});
-
-        stream->parallel_for(
-            sycl::nd_range<3>(block_nums * block_dims, block_dims),
-            [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-                dequantize_mul_mat_vec<QK8_0, QR8_0, dequantize_q8_0>(
-                    vx, y, dst, ncols, nrows, item_ct1);
-            });
-    }
-}
-
-static void dequantize_mul_mat_vec_q2_K_sycl(const void *vx, const float *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
+template <int ny, auto kernel>
+static void launch_dequantize_mul_mat_vec_k(const void *vx, const float *y, float *dst,
+                                            const int ncols, const int nrows,
+                                            dpct::queue_ptr stream) {
     GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2; // very slightly faster than 1 even when K_QUANTS_PER_ITERATION = 2
     const int block_num_y = (nrows + ny - 1) / ny;
     const sycl::range<3> block_nums(1, 1, block_num_y);
     const sycl::range<3> block_dims(1, ny, WARP_SIZE);
     stream->parallel_for(
         sycl::nd_range<3>(block_nums * block_dims, block_dims),
         [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q2_k(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
-static void dequantize_mul_mat_vec_q2_K_sycl_reorder(const void *vx, const float *y,
-                                                     float *dst, const int ncols,
-                                                     const int nrows,
-                                                     dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q2_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
-static void dequantize_mul_mat_vec_q3_K_sycl(const void *vx, const float *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q3_k(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
-static void dequantize_mul_mat_vec_q3_K_sycl_reorder(const void *vx, const float *y,
-                                                     float *dst, const int ncols,
-                                                     const int nrows,
-                                                     dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q3_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
-static void dequantize_mul_mat_vec_q4_K_sycl(const void *vx, const float *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q4_k(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
-static void dequantize_mul_mat_vec_q5_K_sycl(const void *vx, const float *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q5_k(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
-static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
-                                             float *dst, const int ncols,
-                                             const int nrows,
-                                             dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q6_k(vx, y, dst, ncols, nrows, item_ct1);
+            kernel(vx, y, dst, ncols, nrows, item_ct1);
         });
 }
 
@@ -2248,54 +1926,6 @@ bool ggml_sycl_dequantize_mul_mat_vec_glu_reorder_esimd(enum ggml_type src0_type
 
 #endif // GGML_SYCL_DMMV_HAS_ESIMD
 
-static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float *y,
-                                                     float *dst, const int ncols,
-                                                     const int nrows,
-                                                     dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q4_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
-static void dequantize_mul_mat_vec_q5_K_sycl_reorder(const void *vx, const float *y,
-                                                     float *dst, const int ncols,
-                                                     const int nrows,
-                                                     dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q5_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
-static void dequantize_mul_mat_vec_q6_K_sycl_reorder(const void *vx, const float *y,
-                                                     float *dst, const int ncols,
-                                                     const int nrows,
-                                                     dpct::queue_ptr stream) {
-    GGML_ASSERT(ncols % QK_K == 0);
-    const int ny = 2 / K_QUANTS_PER_ITERATION;
-    const int block_num_y = (nrows + ny - 1) / ny;
-    const sycl::range<3> block_nums(1, 1, block_num_y);
-    const sycl::range<3> block_dims(1, ny, WARP_SIZE);
-    stream->parallel_for(
-        sycl::nd_range<3>(block_nums * block_dims, block_dims),
-        [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
-            dequantize_mul_mat_vec_q6_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
-        });
-}
-
 void ggml_sycl_op_dequantize_mul_mat_vec(
     ggml_backend_sycl_context & ctx,
     const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst,
@@ -2344,27 +1974,27 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
         case GGML_TYPE_Q1_0:
             if ((ggml_tensor_extra_gpu*)dst->src[0]->extra &&
                 ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
-                dequantize_mul_mat_vec_q1_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec_reorder<QK1_0, QR1_0, QK1_0/8, dequantize_q1_0_reorder>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             } else {
-                dequantize_mul_mat_vec_q1_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec<QK1_0, QR1_0, dequantize_q1_0>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             }
             break;
         case GGML_TYPE_Q4_0:
             if ((ggml_tensor_extra_gpu*)dst->src[0]->extra &&
                 ((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
-                dequantize_mul_mat_vec_q4_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec_reorder<QK4_0, QR4_0, QK4_0/2, dequantize_q4_0_reorder>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             } else {
-                dequantize_mul_mat_vec_q4_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec<QK4_0, QR4_0, dequantize_q4_0>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             }
             break;
         case GGML_TYPE_Q4_1:
-            dequantize_mul_mat_vec_q4_1_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+            launch_dequantize_mul_mat_vec<QK4_1, QR4_1, dequantize_q4_1>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             break;
         case GGML_TYPE_Q5_0:
-            dequantize_mul_mat_vec_q5_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+            launch_dequantize_mul_mat_vec<QK5_0, QR5_0, dequantize_q5_0>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             break;
         case GGML_TYPE_Q5_1:
-            dequantize_mul_mat_vec_q5_1_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+            launch_dequantize_mul_mat_vec<QK5_1, QR5_1, dequantize_q5_1>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             break;
         case GGML_TYPE_Q8_0:
             if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
@@ -2375,10 +2005,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
                 } else
 #endif
                 {
-                    dequantize_mul_mat_vec_q8_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+                    launch_dequantize_mul_mat_vec_reorder<QK8_0, QR8_0, QK8_0, dequantize_q8_0_reorder>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
                 }
             } else {
-                dequantize_mul_mat_vec_q8_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec<QK8_0, QR8_0, dequantize_q8_0>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
             }
             break;
         case GGML_TYPE_Q2_K:
@@ -2391,10 +2021,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
                 else
 #endif
                 {
-                    dequantize_mul_mat_vec_q2_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                    launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q2_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
                 }
             } else {
-                dequantize_mul_mat_vec_q2_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec_k<2, dequantize_mul_mat_vec_q2_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }
             break;
         case GGML_TYPE_Q3_K:
@@ -2407,10 +2037,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
                 else
 #endif
                 {
-                    dequantize_mul_mat_vec_q3_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                    launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q3_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
                 }
             } else {
-                dequantize_mul_mat_vec_q3_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q3_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }
             break;
         case GGML_TYPE_Q4_K:
@@ -2423,10 +2053,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
                 else
 #endif
                 {
-                    dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                    launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q4_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
                 }
             } else {
-                dequantize_mul_mat_vec_q4_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q4_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }
             break;
         case GGML_TYPE_Q5_K:
@@ -2439,10 +2069,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
                 else
 #endif
                 {
-                    dequantize_mul_mat_vec_q5_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                    launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q5_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
                 }
             } else {
-                dequantize_mul_mat_vec_q5_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q5_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }
             break;
         case GGML_TYPE_Q6_K:
@@ -2455,10 +2085,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
                 else
 #endif
                 {
-                    dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                    launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q6_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
                 }
             } else {
-                dequantize_mul_mat_vec_q6_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q6_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }
             break;
         case GGML_TYPE_F16:

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -7829,9 +7829,15 @@ struct ggml_backend_sycl_comm_context {
 
     // pinned host staging per device context, reused across calls (queues are in-order)
     uint8_t * host_out[2] = { nullptr, nullptr };
-    uint8_t * host_in[2]  = { nullptr, nullptr };
+    uint8_t * host_in[2]  = { nullptr, nullptr };  // only used by the 4-buffer path (separate SYCL contexts)
     size_t    host_bytes  = 0;
 
+    // SYCL_ENABLE_DEFAULT_CONTEXTS=1 sets comm_exchange_fast on Linux - PR #30956
+    using exchange_fn_t = void (*)(ggml_backend_sycl_comm_context * comm, queue_ptr q0, queue_ptr q1,
+                                   void * dst0, const void * src1, void * dst1, const void * src0,
+                                   size_t nbytes);
+    exchange_fn_t comm_exchange = nullptr; // Moves same context check out of hot path; tests only once
+
     void free_host() {
         for (int i = 0; i < 2; ++i) {
             const sycl::context sctx = ((ggml_backend_sycl_context *) backends[i]->context)->stream()->get_context();
@@ -7842,6 +7848,41 @@ struct ggml_backend_sycl_comm_context {
     }
 };
 
+// dst0 <- src1 and dst1 <- src0 through pinned host memory.
+static void ggml_sycl_comm_exchange_fast(ggml_backend_sycl_comm_context * comm, queue_ptr q0, queue_ptr q1,
+                                         void * dst0, const void * src1, void * dst1, const void * src0,
+                                         size_t nbytes) {
+
+    // Increases Linux tensor split prompt eval speeds on par with layer split speeds
+    // Reallocate host buffers if they are too small
+    if (comm->host_bytes < nbytes) {
+        q0->wait();
+        q1->wait();
+        comm->free_host();
+        comm->host_out[0] = sycl::malloc_host<uint8_t>(nbytes, q0->get_context());
+        comm->host_out[1] = sycl::malloc_host<uint8_t>(nbytes, q1->get_context());
+
+        GGML_ASSERT(comm->host_out[0] && comm->host_out[1]);
+        comm->host_bytes = nbytes;
+    }
+
+    // Step 1: Concurrently copy from GPUs into host buffers (Asynchronous)
+    sycl::event e_d2h_0 = q0->memcpy(comm->host_out[0], src0, nbytes);
+    sycl::event e_d2h_1 = q1->memcpy(comm->host_out[1], src1, nbytes);
+
+    // Step 2: Asynchronous Barrier Interlocking (NO HOST BLOCKING), wait() causes MTP conflicts
+    q0->ext_oneapi_submit_barrier({e_d2h_1}); // q0 waits for q1's device-to-host copy
+    q1->ext_oneapi_submit_barrier({e_d2h_0}); // q1 waits for q0's device-to-host copy
+
+    // Step 3: Cross-upload directly from the opposing host buffer.
+    sycl::event e_h2d_0 = q0->memcpy(dst0, comm->host_out[1], nbytes); // q0 pulls from host_out[1]
+    sycl::event e_h2d_1 = q1->memcpy(dst1, comm->host_out[0], nbytes); // q1 pulls from host_out[0]
+
+    // Step 4: Finalize the exchange before returning to host code
+    e_h2d_0.wait();
+    e_h2d_1.wait();
+}
+
 // dst0 <- src1 and dst1 <- src0 through pinned host memory, with a single host wait
 static void ggml_sycl_comm_exchange(ggml_backend_sycl_comm_context * comm, queue_ptr q0, queue_ptr q1,
                                     void * dst0, const void * src1, void * dst1, const void * src0,
@@ -7887,6 +7928,15 @@ void * ggml_backend_sycl_comm_init(ggml_backend_t * backends, size_t n_backends)
     auto * sctx1 = (ggml_backend_sycl_context *) backends[1]->context;
     ctx->buf0 = std::make_unique<ggml_sycl_pool_alloc<uint8_t>>(sctx0->pool());
     ctx->buf1 = std::make_unique<ggml_sycl_pool_alloc<uint8_t>>(sctx1->pool());
+
+#if defined(_WIN32) || defined(_WIN64) // Not supported on Windows
+    ctx->comm_exchange = ggml_sycl_comm_exchange;
+#else
+    // SYCL_ENABLE_DEFAULT_CONTEXTS=1 sets comm_exchange_fast on Linux - PR #30956
+    ctx->comm_exchange = (sctx0->stream()->get_context() == sctx1->stream()->get_context())
+                       ? ggml_sycl_comm_exchange_fast
+                       : ggml_sycl_comm_exchange;
+#endif
     return ctx;
 }
 catch (const sycl::exception &) { return nullptr; }
@@ -7970,7 +8020,7 @@ bool ggml_backend_sycl_comm_allreduce_tensor(void * comm_ctx_v, struct ggml_tens
         sycl::half * f16_tmp0 = (sycl::half *) buf0;
         sycl::half * f16_tmp1 = (sycl::half *) buf1;
 
-        ggml_sycl_comm_exchange(comm_ctx, q0, q1, f16_tmp0, tensors[1]->data, f16_tmp1, tensors[0]->data, nbytes);
+        comm_ctx->comm_exchange(comm_ctx, q0, q1, f16_tmp0, tensors[1]->data, f16_tmp1, tensors[0]->data, nbytes);
 
         q0->submit([&](sycl::handler & h) {
             h.parallel_for(sycl::range<1>(nelem), [=](sycl::id<1> i) {
@@ -8000,7 +8050,7 @@ bool ggml_backend_sycl_comm_allreduce_tensor(void * comm_ctx_v, struct ggml_tens
         float * tmp1 = (float *) buf1;
 
         // separate SYCL contexts: a raw peer-USM q->memcpy is a silent no-op, so stage through host memory
-        ggml_sycl_comm_exchange(comm_ctx, q0, q1, tmp0, tensors[1]->data, tmp1, tensors[0]->data, nbytes);
+        comm_ctx->comm_exchange(comm_ctx, q0, q1, tmp0, tensors[1]->data, tmp1, tensors[0]->data, nbytes);
 
         q0->submit([&](sycl::handler & h) {
             h.parallel_for(sycl::range<1>(nelem), [=](sycl::id<1> i) {
@@ -8034,7 +8084,7 @@ bool ggml_backend_sycl_comm_allreduce_tensor(void * comm_ctx_v, struct ggml_tens
 
     // Phase B: exchange the compressed bytes (in-order queues keep this after the compress kernels).
     const size_t bf16_bytes = nelem * sizeof(uint16_t);
-    ggml_sycl_comm_exchange(comm_ctx, q0, q1, inbox0, outbox1, inbox1, outbox0, bf16_bytes);
+    comm_ctx->comm_exchange(comm_ctx, q0, q1, inbox0, outbox1, inbox1, outbox0, bf16_bytes);
 
     // Phase C: decompress + add into local FP32 partial.
     q0->submit([&](sycl::handler & h) {

ggml/src/ggml-vulkan/CMakeLists.txt

@@ -66,6 +66,7 @@ if (Vulkan_FOUND)
                              ggml-vulkan-push-constants.h
                              ggml-vulkan-common.h
                              ggml-vulkan-buffers.cpp
+                             ggml-vulkan-comm.cpp
                              ggml-vulkan-debug.cpp
                             )
 

ggml/src/ggml-vulkan/ggml-vulkan-comm.cpp

@@ -0,0 +1,591 @@
+#include "ggml-vulkan-common.h"
+
+#if !defined(_WIN32)
+#include <fcntl.h>
+#include <unistd.h>
+#if defined(__linux__)
+#include <pthread.h>
+#endif
+
+static constexpr size_t ggml_vk_comm_f16_min_size = 1024 * 1024;
+
+// follows ops ggml-backend-meta may delay a reduction across
+// TODO: the required precision should come from the meta backend instead.
+static bool ggml_vk_comm_needs_f32(const ggml_tensor * t, int depth = 0) {
+    if (t == nullptr || depth > 64) {
+        return false;
+    }
+    switch (t->op) {
+        case GGML_OP_MUL_MAT:
+        case GGML_OP_MUL_MAT_ID:
+            return ggml_get_op_params_i32(t, 0) == GGML_PREC_F32;
+        case GGML_OP_ADD:
+            return ggml_vk_comm_needs_f32(t->src[0], depth + 1) || ggml_vk_comm_needs_f32(t->src[1], depth + 1);
+        case GGML_OP_MUL:
+        case GGML_OP_ADD_ID:
+            return ggml_vk_comm_needs_f32(t->src[0], depth + 1);
+        case GGML_OP_VIEW:
+            return ggml_vk_comm_needs_f32(t->src[0], depth + 1);
+        default:
+            return false;
+    }
+}
+
+// host_bridge: device i signals done[i], a host thread signals go[i] once all peers of i are done
+// sync_fd: a host thread exchanges SYNC_FD semaphores between peers, the wait resolves in the kernel
+enum class vk_comm_sync {
+    host_bridge,
+    sync_fd,
+};
+
+struct vk_comm_device {
+    ggml_backend_t                     backend  = nullptr;
+    ggml_backend_vk_context *          ctx      = nullptr;
+    void *                             host_ptr = nullptr;
+    std::vector<ggml_backend_buffer_t> host_buf; // host_ptr imported into each peer
+
+    struct {
+        PFN_vkGetSemaphoreFdKHR    get_fd    = nullptr;
+        PFN_vkImportSemaphoreFdKHR import_fd = nullptr;
+        vk::Semaphore              sem;
+        std::vector<vk::Semaphore> imp;      // [peer], pooled, reused only after ggml_vk_graph_cleanup
+    } sfd;
+
+    struct {
+        vk::Semaphore done;
+        vk::Semaphore go;
+    } bridge;
+};
+
+struct ggml_backend_vk_comm_context {
+    std::vector<vk_comm_device> devs;
+    vk_comm_sync                mode     = vk_comm_sync::host_bridge;
+    size_t                      align    = 4096;
+    size_t                      cap      = 0;
+    uint64_t                    round    = 0;
+    ggml_context *              tctx     = nullptr;
+    ggml_tensor *               slot     = nullptr;
+    bool                        disabled = false;
+
+    std::thread                 thread;
+    std::chrono::nanoseconds    spin{ 200'000 };
+    std::atomic<uint64_t>       submitted{ 0 };
+    std::atomic<uint64_t>       imported{ 0 };
+    std::atomic<bool>           sleeping{ false };
+    std::atomic<bool>           stop{ false };
+    std::mutex                  mutex;
+    std::condition_variable     cv;
+};
+
+static void ggml_vk_comm_free_buffers(ggml_backend_vk_comm_context * comm) {
+    for (vk_comm_device & d : comm->devs) {
+        for (ggml_backend_buffer_t & buf : d.host_buf) {
+            ggml_backend_buffer_free(buf);
+            buf = nullptr;
+        }
+        std::free(d.host_ptr);
+        d.host_ptr = nullptr;
+    }
+    comm->cap = 0;
+}
+
+static bool ggml_backend_vk_comm_ensure(ggml_backend_vk_comm_context * comm, size_t nbytes) {
+    if (comm->disabled) {
+        return false;
+    }
+    if (nbytes <= comm->cap) {
+        return true;
+    }
+    auto fail = [comm]() {
+        GGML_LOG_WARN("ggml_vulkan: shared host buffer allocation failed, using generic all-reduce\n");
+        comm->disabled = true;
+        return false;
+    };
+    for (vk_comm_device & d : comm->devs) {
+        ggml_backend_synchronize(d.backend);
+    }
+    ggml_vk_comm_free_buffers(comm);
+
+    const size_t n   = comm->devs.size();
+    const size_t cap = GGML_PAD(nbytes, comm->align);
+    for (size_t i = 0; i < n; i++) {
+        vk_comm_device & d = comm->devs[i];
+        if (posix_memalign(&d.host_ptr, comm->align, 2 * cap) != 0) {
+            d.host_ptr = nullptr;
+            return fail();
+        }
+        for (size_t j = 0; j < n; j++) {
+            d.host_buf[j] = ggml_backend_vk_device_buffer_from_host_ptr(
+                ggml_backend_get_device(comm->devs[j].backend), d.host_ptr, 2 * cap, 2 * cap);
+            if (!d.host_buf[j]) {
+                return fail();
+            }
+        }
+    }
+    comm->cap = cap;
+    return true;
+}
+
+static void ggml_vk_comm_set_slot(ggml_tensor * slot, const ggml_tensor * src, ggml_type type, ggml_backend_buffer_t buf, size_t offs) {
+    slot->type  = type;
+    slot->nb[0] = ggml_type_size(type);
+    for (int d = 0; d < GGML_MAX_DIMS; d++) {
+        slot->ne[d] = src->ne[d];
+        if (d > 0) {
+            slot->nb[d] = slot->nb[d - 1] * slot->ne[d - 1];
+        }
+    }
+    slot->buffer = buf;
+    slot->data   = (char *) ggml_backend_buffer_get_base(buf) + offs;
+}
+
+static void ggml_vk_comm_submit(ggml_backend_vk_context * ctx, vk_context & c, vk::Semaphore signal, uint64_t value) {
+    ggml_vk_ctx_end(c);
+    for (auto & cpy : c->in_memcpys) {
+        memcpy(cpy.dst, cpy.src, cpy.n);
+    }
+    for (auto & mset : c->memsets) {
+        memset(mset.dst, mset.val, mset.n);
+    }
+    c->in_memcpys.clear();
+    c->memsets.clear();
+    if (signal) {
+        c->seqs.back().back().signal_semaphores.push_back({ signal, value });
+    }
+    ggml_vk_submit(c, {});
+    ctx->submit_pending = true;
+    ctx->compute_ctx.reset();
+}
+
+static void ggml_vk_comm_notify_submitted(ggml_backend_vk_comm_context * comm, uint64_t round) {
+    comm->submitted.store(round);
+    if (comm->sleeping.load()) {
+        std::lock_guard<std::mutex> lock(comm->mutex);
+        comm->cv.notify_one();
+    }
+}
+
+// false on stop
+static bool ggml_vk_comm_wait_submitted(ggml_backend_vk_comm_context * comm, uint64_t r) {
+    using clock = std::chrono::steady_clock;
+    for (const auto t0 = clock::now(); comm->submitted.load() < r;) {
+        if (comm->stop.load()) {
+            return false;
+        }
+        if (clock::now() - t0 < comm->spin) {
+            YIELD();
+            continue;
+        }
+        comm->sleeping.store(true);
+        if (comm->submitted.load() < r) {
+            std::unique_lock<std::mutex> lock(comm->mutex);
+            comm->cv.wait(lock, [&] { return comm->submitted.load() >= r || comm->stop.load(); });
+        }
+        comm->sleeping.store(false);
+    }
+    return true;
+}
+
+// signals go[i] once all peers of i are done; false on stop
+static bool ggml_vk_comm_bridge_release(ggml_backend_vk_comm_context * comm, uint64_t r, std::vector<uint64_t> & seen, std::vector<bool> & sent) {
+    using clock = std::chrono::steady_clock;
+    const size_t n = comm->devs.size();
+    std::fill(sent.begin(), sent.end(), false);
+    size_t left = n;
+    for (const auto t0 = clock::now();;) {
+        size_t pending = n;
+        for (size_t i = 0; i < n; i++) {
+            if (sent[i]) {
+                continue;
+            }
+            bool ready = true;
+            for (size_t j = 0; j < n && ready; j++) {
+                if (j != i && seen[j] < r) {
+                    const vk_comm_device & p = comm->devs[j];
+                    seen[j] = p.ctx->device->device.getSemaphoreCounterValue(p.bridge.done);
+                    ready   = seen[j] >= r;
+                    pending = ready ? pending : j;
+                }
+            }
+            if (ready) {
+                const vk_comm_device & d = comm->devs[i];
+                d.ctx->device->device.signalSemaphore({ d.bridge.go, r });
+                sent[i] = true;
+                left--;
+            }
+        }
+        if (left == 0) {
+            return true;
+        }
+        if (clock::now() - t0 < comm->spin) {
+            YIELD();
+            continue;
+        }
+        const vk_comm_device &      p = comm->devs[pending];
+        const vk::SemaphoreWaitInfo wi{ {}, 1, &p.bridge.done, &r };
+        while (p.ctx->device->device.waitSemaphores(wi, 100'000'000) != vk::Result::eSuccess) {
+            if (comm->stop.load()) {
+                return false;
+            }
+        }
+    }
+}
+
+static void ggml_vk_comm_bridge_thread(ggml_backend_vk_comm_context * comm) {
+#if defined(__linux__)
+    pthread_setname_np(pthread_self(), "ggml-vk-comm");
+#endif
+    std::vector<uint64_t> seen(comm->devs.size(), 0);
+    std::vector<bool>     sent(comm->devs.size());
+    try {
+        for (uint64_t r = 1;; r++) {
+            if (!ggml_vk_comm_wait_submitted(comm, r) || !ggml_vk_comm_bridge_release(comm, r, seen, sent)) {
+                return;
+            }
+        }
+    } catch (const vk::SystemError & e) {
+        GGML_ABORT("ggml_vulkan: all-reduce bridge failed: %s", e.what());
+    }
+}
+
+static void ggml_vk_comm_bridge_init(ggml_backend_vk_comm_context * comm) {
+    vk::SemaphoreTypeCreateInfo stci{ vk::SemaphoreType::eTimeline, 0 };
+    vk::SemaphoreCreateInfo     sci{};
+    sci.pNext = &stci;
+    for (vk_comm_device & d : comm->devs) {
+        d.bridge.done = d.ctx->device->device.createSemaphore(sci);
+        d.bridge.go   = d.ctx->device->device.createSemaphore(sci);
+    }
+    comm->thread = std::thread(ggml_vk_comm_bridge_thread, comm);
+}
+
+static void ggml_vk_comm_bridge_free(vk_comm_device & d) {
+    if (d.bridge.done) {
+        d.ctx->device->device.destroySemaphore(d.bridge.done);
+    }
+    if (d.bridge.go) {
+        d.ctx->device->device.destroySemaphore(d.bridge.go);
+    }
+}
+
+static void ggml_vk_comm_barrier_to_host(vk_context & c) {
+    const vk::MemoryBarrier to_host{ vk::AccessFlagBits::eTransferWrite | vk::AccessFlagBits::eShaderWrite, vk::AccessFlagBits::eHostRead };
+    c->s->buffer->buf.pipelineBarrier(c->p->q->stage_flags, vk::PipelineStageFlagBits::eHost, {}, { to_host }, {}, {});
+}
+
+// the peers' slots were made available to the host before the wait resolved, the submit predates that
+static void ggml_vk_comm_barrier_from_host(vk_context & c) {
+    const vk::MemoryBarrier from_host{ vk::AccessFlagBits::eHostWrite, vk::AccessFlagBits::eShaderRead | vk::AccessFlagBits::eTransferRead };
+    c->s->buffer->buf.pipelineBarrier(vk::PipelineStageFlagBits::eHost, c->p->q->stage_flags, {}, { from_host }, {}, {});
+}
+
+static void ggml_vk_comm_bridge_signal(ggml_backend_vk_comm_context * comm, size_t i, vk_context & c, uint64_t round) {
+    const vk_comm_device & d = comm->devs[i];
+    ggml_vk_comm_barrier_to_host(c);
+    ggml_vk_comm_submit(d.ctx, c, d.bridge.done, round);
+}
+
+static void ggml_vk_comm_bridge_wait(ggml_backend_vk_comm_context * comm, size_t i, vk_context & c, uint64_t round) {
+    c->s->wait_semaphores.push_back({ comm->devs[i].bridge.go, round });
+    ggml_vk_comm_barrier_from_host(c);
+}
+
+static bool ggml_vk_comm_sfd_supported(vk_comm_device & d) {
+    const vk_device & dev = d.ctx->device;
+    const bool mesa = dev->driver_id == vk::DriverId::eMesaRadv ||
+                      dev->driver_id == vk::DriverId::eIntelOpenSourceMESA ||
+                      dev->driver_id == vk::DriverId::eMesaNvk;
+    if (!mesa || !dev->external_semaphore) {
+        return false;
+    }
+    d.sfd.get_fd    = (PFN_vkGetSemaphoreFdKHR)    dev->device.getProcAddr("vkGetSemaphoreFdKHR");
+    d.sfd.import_fd = (PFN_vkImportSemaphoreFdKHR) dev->device.getProcAddr("vkImportSemaphoreFdKHR");
+    if (!d.sfd.get_fd || !d.sfd.import_fd) {
+        return false;
+    }
+    vk::PhysicalDeviceExternalSemaphoreInfo esi{ vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd };
+    vk::ExternalSemaphoreProperties esp = dev->physical_device.getExternalSemaphoreProperties(esi);
+    const vk::ExternalSemaphoreFeatureFlags need =
+        vk::ExternalSemaphoreFeatureFlagBits::eExportable | vk::ExternalSemaphoreFeatureFlagBits::eImportable;
+    return (esp.externalSemaphoreFeatures & need) == need &&
+           (esp.compatibleHandleTypes & vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd);
+}
+
+// Mesa moves a queue to its submit thread on the first wait-before-signal, taking the CS ioctls off this thread
+static void ggml_vk_comm_sfd_enable_submit_thread(const vk_device & dev) {
+    vk::SemaphoreTypeCreateInfo stci{ vk::SemaphoreType::eTimeline, 0 };
+    vk::SemaphoreCreateInfo     sci{};
+    sci.pNext = &stci;
+    const vk::Semaphore sem   = dev->device.createSemaphore(sci);
+    const vk::Fence     fence = dev->device.createFence({});
+    const uint64_t      one   = 1;
+    vk::TimelineSemaphoreSubmitInfo tsi{ 1, &one, 0, nullptr };
+    const vk::PipelineStageFlags stage = vk::PipelineStageFlagBits::eAllCommands;
+    vk::SubmitInfo si{ 1, &sem, &stage, 0, nullptr, 0, nullptr };
+    si.setPNext(&tsi);
+    dev->compute_queue->handle->submit(si, fence);
+    dev->device.signalSemaphore({ sem, 1 });
+    VK_CHECK(dev->device.waitForFences({ fence }, true, UINT64_MAX), "ggml_vk_comm_sfd_enable_submit_thread", dev);
+    dev->device.destroyFence(fence);
+    dev->device.destroySemaphore(sem);
+}
+
+// Requires a submitted signal on sem; the export resets it.
+static int ggml_vk_comm_sfd_export(vk_comm_device & d) {
+    int fd = -1;
+    VkSemaphoreGetFdInfoKHR gi{};
+    gi.sType      = VK_STRUCTURE_TYPE_SEMAPHORE_GET_FD_INFO_KHR;
+    gi.semaphore  = d.sfd.sem;
+    gi.handleType = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT;
+    const VkResult result = d.sfd.get_fd((VkDevice) d.ctx->device->device, &gi, &fd);
+    if (result != VK_SUCCESS) {
+        GGML_ABORT("ggml_vulkan: vkGetSemaphoreFdKHR failed: %d", (int) result);
+    }
+    return fd;
+}
+
+// fd -1 imports an already signalled payload
+static void ggml_vk_comm_sfd_import(vk_comm_device & d, vk::Semaphore sem, int fd) {
+    VkImportSemaphoreFdInfoKHR isi{};
+    isi.sType      = VK_STRUCTURE_TYPE_IMPORT_SEMAPHORE_FD_INFO_KHR;
+    isi.semaphore  = sem;
+    isi.flags      = VK_SEMAPHORE_IMPORT_TEMPORARY_BIT;
+    isi.handleType = VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_SYNC_FD_BIT;
+    isi.fd         = fd;
+    const VkResult result = d.sfd.import_fd((VkDevice) d.ctx->device->device, &isi);
+    if (result != VK_SUCCESS) {
+        GGML_ABORT("ggml_vulkan: vkImportSemaphoreFdKHR failed: %d", (int) result);
+    }
+}
+
+static void ggml_vk_comm_sfd_exchange(ggml_backend_vk_comm_context * comm, uint64_t r) {
+    const size_t n = comm->devs.size();
+    for (size_t j = 0; j < n; j++) {
+        const int    fd   = ggml_vk_comm_sfd_export(comm->devs[j]);
+        const size_t last = j == n - 1 ? n - 2 : n - 1;
+        for (size_t i = 0; i < n; i++) {
+            if (i == j) {
+                continue;
+            }
+            const int imp_fd = fd < 0 || i == last ? fd : fcntl(fd, F_DUPFD_CLOEXEC, 0);
+            GGML_ASSERT(imp_fd >= 0 || fd < 0);
+            ggml_vk_comm_sfd_import(comm->devs[i], comm->devs[i].sfd.imp[j], imp_fd);
+        }
+    }
+    comm->imported.store(r, std::memory_order_release);
+}
+
+static void ggml_vk_comm_sfd_thread(ggml_backend_vk_comm_context * comm) {
+#if defined(__linux__)
+    pthread_setname_np(pthread_self(), "ggml-vk-comm");
+#endif
+    for (uint64_t r = 1; ggml_vk_comm_wait_submitted(comm, r); r++) {
+        ggml_vk_comm_sfd_exchange(comm, r);
+    }
+}
+
+static void ggml_vk_comm_sfd_init(ggml_backend_vk_comm_context * comm) {
+    vk::ExportSemaphoreCreateInfo esci{ vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd };
+    vk::SemaphoreCreateInfo       sci{};
+    sci.pNext = &esci;
+    for (vk_comm_device & d : comm->devs) {
+        d.sfd.sem = d.ctx->device->device.createSemaphore(sci);
+        d.sfd.imp.assign(comm->devs.size(), {});
+        ggml_vk_comm_sfd_enable_submit_thread(d.ctx->device);
+    }
+    comm->thread = std::thread(ggml_vk_comm_sfd_thread, comm);
+}
+
+static void ggml_vk_comm_sfd_free(vk_comm_device & d) {
+    if (d.sfd.sem) {
+        d.ctx->device->device.destroySemaphore(d.sfd.sem);
+    }
+}
+
+static void ggml_vk_comm_sfd_signal(ggml_backend_vk_comm_context * comm, size_t i, vk_context & c) {
+    vk_comm_device & d = comm->devs[i];
+    ggml_vk_comm_barrier_to_host(c);
+    ggml_vk_comm_submit(d.ctx, c, d.sfd.sem, 0);
+    for (size_t j = 0; j < comm->devs.size(); j++) {
+        d.sfd.imp[j] = j == i ? vk::Semaphore{} : ggml_vk_create_binary_semaphore(d.ctx);
+    }
+}
+
+static void ggml_vk_comm_sfd_wait(ggml_backend_vk_comm_context * comm, size_t i, vk_context & c, uint64_t round) {
+    for (size_t j = 0; j < comm->devs.size(); j++) {
+        if (j != i) {
+            c->s->wait_semaphores.push_back({ comm->devs[i].sfd.imp[j], 0 });
+        }
+    }
+    c->s->ready       = &comm->imported;
+    c->s->ready_value = round;
+    ggml_vk_comm_barrier_from_host(c);
+}
+
+void ggml_backend_vk_comm_free(void * comm_ctx) {
+    ggml_backend_vk_comm_context * comm = static_cast<ggml_backend_vk_comm_context *>(comm_ctx);
+    // pending reductions need the thread
+    for (vk_comm_device & d : comm->devs) {
+        ggml_backend_synchronize(d.backend);
+    }
+    if (comm->thread.joinable()) {
+        {
+            std::lock_guard<std::mutex> lock(comm->mutex);
+            comm->stop.store(true);
+        }
+        comm->cv.notify_all();
+        comm->thread.join();
+    }
+    for (vk_comm_device & d : comm->devs) {
+        ggml_vk_comm_sfd_free(d);
+        ggml_vk_comm_bridge_free(d);
+    }
+    ggml_vk_comm_free_buffers(comm);
+    if (comm->tctx) {
+        ggml_free(comm->tctx);
+    }
+    delete comm;
+}
+
+void * ggml_backend_vk_comm_init(ggml_backend_t * backends, size_t n_backends) {
+    if (vk_instance.allreduce_disabled || n_backends < 2) {
+        return nullptr;
+    }
+    for (size_t i = 0; i < n_backends; i++) {
+        if (!ggml_backend_is_vk(backends[i])) {
+            return nullptr;
+        }
+    }
+    ggml_backend_vk_comm_context * comm = new ggml_backend_vk_comm_context;
+    comm->devs.resize(n_backends);
+    for (size_t i = 0; i < n_backends; i++) {
+        vk_comm_device & d = comm->devs[i];
+        d.backend = backends[i];
+        d.ctx     = (ggml_backend_vk_context *) backends[i]->context;
+        d.host_buf.assign(n_backends, nullptr);
+    }
+
+    for (size_t i = 0; i < n_backends; i++) {
+        const vk_device & dev = comm->devs[i].ctx->device;
+        comm->align = std::max(comm->align, (size_t) dev->min_imported_host_pointer_alignment);
+        if (!dev->external_memory_host) {
+            GGML_LOG_WARN("ggml_vulkan: device %zu: VK_EXT_external_memory_host missing, using generic all-reduce\n", i);
+            ggml_backend_vk_comm_free(comm);
+            return nullptr;
+        }
+        if (dev->serialize_submissions) {
+            GGML_LOG_WARN("ggml_vulkan: device %zu: GGML_VK_SERIALIZE_SUBMISSIONS set, using generic all-reduce\n", i);
+            ggml_backend_vk_comm_free(comm);
+            return nullptr;
+        }
+    }
+
+    if (std::all_of(comm->devs.begin(), comm->devs.end(), ggml_vk_comm_sfd_supported)) {
+        comm->mode = vk_comm_sync::sync_fd;
+    }
+
+    if (const char * spin = getenv("GGML_VK_ALLREDUCE_SPIN_US")) {
+        comm->spin = std::chrono::microseconds(std::max(0, atoi(spin)));
+    }
+    const ggml_init_params ip = { ggml_tensor_overhead(), nullptr, true };
+    comm->tctx = ggml_init(ip);
+    comm->slot = ggml_new_tensor_1d(comm->tctx, GGML_TYPE_F32, 1);
+
+    try {
+        if (comm->mode == vk_comm_sync::sync_fd) {
+            ggml_vk_comm_sfd_init(comm);
+        } else {
+            ggml_vk_comm_bridge_init(comm);
+        }
+    } catch (const vk::SystemError & e) {
+        GGML_LOG_WARN("ggml_vulkan: semaphore initialization failed: %s, using generic all-reduce\n", e.what());
+        ggml_backend_vk_comm_free(comm);
+        return nullptr;
+    }
+    return comm;
+}
+
+// A: partial to own host slot (round & 1), signal; B: wait on peers' A, add their slots; all A precede any B
+// slot reuse: A(k) follows own B(k-1), which waited on peers' A(k-1), submitted after their B(k-2), the last reader
+bool ggml_backend_vk_comm_allreduce_tensor(void * comm_ctx, ggml_tensor ** tensors) {
+    ggml_backend_vk_comm_context * comm = static_cast<ggml_backend_vk_comm_context *>(comm_ctx);
+    const size_t n = comm->devs.size();
+
+    const int64_t ne = ggml_nelements(tensors[0]);
+    for (size_t i = 0; i < n; i++) {
+        const ggml_tensor * t = tensors[i];
+        if (t->type != GGML_TYPE_F32 || ggml_nelements(t) != ne || !ggml_is_contiguous(t) || !ggml_backend_buffer_is_vk(t->buffer)) {
+            return false;
+        }
+        const ggml_backend_vk_context * ctx = comm->devs[i].ctx;
+        if (!ctx->compute_ctx.expired() && !ctx->compute_ctx.lock()->out_memcpys.empty()) {
+            return false;
+        }
+    }
+    if (ne == 0) {
+        return true;
+    }
+    const size_t nbytes = ggml_nbytes(tensors[0]);
+    if (!ggml_backend_vk_comm_ensure(comm, nbytes)) {
+        return false;
+    }
+
+    const bool      f32_wire  = nbytes < ggml_vk_comm_f16_min_size || ggml_vk_comm_needs_f32(tensors[0]);
+    const ggml_type wire      = f32_wire ? GGML_TYPE_F32 : GGML_TYPE_F16;
+    const uint64_t  round     = ++comm->round;
+    const size_t    slot_offs = (round & 1) * comm->cap;
+    const size_t    slot_size = ne * ggml_type_size(wire);
+
+    for (size_t i = 0; i < n; i++) {
+        const vk_comm_device &    d   = comm->devs[i];
+        ggml_backend_vk_context * ctx = d.ctx;
+        ggml_tensor *             t   = tensors[i];
+        vk_buffer & tbuf = ((ggml_backend_vk_buffer_context *) t->buffer->context)->dev_buffer;
+        vk_buffer & hbuf = ((ggml_backend_vk_buffer_context *) d.host_buf[i]->context)->dev_buffer;
+        const size_t toffs = vk_tensor_offset(t) + t->view_offs;
+
+        vk_context c = ggml_vk_get_compute_ctx(ctx);
+        ggml_vk_sync_buffers(ctx, c);
+        if (!(t->flags & GGML_TENSOR_FLAG_COMPUTE)) {
+            c->s->buffer->buf.fillBuffer(hbuf->buffer, slot_offs, GGML_PAD(slot_size, 4), 0);
+            c->s->buffer->buf.fillBuffer(tbuf->buffer, toffs, nbytes, 0);
+        } else if (wire == GGML_TYPE_F32) {
+            ggml_vk_buffer_copy_async(c, hbuf, slot_offs, tbuf, toffs, nbytes);
+        } else {
+            ggml_vk_comm_set_slot(comm->slot, t, wire, d.host_buf[i], slot_offs);
+            ggml_vk_cpy(ctx, c, t, comm->slot);
+        }
+        if (comm->mode == vk_comm_sync::sync_fd) {
+            ggml_vk_comm_sfd_signal(comm, i, c);
+        } else {
+            ggml_vk_comm_bridge_signal(comm, i, c, round);
+        }
+    }
+    ggml_vk_comm_notify_submitted(comm, round);
+
+    for (size_t i = 0; i < n; i++) {
+        ggml_backend_vk_context * ctx = comm->devs[i].ctx;
+        vk_context c = ggml_vk_get_compute_ctx(ctx);
+        ctx->do_add_rms_partials = false;
+        if (comm->mode == vk_comm_sync::sync_fd) {
+            ggml_vk_comm_sfd_wait(comm, i, c, round);
+        } else {
+            ggml_vk_comm_bridge_wait(comm, i, c, round);
+        }
+        ggml_vk_sync_buffers(ctx, c);
+        for (size_t j = 0; j < n; j++) {
+            if (j == i) {
+                continue;
+            }
+            ggml_vk_comm_set_slot(comm->slot, tensors[i], wire, comm->devs[j].host_buf[i], slot_offs);
+            ggml_vk_add(ctx, c, tensors[i], comm->slot, tensors[i]);
+            ggml_vk_sync_buffers(ctx, c);
+        }
+        // small reductions go out with the next graph's first submit
+        if (wire != GGML_TYPE_F32 || vk_perf_logger_enabled) {
+            ggml_vk_comm_submit(ctx, c, {}, 0);
+        }
+    }
+    return true;
+}
+#endif

ggml/src/ggml-vulkan/ggml-vulkan-common.h

@@ -86,6 +86,7 @@ void ggml_vk_cmd_label_begin(vk::CommandBuffer buf, const char * name);
 void ggml_vk_ctx_end(vk_context& ctx);
 void ggml_vk_ctx_begin(vk_device& device, vk_context& subctx);
 vk_context ggml_vk_get_compute_ctx(ggml_backend_vk_context * ctx);
+vk::Semaphore ggml_vk_create_binary_semaphore(ggml_backend_vk_context * ctx);
 vk_context ggml_vk_get_transfer_ctx(ggml_backend_vk_context * ctx);
 bool ggml_vk_submit_transfer_ctx(ggml_backend_vk_context * ctx);
 size_t ggml_vk_align_size(size_t width, size_t align);
@@ -233,8 +234,16 @@ size_t ggml_backend_vk_buffer_type_get_alignment(ggml_backend_buffer_type_t buft
 size_t ggml_backend_vk_buffer_type_get_max_size(ggml_backend_buffer_type_t buft);
 size_t ggml_backend_vk_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft, const ggml_tensor * tensor);
 void ggml_backend_vk_free(ggml_backend_t backend);
+ggml_backend_buffer_t ggml_backend_vk_device_buffer_from_host_ptr(ggml_backend_dev_t dev, void * ptr, size_t size, size_t max_tensor_size);
 ggml_backend_reg_t ggml_backend_vk_reg();
 
+// comm
+#if !defined(_WIN32)
+void * ggml_backend_vk_comm_init(ggml_backend_t * backends, size_t n_backends);
+void ggml_backend_vk_comm_free(void * comm_ctx);
+bool ggml_backend_vk_comm_allreduce_tensor(void * comm_ctx, ggml_tensor ** tensors);
+#endif
+
 // debug
 int64_t ggml_vk_get_op_batch_size(const ggml_tensor * op);
 #ifdef GGML_VULKAN_RUN_TESTS

ggml/src/ggml-vulkan/ggml-vulkan-types.h

@@ -679,6 +679,7 @@ struct vk_device_struct {
     uint64_t suballocation_block_size;
     uint64_t min_imported_host_pointer_alignment;
     bool external_memory_host {};
+    bool external_semaphore {};
     bool fp16;
     bool bf16;
     bool pipeline_robustness;
@@ -1087,6 +1088,9 @@ struct vk_submission {
     vk_command_buffer* buffer = nullptr;
     std::vector<vk_semaphore> wait_semaphores;
     std::vector<vk_semaphore> signal_semaphores;
+    // submit waits until *ready >= ready_value
+    const std::atomic<uint64_t> * ready = nullptr;
+    uint64_t ready_value = 0;
 };
 
 typedef std::vector<vk_submission> vk_sequence;
@@ -1225,7 +1229,7 @@ struct ggml_backend_vk_context {
 
     vk_device device;
 
-    size_t semaphore_idx, event_idx;
+    size_t semaphore_idx, binary_semaphore_idx, event_idx;
     ggml_vk_garbage_collector gc;
     size_t prealloc_size_x, prealloc_size_y, prealloc_size_split_k, prealloc_size_add_rms_partials, prealloc_size_add_rms_partials_offset;
     vk_buffer prealloc_x, prealloc_y, prealloc_split_k, prealloc_add_rms_partials, sync_staging;
@@ -1341,6 +1345,7 @@ struct vk_instance_t {
     uint32_t host_import_failed = 0;  // bitmask by device->idx, guarded by host_allocations_mutex
 
     bool multi_device = false;
+    bool allreduce_disabled = false;
 };
 
 typedef void (*ggml_vk_func_t)(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst);

ggml/src/ggml-vulkan/ggml-vulkan.cpp

@@ -934,6 +934,11 @@ void ggml_vk_submit(vk_context& ctx, vk::Fence fence) {
 
     for (const auto& sequence : ctx->seqs) {
         for (const auto& submission : sequence) {
+            if (submission.ready) {
+                while (submission.ready->load(std::memory_order_acquire) < submission.ready_value) {
+                    YIELD();
+                }
+            }
             stage_flags.push_back({});
             idx++;
             tl_wait_vals.push_back({});
@@ -1082,14 +1087,16 @@ vk_context ggml_vk_create_temporary_context(vk_command_pool& p) {
     return result;
 }
 
-static vk_semaphore * ggml_vk_create_binary_semaphore(ggml_backend_vk_context * ctx) {
-    VK_LOG_DEBUG("ggml_vk_create_timeline_semaphore()");
-    vk::SemaphoreTypeCreateInfo tci{ vk::SemaphoreType::eBinary, 0 };
-    vk::SemaphoreCreateInfo ci{};
-    ci.setPNext(&tci);
-    vk::Semaphore semaphore = ctx->device->device.createSemaphore(ci);
-    ctx->gc.semaphores.push_back({ semaphore, 0 });
-    return &ctx->gc.semaphores[ctx->gc.semaphores.size() - 1];
+vk::Semaphore ggml_vk_create_binary_semaphore(ggml_backend_vk_context * ctx) {
+    VK_LOG_DEBUG("ggml_vk_create_binary_semaphore()");
+    if (ctx->binary_semaphore_idx >= ctx->gc.semaphores.size()) {
+        vk::SemaphoreTypeCreateInfo tci{ vk::SemaphoreType::eBinary, 0 };
+        vk::SemaphoreCreateInfo ci{};
+        ci.setPNext(&tci);
+        vk::Semaphore semaphore = ctx->device->device.createSemaphore(ci);
+        ctx->gc.semaphores.push_back({ semaphore, 0 });
+    }
+    return ctx->gc.semaphores[ctx->binary_semaphore_idx++].s;
 }
 
 static vk_semaphore * ggml_vk_create_timeline_semaphore(ggml_backend_vk_context * ctx) {
@@ -4128,6 +4135,8 @@ vk_device ggml_vk_get_device(size_t idx) {
                 device->memory_priority = true;
             } else if (strcmp("VK_EXT_external_memory_host", properties.extensionName) == 0) {
                 device->external_memory_host = true;
+            } else if (strcmp("VK_KHR_external_semaphore_fd", properties.extensionName) == 0) {
+                device->external_semaphore = true;
 #if defined(VK_EXT_shader_64bit_indexing)
             } else if (strcmp("VK_EXT_shader_64bit_indexing", properties.extensionName) == 0) {
                 device->shader_64b_indexing = true;
@@ -4475,6 +4484,9 @@ vk_device ggml_vk_get_device(size_t idx) {
         if (device->external_memory_host) {
             device_extensions.push_back("VK_EXT_external_memory_host");
         }
+        if (device->external_semaphore) {
+            device_extensions.push_back("VK_KHR_external_semaphore_fd");
+        }
 
 #if defined(VK_EXT_shader_64bit_indexing)
         VkPhysicalDeviceShader64BitIndexingFeaturesEXT shader_64bit_indexing_features {};
@@ -5287,6 +5299,7 @@ void ggml_vk_instance_init() {
     vk_perf_logger_concurrent = getenv("GGML_VK_PERF_LOGGER_CONCURRENT") != nullptr;
     vk_enable_sync_logger = getenv("GGML_VK_SYNC_LOGGER") != nullptr;
     vk_memory_logger_enabled = getenv("GGML_VK_MEMORY_LOGGER") != nullptr;
+    vk_instance.allreduce_disabled = getenv("GGML_VK_DISABLE_ALLREDUCE") != nullptr;
     const char* GGML_VK_PIPELINE_STATS = getenv("GGML_VK_PIPELINE_STATS");
     if (GGML_VK_PIPELINE_STATS != nullptr) {
         vk_pipeline_stats_filter = GGML_VK_PIPELINE_STATS;
@@ -5493,6 +5506,7 @@ void ggml_vk_init(ggml_backend_vk_context * ctx, size_t idx) {
     ctx->device = ggml_vk_get_device(idx);
 
     ctx->semaphore_idx = 0;
+    ctx->binary_semaphore_idx = 0;
     ctx->event_idx = 0;
 
     ctx->prealloc_size_x = 0;
@@ -5920,8 +5934,18 @@ static uint32_t ggml_vk_guess_split_k(ggml_backend_vk_context * ctx, uint32_t m,
         return 1;
     }
 
+    const bool rdna = ctx->device->architecture == AMD_RDNA3 || ctx->device->architecture == AMD_RDNA4;
     uint32_t split_k = 1;
-    if (ctx->device->shader_core_count != 0 && n >= pipeline->wg_denoms[1]) {
+    if (rdna && ctx->device->shader_core_count != 0 && k >= 2048 && n <= 32) {
+        // Small n gives too few workgroups to hide latency, aim for 8 per CU. Small tiles are one subgroup each.
+        const uint32_t tiles = CEIL_DIV(m, pipeline->wg_denoms[0]) * CEIL_DIV(n, pipeline->wg_denoms[1]);
+        if (tiles < 4 * ctx->device->shader_core_count) {
+            split_k = std::min(8u, CEIL_DIV(8 * ctx->device->shader_core_count, tiles));
+        }
+    }
+
+    // On RDNA, also split when n <= 64 only partly fills one tile column.
+    if (split_k == 1 && ctx->device->shader_core_count != 0 && (n >= pipeline->wg_denoms[1] || (rdna && n <= 64))) {
         // If k is 'large' and the SMs will fill less than halfway, use split_k.
         uint32_t m_tiles = CEIL_DIV(m, pipeline->wg_denoms[0]);
         uint32_t n_tiles = CEIL_DIV(n, pipeline->wg_denoms[1]);
@@ -5934,24 +5958,24 @@ static uint32_t ggml_vk_guess_split_k(ggml_backend_vk_context * ctx, uint32_t m,
             }
             // Cap the split at 8x. Unless k is huge this is a lot of overhead.
             split_k = std::min(split_k, 8u);
-
-            // ggml_vk_matmul will align the splits to be a multiple of 256.
-            // If this rounded up size would cause the last split to be empty,
-            // then reduce the split count.
-            while (true) {
-                if (split_k == 1) {
-                    break;
-                }
-                uint32_t k_split = CEIL_DIV(k, split_k);
-                k_split = ROUNDUP_POW2(k_split, 256);
-                if (k_split * (split_k - 1) < k) {
-                    break;
-                }
-                split_k--;
-            }
         }
     }
 
+    // ggml_vk_matmul will align the splits to be a multiple of 256.
+    // If this rounded up size would cause the last split to be empty,
+    // then reduce the split count.
+    while (true) {
+        if (split_k == 1) {
+            break;
+        }
+        uint32_t k_split = CEIL_DIV(k, split_k);
+        k_split = ROUNDUP_POW2(k_split, 256);
+        if (k_split * (split_k - 1) < k) {
+            break;
+        }
+        split_k--;
+    }
+
     return split_k;
 }
 
@@ -13003,10 +13027,7 @@ void ggml_vk_graph_cleanup(ggml_backend_vk_context * ctx) {
         ggml_vk_command_pool_cleanup(ctx->device, ctx->transfer_cmd_pool);
     }
 
-    for (size_t i = 0; i < ctx->gc.semaphores.size(); i++) {
-        ctx->device->device.destroySemaphore({ ctx->gc.semaphores[i].s });
-    }
-    ctx->gc.semaphores.clear();
+    ctx->binary_semaphore_idx = 0;
 
     for (size_t i = 0; i < ctx->gc.tl_semaphores.size(); i++) {
         ctx->device->device.destroySemaphore({ ctx->gc.tl_semaphores[i].s });
@@ -13054,6 +13075,11 @@ void ggml_vk_cleanup(ggml_backend_vk_context * ctx) {
     }
     ctx->gc.events.clear();
 
+    for (auto & sem : ctx->gc.semaphores) {
+        ctx->device->device.destroySemaphore(sem.s);
+    }
+    ctx->gc.semaphores.clear();
+
     ctx->device->device.destroyFence(ctx->fence);
     ctx->device->device.destroyFence(ctx->almost_ready_fence);
 
@@ -14454,7 +14480,9 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
             flops_cap = 2'000'000'000ULL * ctx->device->shader_core_count;
         }
     }
-    uint64_t flops_per_submit = std::min(flops_cap, ctx->last_total_flops / 40u);
+    // small graphs gain nothing from splitting, each submit costs CPU time
+    uint64_t flops_per_submit = ctx->last_total_flops < 1'000'000'000ULL ? 0 :
+        std::clamp(ctx->last_total_flops / 40u, std::min<uint64_t>(100'000'000ULL, flops_cap), flops_cap);
 
     auto const submit_after = [&](int start, int end) {
         if (ctx->device->serialize_submissions) {
@@ -16256,7 +16284,7 @@ static void ggml_backend_vk_device_event_synchronize(ggml_backend_dev_t dev, ggm
     }
 }
 
-static ggml_backend_buffer_t ggml_backend_vk_device_buffer_from_host_ptr(ggml_backend_dev_t dev, void * ptr, size_t size, size_t max_tensor_size) {
+ggml_backend_buffer_t ggml_backend_vk_device_buffer_from_host_ptr(ggml_backend_dev_t dev, void * ptr, size_t size, size_t max_tensor_size) {
     VK_LOG_DEBUG("ggml_backend_vk_device_buffer_from_host_ptr(backend=" << dev << ", ptr=" << ptr << ", size=" << size << ")");
     GGML_UNUSED(max_tensor_size);
 
@@ -16338,11 +16366,29 @@ static ggml_backend_dev_t ggml_backend_vk_reg_get_device(ggml_backend_reg_t reg,
     return devices[device];
 }
 
+static void * ggml_backend_vk_reg_get_proc_address(ggml_backend_reg_t reg, const char * name) {
+    UNUSED(reg);
+#if !defined(_WIN32)
+    if (strcmp(name, "ggml_backend_comm_init") == 0) {
+        return (void *) ggml_backend_vk_comm_init;
+    }
+    if (strcmp(name, "ggml_backend_comm_free") == 0) {
+        return (void *) ggml_backend_vk_comm_free;
+    }
+    if (strcmp(name, "ggml_backend_comm_allreduce_tensor") == 0) {
+        return (void *) ggml_backend_vk_comm_allreduce_tensor;
+    }
+#else
+    UNUSED(name);
+#endif
+    return nullptr;
+}
+
 static const struct ggml_backend_reg_i ggml_backend_vk_reg_i = {
     /* .get_name         = */ ggml_backend_vk_reg_get_name,
     /* .get_device_count = */ ggml_backend_vk_reg_get_device_count,
     /* .get_device       = */ ggml_backend_vk_reg_get_device,
-    /* .get_proc_address = */ NULL,
+    /* .get_proc_address = */ ggml_backend_vk_reg_get_proc_address,
 };
 
 ggml_backend_reg_t ggml_backend_vk_reg() {

ggml/src/ggml-vulkan/vulkan-shaders/count_experts.comp

@@ -40,6 +40,8 @@ layout (binding = 2) writeonly buffer DST {float data_dst[];};
 shared uint vals[EXPERTS_PER_CHUNK];
 shared uint offsets[EXPERTS_PER_CHUNK];
 shared uint cursors[EXPERTS_PER_CHUNK];
+shared uint skipped_rows;
+shared uint skip_count;
 
 // data_d[0] is the tile count. Each tile has (expert, valid rows).
 // Packed (token << 16) | slot row ids start at p.row_ids_offset, with each expert on a tile boundary.
@@ -51,23 +53,38 @@ void write_tiles(uint expert, uint count, uint tile_begin, uint n_tiles) {
     }
 }
 
-// the matmul writes only the rows that an expert owns, so zero the rows of the skipped slots
+// The matmul does not write skipped rows, so collect their offsets and zero them cooperatively.
 void zero_skipped_rows() {
     const uint num_elements = p.ne00 * p.ne01;
     const uint tid = gl_LocalInvocationID.x;
 
-    for (uint idx = 0; idx < num_elements; idx++) {
-        const uint i01 = fastdiv(idx, p.ne00mp, p.ne00L);
-        const uint i00 = idx - i01 * p.ne00;
-
-        if (int(data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00]) != -1) {
-            continue;
+    // Reuse offsets[] as a queue, so each batch must fit in EXPERTS_PER_CHUNK.
+    for (uint base = 0; base < num_elements; base += EXPERTS_PER_CHUNK) {
+        const uint end = min(base + EXPERTS_PER_CHUNK, num_elements);
+        if (tid == 0) {
+            skip_count = 0;
         }
+        barrier();
 
-        const uint d_off = i00 * p.stride_dst + i01 * p.batch_stride_dst;
-        for (uint i = tid; i < p.ne0_dst; i += BLOCK_SIZE) {
-            data_dst[d_off + i] = 0.0f;
+        for (uint idx = base + tid; idx < end; idx += BLOCK_SIZE) {
+            const uint i01 = fastdiv(idx, p.ne00mp, p.ne00L);
+            const uint i00 = idx - i01 * p.ne00;
+
+            if (int(data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00]) == -1) {
+                const uint slot = atomicAdd(skip_count, 1);
+                offsets[slot] = i00 * p.stride_dst + i01 * p.batch_stride_dst;
+            }
         }
+        barrier();
+
+        const uint n_skipped = skip_count;
+        for (uint r = 0; r < n_skipped; ++r) {
+            const uint d_off = offsets[r];
+            for (uint i = tid; i < p.ne0_dst; i += BLOCK_SIZE) {
+                data_dst[d_off + i] = 0.0f;
+            }
+        }
+        barrier();
     }
 }
 
@@ -76,6 +93,11 @@ void main() {
     const uint tid = gl_LocalInvocationID.x;
     const uint tile_shift = findLSB(p.row_tile_size);
     uint total_tiles = 0;
+    uint local_skipped = 0;
+    if (tid == 0) {
+        skipped_rows = 0;
+    }
+    barrier();
 #ifdef USE_SUBGROUPS
     // Use the subgroup that contains invocation 0.
     const bool prefix_subgroup = subgroupAny(tid == 0);
@@ -135,7 +157,12 @@ void main() {
         for (uint idx = tid; idx < num_elements; idx += BLOCK_SIZE) {
             const uint i01 = fastdiv(idx, p.ne00mp, p.ne00L);
             const uint i00 = idx - i01 * p.ne00;
-            const uint expert = data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00] - expert_base;
+            const uint id = data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00];
+            // Count skipped slots only in the first expert chunk.
+            if (expert_base == 0 && int(id) == -1) {
+                local_skipped++;
+            }
+            const uint expert = id - expert_base;
             if (expert < n_experts) {
                 const uint row = atomicAdd(cursors[expert], 1);
                 data_d[p.row_ids_offset + offsets[expert] + row] = (i01 << 16) | i00;
@@ -144,9 +171,17 @@ void main() {
         // The next chunk clears vals. Scatter uses only offsets and cursors.
         expert_base += n_experts;
     }
+    if (local_skipped != 0) {
+        atomicAdd(skipped_rows, local_skipped);
+    }
+    // All scatter reads must finish before reusing offsets[] as a queue.
+    barrier();
+
     if (tid == 0) {
         data_d[0] = total_tiles;
     }
 
-    zero_skipped_rows();
+    if (skipped_rows != 0) {
+        zero_skipped_rows();
+    }
 }

ggml/src/ggml-webgpu/ggml-webgpu-shader-lib.hpp

@@ -19,6 +19,7 @@
 #define GGML_WEBGPU_I32_SIZE_BYTES                   4
 #define GGML_WEBGPU_FLASH_ATTN_PREFERRED_KV_SG_TILES 8u
 #define GGML_WEBGPU_FLASH_ATTN_VEC_MAX_SEQ_LEN       20u
+#define GGML_WEBGPU_MUL_MAT_VEC_MAX_ROWS             4
 #define GGML_WEBGPU_FLASH_ATTN_VEC_MAX_KV_TILE       32u
 #define GGML_WEBGPU_FLASH_ATTN_TILE_MAX_KV_TILE      64u
 #define GGML_WEBGPU_FLASH_ATTN_PREFERRED_WG_SIZE     128u
@@ -1189,9 +1190,10 @@ struct ggml_webgpu_soft_max_pipeline_key_hash {
 
 inline bool ggml_webgpu_can_use_mmvq(const ggml_tensor * src0,
                                      const ggml_tensor * src1,
+                                     int64_t             n_rows,
                                      bool                supports_dot_product,
                                      const std::string & vendor) {
-    if (src1->ne[1] <= 4) {
+    if (n_rows <= GGML_WEBGPU_MUL_MAT_VEC_MAX_ROWS) {
         bool supports_dp4a = vendor == "amd" || vendor == "intel" || vendor == "nvidia";
         if (supports_dp4a && supports_dot_product) {
             switch (src1->type) {
@@ -2011,13 +2013,13 @@ class ggml_webgpu_shader_lib {
         ggml_webgpu_mul_mat_vec_pipeline_key key = {};
         key.src0_type                            = context.src0->type;
         key.src1_type                            = context.src1->type;
-        key.vectorized = (context.src0->ne[0] % 4 == 0 &&
-                          (context.src0->type == GGML_TYPE_F32 || context.src0->type == GGML_TYPE_F16)) ?
-                             1 :
-                             0;
-        key.num_cols   = context.dst->ne[1];
-        key.use_mmvq =
-            ggml_webgpu_can_use_mmvq(context.src0, context.src1, context.supports_dot_product, context.vendor);
+        key.vectorized  = (context.src0->ne[0] % 4 == 0 &&
+                           (context.src0->type == GGML_TYPE_F32 || context.src0->type == GGML_TYPE_F16)) ?
+                              1 :
+                              0;
+        key.num_cols    = context.dst->ne[1];
+        key.use_mmvq    = ggml_webgpu_can_use_mmvq(context.src0, context.src1, context.src1->ne[1],
+                                                   context.supports_dot_product, context.vendor);
         key.src_overlap = src_overlap;
 
         auto it = mul_mat_vec_pipelines.find(key);

ggml/src/ggml-webgpu/ggml-webgpu.cpp

@@ -1634,11 +1634,11 @@ static webgpu_encoded_op ggml_webgpu_mul_mat(webgpu_context & ctx,
                                              ggml_tensor *    src1,
                                              ggml_tensor *    dst) {
     // Determine if this is a mat-vec operation
-    bool use_mat_vec = (dst->ne[1] <= 4);
+    bool use_mat_vec = (dst->ne[1] <= GGML_WEBGPU_MUL_MAT_VEC_MAX_ROWS);
 
     // use MMVQ path for mat-vec
-    bool use_mmvq = ggml_webgpu_can_use_mmvq(src0, src1, ctx->global_ctx->capabilities.supports_dot_product,
-                                             ctx->global_ctx->vendor);
+    bool use_mmvq = ggml_webgpu_can_use_mmvq(
+        src0, src1, src1->ne[1], ctx->global_ctx->capabilities.supports_dot_product, ctx->global_ctx->vendor);
 
     ggml_webgpu_shader_lib_context shader_lib_ctx = {};
 
@@ -1965,7 +1965,8 @@ struct ggml_webgpu_flash_attn_op {
 static bool ggml_webgpu_flash_attn_use_vec_path(const webgpu_global_context & global_ctx,
                                                 const ggml_tensor *           Q,
                                                 const ggml_tensor *           K,
-                                                const ggml_tensor *           V) {
+                                                const ggml_tensor *           V,
+                                                int64_t                       n_q) {
     const size_t storage_offset_alignment = global_ctx->capabilities.limits.minStorageBufferOffsetAlignment;
 
     const bool k_float_vec4_aligned = (K->type != GGML_TYPE_F16 && K->type != GGML_TYPE_F32) ||
@@ -1979,7 +1980,7 @@ static bool ggml_webgpu_flash_attn_use_vec_path(const webgpu_global_context & gl
         ggml_is_quantized(V->type) ? ggml_blck_size(V->type) : GGML_WEBGPU_FLASH_ATTN_TILE_KV_VEC_WIDTH;
     const bool kv_vec_head_dims_aligned = Q->ne[0] % k_vec_head_align == 0 && V->ne[0] % v_vec_head_align == 0;
 
-    return global_ctx->capabilities.supports_subgroups && (Q->ne[1] < GGML_WEBGPU_FLASH_ATTN_VEC_MAX_SEQ_LEN) &&
+    return global_ctx->capabilities.supports_subgroups && (n_q < GGML_WEBGPU_FLASH_ATTN_VEC_MAX_SEQ_LEN) &&
            kv_vec_head_dims_aligned && k_float_vec4_aligned && v_float_vec4_aligned;
 }
 
@@ -2288,7 +2289,7 @@ static webgpu_encoded_op ggml_webgpu_flash_attn(webgpu_context & ctx,
                                                 ggml_tensor *    sinks,
                                                 ggml_tensor *    dst) {
     ggml_webgpu_flash_attn_op op = ggml_webgpu_flash_attn_prepare(ctx, Q, K, V, mask, sinks, dst);
-    if (ggml_webgpu_flash_attn_use_vec_path(ctx->global_ctx, Q, K, V)) {
+    if (ggml_webgpu_flash_attn_use_vec_path(ctx->global_ctx, Q, K, V, Q->ne[1])) {
         return ggml_webgpu_flash_attn_vec(ctx, Q, K, V, mask, sinks, dst, std::move(op));
     }
     return ggml_webgpu_flash_attn_direct(ctx, op);
@@ -4012,7 +4013,10 @@ static size_t ggml_backend_webgpu_buffer_type_get_alloc_size(ggml_backend_buffer
                 const ggml_tensor * V            = tensor->src[2];
                 const ggml_tensor * mask         = tensor->src[3];
                 const auto &        capabilities = ctx->webgpu_global_ctx->capabilities;
-                if (ggml_webgpu_flash_attn_use_vec_path(ctx->webgpu_global_ctx, Q, K, V)) {
+                // the vec path scratch grows with the query rows, so a longer batch reserves the scratch of the
+                // longest vec batch and the size stays monotonic in the batch length
+                const int64_t       n_q = std::min<int64_t>(Q->ne[1], GGML_WEBGPU_FLASH_ATTN_VEC_MAX_SEQ_LEN - 1);
+                if (ggml_webgpu_flash_attn_use_vec_path(ctx->webgpu_global_ctx, Q, K, V, n_q)) {
                     const bool kv_direct =
                         ggml_webgpu_flash_attn_k_direct(Q, K, GGML_WEBGPU_FLASH_ATTN_TILE_KV_VEC_WIDTH) ||
                         ggml_webgpu_flash_attn_v_direct(Q, V, GGML_WEBGPU_FLASH_ATTN_TILE_KV_VEC_WIDTH);
@@ -4024,7 +4028,7 @@ static size_t ggml_backend_webgpu_buffer_type_get_alloc_size(ggml_backend_buffer
                     uint32_t       nwg = ggml_webgpu_flash_attn_vec_nwg(vec_nwg_cap, kv_tile, (uint32_t) K->ne[1]);
 
                     const size_t   align = capabilities.limits.minStorageBufferOffsetAlignment;
-                    const uint64_t nrows = (uint64_t) Q->ne[1] * Q->ne[2] * Q->ne[3];
+                    const uint64_t nrows = (uint64_t) n_q * Q->ne[2] * Q->ne[3];
                     if (nwg > 1u) {
                         const uint64_t tmp_data_elems  = nrows * (uint64_t) V->ne[0] * nwg;
                         const uint64_t tmp_stats_elems = nrows * 2u * nwg;
@@ -4036,7 +4040,7 @@ static size_t ggml_backend_webgpu_buffer_type_get_alloc_size(ggml_backend_buffer
                     }
                     if (mask != nullptr) {
                         const uint32_t blk_nblk0       = CEIL_DIV((uint32_t) K->ne[1], kv_tile);
-                        const uint32_t blk_nblk1       = CEIL_DIV((uint32_t) Q->ne[1], 1u);
+                        const uint32_t blk_nblk1       = (uint32_t) n_q;
                         const uint32_t stride_mask3    = (uint32_t) (mask->nb[3] / ggml_type_size(mask->type));
                         const uint32_t blk_batch_count = stride_mask3 > 0 ? (uint32_t) Q->ne[3] : 1u;
                         const uint64_t blk_elems       = (uint64_t) blk_nblk0 * blk_nblk1 * blk_batch_count;
@@ -4052,11 +4056,15 @@ static size_t ggml_backend_webgpu_buffer_type_get_alloc_size(ggml_backend_buffer
             {
                 const ggml_tensor * src0 = tensor->src[0];
                 const ggml_tensor * src1 = tensor->src[1];
-                bool                use_mmvq =
-                    ggml_webgpu_can_use_mmvq(src0, src1, ctx->webgpu_global_ctx->capabilities.supports_dot_product,
-                                             ctx->webgpu_global_ctx->vendor);
+
+                // the q8 copy of src1 grows with its rows, so a longer batch reserves the copy of the longest
+                // mmvq batch and the size stays monotonic in the batch length
+                const int64_t n_rows   = std::min<int64_t>(src1->ne[1], GGML_WEBGPU_MUL_MAT_VEC_MAX_ROWS);
+                const bool    use_mmvq = ggml_webgpu_can_use_mmvq(
+                    src0, src1, n_rows, ctx->webgpu_global_ctx->capabilities.supports_dot_product,
+                    ctx->webgpu_global_ctx->vendor);
                 if (use_mmvq) {
-                    const size_t q8_src1_size = src1->ne[3] * src1->ne[2] * src1->ne[1] *
+                    const size_t q8_src1_size = src1->ne[3] * src1->ne[2] * n_rows *
                                                 (36 /* sizeof(q8_1) */ * (src1->ne[0] / /* block_size */ 32));
                     res = ROUNDUP_POW2(res + q8_src1_size +
                                            ctx->webgpu_global_ctx->capabilities.limits.minStorageBufferOffsetAlignment,

ggml/src/ggml-zdnn/common.hpp

@@ -34,6 +34,8 @@ struct ggml_backend_zdnn_device_context {
 
     size_t max_size;
 
+    int32_t n_threads;  // CPU threads
+
     char name[128];
 };
 

ggml/src/ggml-zdnn/ggml-zdnn.cpp

@@ -12,6 +12,7 @@
 #include <vector>
 #include <memory>
 #include <csignal>  // raise(SIGTRAP)
+#include <thread>
 #include <unistd.h>
 
 static void ggml_zdnn_compute_forward_mul_mat(
@@ -85,6 +86,15 @@ static enum ggml_status ggml_zdnn_graph_compute(ggml_backend_t backend, ggml_cgr
     GGML_UNUSED(ctx_dev);
 }
 
+static int64_t ggml_zdnn_mul_mat_min_batch(ggml_type type, int32_t n_threads) {
+    switch (type) {
+        case GGML_TYPE_Q8_0:
+            return std::max<int64_t>(1, n_threads - 2);
+        default:
+            return 1;
+    }
+}
+
 static bool ggml_zdnn_supports_op(const ggml_backend_zdnn_device_context * ctx_dev, const ggml_tensor * op) {
     switch (op->op) {
         case GGML_OP_NONE:
@@ -112,6 +122,11 @@ static bool ggml_zdnn_supports_op(const ggml_backend_zdnn_device_context * ctx_d
                         return false;
                 }
 
+                // offload smaller batches to the CPU
+                if (ne1 < ggml_zdnn_mul_mat_min_batch(weights->type, ctx_dev->n_threads)) {
+                    return false;
+                }
+
                 switch (weights->type) {
                     case GGML_TYPE_F32:
                     case GGML_TYPE_F16:
@@ -150,6 +165,7 @@ static ggml_backend_zdnn_device_context g_ggml_ctx_dev_main = {
     /* .has_parmblkformat_0   = */ false,
     /* .has_parmblkformat_1   = */ false,
     /* .max_size              = */ 0,
+    /* .n_threads             = */ 1,
     /* .name                  = */ "",
 };
 
@@ -164,6 +180,8 @@ static int ggml_backend_zdnn_device_acq(ggml_backend_zdnn_device_context * ctx)
         ctx->has_parmblkformat_0 = zdnn_is_nnpa_parmblk_fmt_installed(1, NNPA_PARMBLKFORMAT_0);
         ctx->has_parmblkformat_1 = zdnn_is_nnpa_parmblk_fmt_installed(1, NNPA_PARMBLKFORMAT_1);
         ctx->max_size = zdnn_get_nnpa_max_dim_idx_size();
+        // llama.cpp uses every core by default, and splits the first graph before it sets the threads
+        ctx->n_threads = std::max<int32_t>(1, (int32_t)std::thread::hardware_concurrency());
         strncpy(ctx->name, GGML_ZDNN_NAME, sizeof(ctx->name) - 1);
     }
 
@@ -452,8 +470,7 @@ static size_t ggml_backend_zdnn_buffer_type_get_alignment(ggml_backend_buffer_ty
 }
 
 static bool ggml_backend_zdnn_buffer_type_is_host(ggml_backend_buffer_type_t buft) {
-    /* while it resides in host memory, additional transformation is needed */
-    return false;
+    return true;
 
     GGML_UNUSED(buft);
 }
@@ -673,6 +690,9 @@ static void ggml_backend_zdnn_set_n_threads(ggml_backend_t backend, int n_thread
 
     ggml_backend_zdnn_context * ctx = (ggml_backend_zdnn_context *)backend->context;
     ctx->n_threads = n_threads;
+
+    ggml_backend_zdnn_device_context * ctx_dev = (ggml_backend_zdnn_device_context *)backend->device->context;
+    ctx_dev->n_threads = n_threads;
 }
 
 static void * ggml_backend_zdnn_get_proc_address(ggml_backend_reg_t reg, const char * name) {

ggml/src/ggml-zdnn/mmq.cpp

@@ -1,6 +1,7 @@
 #include "ggml.h"
 #include "mmq.hpp"
 #include "quantize.hpp"
+#include "utils.hpp"
 
 #include <algorithm>
 #include <future>
@@ -189,6 +190,11 @@ void ggml_zdnn_mul_mat_q(
     ggml_backend_zdnn_buffer * weights_extra = (ggml_backend_zdnn_buffer *)weights->extra;
     ggml_backend_zdnn_buffer * output_extra  = (ggml_backend_zdnn_buffer *)output->extra;
 
+    // weights read straight into a host buffer skip set_tensor, so transform them on first use
+    if (!weights_extra->ztensor.is_transformed) {
+        ggml_zdnn_load_tensor_q(weights_extra, weights);
+    }
+
     // starting threads costs more than it saves for small inputs, e.g. during token generation
     const int32_t n_threads = ne10*ne11 >= 65536 ? ctx->n_threads : 1;
 

ggml/src/ggml.c

@@ -71,12 +71,12 @@ uint64_t ggml_graph_next_uid(void) {
 
 // Needed for ggml_fp32_to_bf16_row()
 #if defined(__AVX512BF16__)
-#if defined(_MSC_VER)
+#if defined(_MSC_VER) && !defined(__clang__)
 #define m512i(p) p
 #else
 #include <immintrin.h>
 #define m512i(p) (__m512i)(p)
-#endif // defined(_MSC_VER)
+#endif // defined(_MSC_VER) && !defined(__clang__)
 #endif // defined(__AVX512BF16__)
 
 #if defined(__linux__) || \
@@ -357,7 +357,8 @@ void * ggml_aligned_malloc(size_t size) {
     void * aligned_memory = NULL;
   #ifdef GGML_USE_CPU_HBM
     int result = hbw_posix_memalign(&aligned_memory, alignment, size);
-  #elif TARGET_OS_OSX && !defined(GGML_TSAN_ENABLED)
+  #elif defined(__APPLE__) && !defined(GGML_TSAN_ENABLED)
+    // libmalloc can keep freed pages in the process footprint, vm_deallocate returns them at once
     GGML_UNUSED(alignment);
     kern_return_t alloc_status = vm_allocate((vm_map_t) mach_task_self(), (vm_address_t *) &aligned_memory, size, VM_FLAGS_ANYWHERE);
     int result = EFAULT;
@@ -404,7 +405,7 @@ void ggml_aligned_free(void * ptr, size_t size) {
     if (ptr != NULL) {
         hbw_free(ptr);
     }
-#elif TARGET_OS_OSX && !defined(GGML_TSAN_ENABLED)
+#elif defined(__APPLE__) && !defined(GGML_TSAN_ENABLED)
     if (ptr != NULL) {
         vm_deallocate((vm_map_t)mach_task_self(), (vm_address_t)ptr, size);
     }

requirements/requirements-gguf_editor_gui.txt

@@ -1,3 +1,4 @@
-numpy~=2.2.6
+numpy~=2.2.6; python_version < "3.14"
+numpy~=2.4.6; python_version >= "3.14"
 PySide6~=6.9.0
 gguf>=0.17.0

requirements/requirements-server-bench.txt

@@ -1,5 +1,6 @@
 datasets~=4.8.0
 matplotlib~=3.10.0
-numpy~=2.2.6
+numpy~=2.2.6; python_version < "3.14"
+numpy~=2.4.6; python_version >= "3.14"
 requests~=2.32.3
 tqdm~=4.67.1

requirements/requirements-tool_bench.txt

@@ -1,9 +1,9 @@
-aiohttp~=3.9.3
 pytest~=8.3.3
 matplotlib~=3.10.0
-numpy~=2.2.6
+numpy~=2.2.6; python_version < "3.14"
+numpy~=2.4.6; python_version >= "3.14"
 openai~=2.14.0
-pandas~=2.2.3
+pandas~=2.3.3
 prometheus-client~=0.20.0
 requests~=2.32.3
 wget~=3.2

src/llama-context.cpp

@@ -1379,14 +1379,9 @@ void llama_context::set_nextn_layer_offset(int32_t offset) {
 void llama_context::set_causal_attn(bool value) {
     LLAMA_LOG_DEBUG("%s: value = %d\n", __func__, value);
 
-    if (cparams.causal_attn == value) {
-        return;
-    }
+    // no scheduler reserve needed because graph shapes must not depend on causal_attn, a flip only rebuilds the graph
 
     cparams.causal_attn = value;
-
-    // no scheduler reserve needed because graph shapes must not depend on causal_attn, a flip only rebuilds the graph
-    //sched_need_reserve = true;
 }
 
 bool llama_context::get_causal_attn() const {
@@ -1396,14 +1391,9 @@ bool llama_context::get_causal_attn() const {
 void llama_context::set_warmup(bool value) {
     LLAMA_LOG_DEBUG("%s: value = %d\n", __func__, value);
 
-    if (cparams.warmup == value) {
-        return;
-    }
+    // no scheduler reserve needed because warmups are usually with small batches
 
     cparams.warmup = value;
-
-    // warmups are usually with small batches, so no need to reserve
-    //sched_need_reserve = true;
 }
 
 bool llama_context::set_sampler(llama_seq_id seq_id, llama_sampler * sampler) {

tests/CMakeLists.txt

@@ -1,3 +1,5 @@
+find_package(Threads REQUIRED)
+
 llama_add_compile_flags()
 
 # Optional args:
@@ -351,7 +353,7 @@ if (GGML_METAL AND NOT GGML_BACKEND_DL)
     target_include_directories(test-metal-graph-optimize PRIVATE ${PROJECT_SOURCE_DIR}/ggml/src ${PROJECT_SOURCE_DIR}/ggml/src/ggml-metal)
 endif()
 
-llama_build(test-backend-ops.cpp LINK ggml)
+llama_build(test-backend-ops.cpp LINK ggml Threads::Threads)
 
 llama_build_and_test(test-backend-tensor.cpp LINK ggml)
 
@@ -362,6 +364,7 @@ llama_build_and_test(test-backend-sampler.cpp   LABEL "model behavior backend")
 
 if (APPLE)
     llama_build(test-rset-release.cpp)
+    llama_build_and_test(test-host-buffer-release.cpp LINK ggml)
 endif()
 
 if (NOT GGML_BACKEND_DL AND GGML_CPU)

tests/test-host-buffer-release.cpp

@@ -0,0 +1,47 @@
+// ref: https://github.com/ggml-org/llama.cpp/issues/30556
+
+#include "ggml-backend.h"
+
+#include <cstdint>
+#include <cstdio>
+#include <cstring>
+#include <mach/mach.h>
+
+static uint64_t phys_footprint() {
+    task_vm_info_data_t info;
+    mach_msg_type_number_t count = TASK_VM_INFO_COUNT;
+    const kern_return_t ret = task_info(mach_task_self(), TASK_VM_INFO, (task_info_t) &info, &count);
+    GGML_ASSERT(ret == KERN_SUCCESS);
+    return info.phys_footprint;
+}
+
+static bool test_release(ggml_backend_buffer_type_t buft) {
+    const size_t size = 256ull*1024*1024;
+
+    ggml_backend_buffer_t buf = ggml_backend_buft_alloc_buffer(buft, size);
+    GGML_ASSERT(buf != nullptr);
+    ggml_backend_buffer_clear(buf, 1);
+
+    const uint64_t footprint_alloc = phys_footprint();
+    ggml_backend_buffer_free(buf);
+    const uint64_t footprint_free = phys_footprint();
+
+    printf("%s: phys_footprint %.1f MiB -> %.1f MiB after freeing %.1f MiB\n", ggml_backend_buft_name(buft), footprint_alloc/1048576.0, footprint_free/1048576.0, size/1048576.0);
+
+    // do not poll, the footprint must drop at free time
+    return footprint_free < footprint_alloc - 0.75*size;
+}
+
+int main() {
+    ggml_backend_load_all();
+
+    bool ok = test_release(ggml_backend_cpu_buffer_type());
+
+    // private Metal buffers are allocated by Metal, not by ggml_metal_host_malloc
+    ggml_backend_dev_t dev = ggml_backend_dev_by_name("MTL0");
+    if (dev && strcmp(ggml_backend_buft_name(ggml_backend_dev_buffer_type(dev)), "MTL0") == 0) {
+        ok = test_release(ggml_backend_dev_buffer_type(dev)) && ok;
+    }
+
+    return ok ? 0 : 1;
+}

tests/test-jinja.cpp

@@ -677,6 +677,24 @@ static void test_set_statement(testing & t) {
 }
 
 static void test_filters(testing & t) {
+    for (const std::string filter : {"upper", "lower", "title", "capitalize"}) {
+        for (const std::string input : {"'hello'", "missing"}) {
+            test_template_error(t, filter + " rejects unknown keyword on " + input,
+                "{{ " + input + "|" + filter + "(attribute='name') }}",
+                "unexpected keyword argument 'attribute'"
+            );
+            test_template_error(t, filter + " rejects extra positional argument on " + input,
+                "{{ " + input + "|" + filter + "('extra') }}",
+                "expected between 1 and 1 arguments, got 2"
+            );
+        }
+    }
+
+    test_template_error(t, "string method rejects unknown keyword",
+        "{{ 'hello'.upper(attribute='name') }}",
+        "unexpected keyword argument 'attribute'"
+    );
+
     test_template(t, "upper",
         "{{ 'hello'|upper }}",
         json::object(),
@@ -725,6 +743,23 @@ static void test_filters(testing & t) {
         "hello"
     );
 
+    test_template(t, "trim chars keyword",
+        "{{ 'xyxhelloxyx'|trim(chars='xy') }}|{{ missing|trim(chars='xy') }}",
+        json::object(),
+        "hello|"
+    );
+
+    for (const std::string input : {"'hello'", "missing"}) {
+        test_template_error(t, "trim rejects unknown keyword on " + input,
+            "{{ " + input + "|trim(attribute='name') }}",
+            "unexpected keyword argument 'attribute'"
+        );
+        test_template_error(t, "trim rejects extra positional argument on " + input,
+            "{{ " + input + "|trim('x', 'y') }}",
+            "expected between 1 and 2 arguments, got 3"
+        );
+    }
+
     test_template(t, "length string",
         "{{ 'hello'|length }}",
         json::object(),
@@ -2064,6 +2099,79 @@ static void test_array_methods(testing & t) {
         "{\"x\": 2}"
     );
 
+    test_template(t, "array|map with filter int materialized",
+        "{{ arr|map('int')|list|join(',') }}",
+        {{"arr", json::array({"1", "2", "3"})}},
+        "1,2,3"
+    );
+
+    test_template(t, "array|map with filter upper for schema union type",
+        "{{ item_value|map('upper')|list|join(',') }}",
+        {{"item_value", json::array({"string", "null"})}},
+        "STRING,NULL"
+    );
+
+    test_template_error(t, "array|map rejects unsupported filter attribute",
+        "{{ [{'name': 'alice'}]|map('upper', attribute='name')|join(',') }}",
+        "unexpected keyword argument 'attribute'"
+    );
+
+    test_template_error(t, "array|map rejects unknown keyword on undefined item",
+        "{{ [missing]|map('lower', attribute='name')|join(',') }}",
+        "unexpected keyword argument 'attribute'"
+    );
+
+    test_template(t, "array|map attribute then filter upper",
+        "{{ arr|map(attribute='name')|map('upper')|join(',') }}",
+        {{"arr", json::array({{{"name", "alice"}}, {{"name", "bob"}}})}},
+        "ALICE,BOB"
+    );
+
+    test_template(t, "array|map with filter arguments",
+        "{{ arr|map('replace', 'a', 'o')|join(',') }}",
+        {{"arr", json::array({"cat", "bat"})}},
+        "cot,bot"
+    );
+
+    test_template(t, "array|map with filter keyword arguments",
+        "{{ arr|map('join', '-', attribute='name')|join(';') }}",
+        {{"arr", json::array({
+            json::array({
+                json({{"name", "a"}}),
+                json({{"name", "b"}}),
+            }),
+            json::array({
+                json({{"name", "c"}}),
+            }),
+        })}},
+        "a-b;c"
+    );
+
+    test_template(t, "array|map with count alias",
+        "{{ arr|map('count')|join(',') }}",
+        {{"arr", json::array({"cat", "hello"})}},
+        "3,5"
+    );
+
+    test_template(t, "array|map with join keyword arguments",
+        "{{ [[{'name': 'a'}, {'name': 'b'}]]|map('join', d='-', attribute='name')|join(';') }}",
+        json::object(),
+        "a-b"
+    );
+
+    for (const std::string input : {"['a', 'b']", "[]", "missing"}) {
+        test_template_error(t, "join rejects unknown keyword on " + input,
+            "{{ " + input + "|join(typo='name') }}",
+            "unexpected keyword argument 'typo'"
+        );
+    }
+
+    test_template(t, "array|map with default alias and arguments",
+        "{{ arr|map('d', 'fallback', true)|join(',') }}",
+        {{"arr", json::array({"", "value"})}},
+        "fallback,value"
+    );
+
     // not used by any chat templates
     // test_template(t, "array.insert()",
     //     "{% set _ = arr.insert(1, 'x') %}{{ arr|join(',') }}",
@@ -2131,6 +2239,12 @@ static void test_array_methods(testing & t) {
         ""
     );
 
+    test_template(t, "undefined|join with keyword arguments",
+        "{{ arr|join(d='-', attribute='name') }}",
+        json::object(),
+        ""
+    );
+
     test_template(t, "undefined|sort",
         "{{ arr|sort|string }}",
         json::object(),

tests/test-llama-archs.cpp

@@ -423,13 +423,6 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
 
     if (arch == LLM_ARCH_MAION_CODER) {
         ms.add_kv(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, std::vector<uint32_t>{1, 1, 0, 0});
-    } else if (arch == LLM_ARCH_BEAM) {
-        std::vector<uint32_t> pattern;
-        pattern.reserve(n_layer);
-        for (uint32_t il = 0; il < n_layer; il++) {
-            pattern.push_back(il % 2);
-        }
-        ms.add_kv(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, pattern);
     } else if (arch == LLM_ARCH_GEMMA4) {
         ms.add_kv(LLM_KV_EMBEDDING_LENGTH_PER_LAYER,      n_embd/2);
         ms.add_kv(LLM_KV_ATTENTION_SHARED_KV_LAYERS,      uint32_t(5));
@@ -445,7 +438,7 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
         ms.add_kv(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, pattern);
     } else if (arch == LLM_ARCH_COHERE2MOE || arch == LLM_ARCH_MIMO2 || arch == LLM_ARCH_STEP35 || arch == LLM_ARCH_SPARK2_5 ||
             arch == LLM_ARCH_MUSE_GLIMMER || arch == LLM_ARCH_MUSE_SPARK || arch == LLM_ARCH_GRANITE_SWA || arch == LLM_ARCH_DOTS3NOTE ||
-            arch == LLM_ARCH_KOLIBRI1 || arch == LLM_ARCH_MAPLE || arch == LLM_ARCH_FLINT) {
+            arch == LLM_ARCH_KOLIBRI1 || arch == LLM_ARCH_MAPLE || arch == LLM_ARCH_FLINT || arch == LLM_ARCH_BEAM) {
         std::vector<uint32_t> pattern;
         pattern.reserve(n_layer);
         for (uint32_t il = 0; il < n_layer; il++) {

tools/imatrix/imatrix.cpp

@@ -1389,7 +1389,7 @@ bool nextn_collector::decode(llama_context * ctx_trunk, const common_batch & bat
     batch.clear();
     if (!pending_h.empty()) {
         const int32_t idx = batch.add(batch_trunk.tokens[0].id, batch_trunk.tokens[0].pos[0], 0, own_lm_head);
-        batch.set_embd(idx, { pending_h.data(), 1, (size_t) n_embd });
+        batch.set_embd_state(idx, { pending_h.data(), 1, (size_t) n_embd });
     }
 
     const float * h_last = nullptr;
@@ -1404,7 +1404,7 @@ bool nextn_collector::decode(llama_context * ctx_trunk, const common_batch & bat
         if (i == n_last) { h_last = h; break; }
 
         const int32_t idx = batch.add(batch_trunk.tokens[i + 1].id, batch_trunk.tokens[i + 1].pos[0], 0, own_lm_head);
-        batch.set_embd(idx, { h, 1, (size_t) n_embd });
+        batch.set_embd_state(idx, { h, 1, (size_t) n_embd });
     }
 
     if (batch.size() > 0) {

tools/server/tests/requirements.txt

@@ -1,4 +1,3 @@
-aiohttp~=3.9.3
 pytest~=8.3.3
 pytest-xdist~=3.6
 filelock~=3.16

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ChatMessageToolCallBlockExecShellCommand.svelte

@@ -8,7 +8,7 @@
 
 	import { parseExecShellCommandMeta } from './parsers/exec-shell-command';
 	import ToolCallBlock from './ToolCallBlock.svelte';
-	import { AlertTriangle, Check, Loader2, XCircle } from '@lucide/svelte';
+	import { AlertTriangle, Check, CircleStop, Loader2, XCircle } from '@lucide/svelte';
 	import { CollapsibleTerminalBlock } from '$lib/components/app';
 	import { SETTINGS_KEYS, TOOL_RUNTIME_SCROLL_AT_BOTTOM_THRESHOLD_PX } from '$lib/constants';
 	import { AgenticSectionType, AttachmentType } from '$lib/enums';
@@ -229,6 +229,7 @@
 {/snippet}
 
 <ToolCallBlock
+	emptyResultIsInterrupted
 	extraLiveStreaming={isLive}
 	{isStreaming}
 	meta={execShellMeta ? { errorMessage: execShellError } : null}
@@ -248,6 +249,12 @@
 				<Loader2 class="h-3 w-3 animate-spin" />
 				Executing command...
 			</div>
+		{:else if ctx.isInterrupted}
+			<div class="flex items-start gap-2 text-xs text-muted-foreground/70 italic">
+				<CircleStop class="mt-0.5 h-3 w-3 shrink-0" />
+
+				<span>Command interrupted before it returned output</span>
+			</div>
 		{:else if execShellError}
 			<div class="flex items-start gap-2 text-xs text-red-600 italic dark:text-red-400">
 				<XCircle class="mt-0.5 h-3 w-3 shrink-0" />

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessageToolCall/ToolCallBlock.svelte

@@ -25,6 +25,7 @@
 		isPending: boolean;
 		isStreamingCall: boolean;
 		isCodeStreaming: boolean;
+		isInterrupted: boolean;
 	}
 
 	interface Props {
@@ -50,6 +51,12 @@
 		 * inline within the body.
 		 */
 		spinIconWhenActive?: boolean;
+		/**
+		 * Treat an empty result as interrupted too. For tools that create
+		 * their result message before running (`exec_shell_command`), so a
+		 * stop mid-run leaves an empty result rather than none.
+		 */
+		emptyResultIsInterrupted?: boolean;
 		/**
 		 * Wrapper component that renders the title row and the body
 		 * children. Defaults to CollapsibleContentBlock;
@@ -65,6 +72,7 @@
 
 	let {
 		children,
+		emptyResultIsInterrupted = false,
 		extraLiveStreaming = false,
 		isStreaming,
 		meta,
@@ -81,6 +89,13 @@
 	const isStreamingCall = $derived(section.type === AgenticSectionType.TOOL_CALL_STREAMING);
 	const showSpinner = $derived(isPending || (isStreamingCall && isStreaming) || extraLiveStreaming);
 	const isCodeStreaming = $derived(isStreaming && (isPending || isStreamingCall));
+	// A finished call with no result was stopped before it returned (user
+	// abort, or the page went away mid-execution).
+	const isInterrupted = $derived(
+		section.type === AgenticSectionType.TOOL_CALL &&
+			(section.toolResult === undefined || (emptyResultIsInterrupted && !section.toolResult)) &&
+			!showSpinner
+	);
 
 	const toolUi: ToolUiEntry | null = $derived(getToolUi(section.toolName));
 	const toolIcon: Component = $derived(
@@ -107,6 +122,8 @@
 
 		if (isStreamingCall && !isStreaming) return 'incomplete';
 
+		if (isInterrupted) return 'interrupted';
+
 		return undefined;
 	}
 
@@ -126,6 +143,7 @@
 >
 	{@render children(meta, {
 		isCodeStreaming,
+		isInterrupted,
 		isPending,
 		isStreaming,
 		isStreamingCall

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessageAgenticContent.svelte

@@ -26,11 +26,16 @@
 
 	let {
 		isLastAssistantMessage = false,
-		isStreaming = false,
+		isStreaming: isChatStreaming = false,
 		message,
 		toolMessages = []
 	}: Props = $props();
 
+	// The chat-level streaming flag is shared by every message; only the last
+	// assistant message is actually live, so earlier ones must not keep their
+	// unresolved tool calls spinning while a newer reply streams.
+	const isStreaming = $derived(isChatStreaming && isLastAssistantMessage);
+
 	let expandedStates: Record<number, boolean> = $state({});
 
 	const showThoughtInProgress = $derived(Boolean(settingsStore.config.showThoughtInProgress));
@@ -48,9 +53,7 @@
 	);
 
 	const pendingPermission = $derived(
-		isStreaming && isLastAssistantMessage
-			? agenticStore.getPendingPermissionRequest(message.convId)
-			: null
+		isStreaming ? agenticStore.getPendingPermissionRequest(message.convId) : null
 	);
 
 	// dismissal applies to the request object, so the next request ( new
@@ -67,9 +70,7 @@
 	}
 
 	const pendingContinue = $derived(
-		isStreaming && isLastAssistantMessage
-			? agenticStore.getPendingContinueRequest(message.convId)
-			: false
+		isStreaming ? agenticStore.getPendingContinueRequest(message.convId) : false
 	);
 
 	let continueDismissed = $state(false);

tools/ui/src/lib/components/app/models/ModelsManager/ModelsManagerModelConfiguration/ModelsManagerModelConfiguration.svelte

@@ -7,6 +7,7 @@
 	import { Trash2 } from '@lucide/svelte';
 	import { DialogConfirmDownload } from '$lib/components/app/dialogs';
 	import { Button } from '$lib/components/ui/button';
+	import * as Tooltip from '$lib/components/ui/tooltip';
 	import { ModelDownloadConfirmAction, ServerModelStatus } from '$lib/enums';
 	import { HuggingFaceService } from '$lib/services';
 	import { deviceStore, modelsStore, uiStore } from '$lib/stores';
@@ -26,6 +27,8 @@
 
 	// the same removal the table row offers, confirmed by the same dialog
 	let deleteOpen = $state(false);
+	// a model the router did not take from its HF cache is not its own to delete
+	let removeBlockedReason = $derived(modelsStore.removeBlockedReason(option.model));
 
 	// the props cache is a plain Map, so the read has to name its version to stay reactive
 	let serverProps = $derived.by(() => {
@@ -130,15 +133,29 @@
 			<ModelsManagerModelConfigurationInformation hub={hubDetails} {option} {serverProps} />
 
 			<div class="mt-4 border-t border-border/30 pt-3 pb-2">
-				<Button
-					class="w-full justify-start gap-2 text-destructive hover:bg-destructive/10 hover:text-destructive dark:hover:bg-destructive/20"
-					onclick={() => (deleteOpen = true)}
-					size="sm"
-					variant="ghost"
-				>
-					<Trash2 class="h-4 w-4" />
-					Delete this model from disk
-				</Button>
+				{#if removeBlockedReason}
+					<!-- a disabled button takes no pointer events, so the wrapper holds the hover -->
+					<Tooltip.Root>
+						<Tooltip.Trigger>
+							{#snippet child({ props })}
+								<div {...props}>
+									{@render deleteButton()}
+								</div>
+							{/snippet}
+						</Tooltip.Trigger>
+
+						<Tooltip.Content class="max-w-64">
+							<p>{removeBlockedReason}</p>
+						</Tooltip.Content>
+					</Tooltip.Root>
+
+					<!-- a phone has no hover, so the reason sits under the button -->
+					{#if deviceStore.isMobile}
+						<p class="px-3 pt-1 text-xs text-muted-foreground">{removeBlockedReason}</p>
+					{/if}
+				{:else}
+					{@render deleteButton()}
+				{/if}
 			</div>
 		{/if}
 	</div>
@@ -178,3 +195,16 @@
 		repoWithTag={option.model}
 	/>
 </div>
+
+{#snippet deleteButton()}
+	<Button
+		class="w-full justify-start gap-2 text-destructive hover:bg-destructive/10 hover:text-destructive dark:hover:bg-destructive/20"
+		disabled={removeBlockedReason !== null}
+		onclick={() => (deleteOpen = true)}
+		size="sm"
+		variant="ghost"
+	>
+		<Trash2 class="h-4 w-4" />
+		Delete this model from disk
+	</Button>
+{/snippet}

tools/ui/src/lib/components/app/models/ModelsManager/row-actions.ts

@@ -49,6 +49,9 @@ export function modelRowActions(
 		];
 	}
 
+	// a model the router did not take from its HF cache is not its own to delete
+	const removeBlockedReason = modelsStore.removeBlockedReason(option.model);
+
 	return [
 		{
 			icon: favorite ? HeartOff : Heart,
@@ -56,10 +59,12 @@ export function modelRowActions(
 			onclick: () => modelsStore.toggleFavorite(option.model)
 		},
 		{
+			disabled: removeBlockedReason !== null,
 			icon: Trash2,
 			label: 'Delete from disk',
 			onclick: () => onDelete(option),
 			separator: true,
+			tooltip: removeBlockedReason ?? undefined,
 			variant: 'destructive' as const
 		},
 		{

tools/ui/src/lib/components/app/navigation/DropdownMenuActions.svelte

@@ -13,6 +13,8 @@
 		onclick: (event: Event) => void;
 		variant?: 'default' | 'destructive';
 		disabled?: boolean;
+		/** Shown on hover, e.g. why a disabled action is not available. */
+		tooltip?: string;
 		shortcut?: string[];
 		separator?: boolean;
 	}
@@ -88,7 +90,14 @@
 					>
 						{@render iconComponent(action.icon, 'h-4 w-4 shrink-0')}
 
-						<span class="flex-1">{action.label}</span>
+						<span class="flex flex-1 flex-col">
+							{action.label}
+
+							<!-- a phone has no hover, so the hint sits under the label -->
+							{#if action.tooltip}
+								<span class="text-xs text-muted-foreground">{action.tooltip}</span>
+							{/if}
+						</span>
 
 						{#if action.shortcut}
 							<KeyboardShortcutInfo keys={action.shortcut} variant={action.variant} />
@@ -131,29 +140,50 @@
 					<DropdownMenu.Separator />
 				{/if}
 
-				<DropdownMenu.Item
-					class="flex items-center justify-between hover:[&>kbd]:opacity-100"
-					disabled={action.disabled}
-					onclick={action.onclick}
-					variant={action.variant}
-				>
-					<div class="flex items-center gap-2">
-						{@render iconComponent(
-							action.icon,
-							`h-4 w-4 ${action.variant === 'destructive' ? 'text-destructive' : ''}`
-						)}
-						{action.label}
-					</div>
+				{#if action.tooltip}
+					<!-- a disabled item takes no pointer events, so the wrapper holds the hover -->
+					<Tooltip.Root>
+						<Tooltip.Trigger>
+							{#snippet child({ props })}
+								<div {...props}>
+									{@render menuItem(action)}
+								</div>
+							{/snippet}
+						</Tooltip.Trigger>
 
-					{#if action.shortcut}
-						<KeyboardShortcutInfo keys={action.shortcut} variant={action.variant} />
-					{/if}
-				</DropdownMenu.Item>
+						<Tooltip.Content class="z-[999999] max-w-64" side="left">
+							<p>{action.tooltip}</p>
+						</Tooltip.Content>
+					</Tooltip.Root>
+				{:else}
+					{@render menuItem(action)}
+				{/if}
 			{/each}
 		</DropdownMenu.Content>
 	</DropdownMenu.Root>
 {/if}
 
+{#snippet menuItem(action: ActionItem)}
+	<DropdownMenu.Item
+		class="flex items-center justify-between hover:[&>kbd]:opacity-100"
+		disabled={action.disabled}
+		onclick={action.onclick}
+		variant={action.variant}
+	>
+		<div class="flex items-center gap-2">
+			{@render iconComponent(
+				action.icon,
+				`h-4 w-4 ${action.variant === 'destructive' ? 'text-destructive' : ''}`
+			)}
+			{action.label}
+		</div>
+
+		{#if action.shortcut}
+			<KeyboardShortcutInfo keys={action.shortcut} variant={action.variant} />
+		{/if}
+	</DropdownMenu.Item>
+{/snippet}
+
 {#snippet iconComponent(IconComponent: Component, className: string)}
 	<IconComponent class={className} />
 {/snippet}

tools/ui/src/lib/components/ui/alert-dialog/alert-dialog-content.svelte

@@ -19,13 +19,13 @@
 	<AlertDialogPrimitive.Content
 		bind:ref
 		class={cn(
-			'fixed z-[999999] grid w-full gap-4 border bg-background p-6 shadow-lg duration-200',
+			'fixed z-[999999] grid w-full gap-4 border border-border bg-popover p-6 text-popover-foreground shadow-md duration-200 dark:border-border/20',
 			// Mobile: Bottom sheet behavior
-			'right-0 bottom-0 left-0 max-h-[100dvh] translate-x-0 translate-y-0 overflow-y-auto rounded-t-lg',
+			'right-0 bottom-0 left-0 max-h-[100dvh] translate-x-0 translate-y-0 overflow-y-auto rounded-t-xl',
 			'data-[state=closed]:animate-out data-[state=closed]:fade-out-0 data-[state=closed]:fill-mode-forwards data-[state=closed]:slide-out-to-bottom-full',
 			'data-[state=open]:animate-in data-[state=open]:fade-in-0 data-[state=open]:slide-in-from-bottom-full',
 			// Desktop: Centered dialog behavior
-			'sm:top-[50%] sm:right-auto sm:bottom-auto sm:left-[50%] sm:max-h-[100vh] sm:max-w-lg sm:translate-x-[-50%] sm:translate-y-[-50%] sm:rounded-lg',
+			'sm:top-[50%] sm:right-auto sm:bottom-auto sm:left-[50%] sm:max-h-[100vh] sm:max-w-lg sm:translate-x-[-50%] sm:translate-y-[-50%] sm:rounded-xl',
 			'sm:data-[state=closed]:slide-out-to-bottom-0 sm:data-[state=closed]:zoom-out-95',
 			'sm:data-[state=open]:slide-in-from-bottom-0 sm:data-[state=open]:zoom-in-95',
 			className

tools/ui/src/lib/components/ui/alert-dialog/alert-dialog-title.svelte

@@ -11,7 +11,7 @@
 
 <AlertDialogPrimitive.Title
 	bind:ref
-	class={cn('text-lg font-semibold', className)}
+	class={cn('text-lg leading-none font-semibold', className)}
 	data-slot="alert-dialog-title"
 	{...restProps}
 />

tools/ui/src/lib/constants/models-manager.constants.ts

@@ -1,6 +1,6 @@
 /** Labels of the models manager table sections. */
 
-import { ModelsTableGroupKind } from '$lib/enums';
+import { ModelsTableGroupKind, ServerModelSource } from '$lib/enums';
 
 export const MODELS_TABLE_GROUP_LABELS: Record<ModelsTableGroupKind, string> = {
 	[ModelsTableGroupKind.DOWNLOADING]: 'Downloading',
@@ -16,3 +16,21 @@ export const MODELS_TABLE_GROUP_LABELS: Record<ModelsTableGroupKind, string> = {
  * so the two stacked sticky rows do not show a gap or overlap by a pixel.
  */
 export const MODELS_TABLE_GROUP_STICKY_OFFSET = 'top: calc(2.25rem - 1px)';
+
+/**
+ * Why a model cannot be deleted from disk, by where the router found it. Only a
+ * model from the HF cache is the router's own to remove.
+ */
+export const MODEL_REMOVE_BLOCKED_REASONS: Record<
+	Exclude<ServerModelSource, ServerModelSource.CACHE>,
+	string
+> = {
+	[ServerModelSource.MODELS_DIR]:
+		'Found in --models-dir: delete the file from that folder to remove it.',
+	[ServerModelSource.PRESET]:
+		'Defined in a --models-preset file: remove its section from the preset to delete it.'
+};
+
+/** Fallback when the router refuses removal without saying where the model came from. */
+export const MODEL_REMOVE_BLOCKED_FALLBACK =
+	'The server does not allow deleting this model: it is not from the Hugging Face cache.';

tools/ui/src/lib/enums/index.ts

@@ -86,7 +86,12 @@ export {
 	ModelsTableSortKey
 } from './model.enums';
 
-export { ServerRole, ServerModelStatus, ServerModelsSseEventType } from './server.enums';
+export {
+	ServerRole,
+	ServerModelSource,
+	ServerModelStatus,
+	ServerModelsSseEventType
+} from './server.enums';
 
 export { ParameterSource, SyncableParameterType, SettingsFieldType } from './settings.enums';
 

tools/ui/src/lib/enums/server.enums.ts

@@ -26,6 +26,16 @@ export enum ServerModelStatus {
 	UNLOADED = 'unloaded'
 }
 
+/**
+ * Where the router found a model - matches server_model_source in
+ * tools/server/server-models.h. Only a cached model can be removed from disk.
+ */
+export enum ServerModelSource {
+	CACHE = 'cache',
+	MODELS_DIR = 'models_dir',
+	PRESET = 'preset'
+}
+
 /**
  * /models/sse event type enum - discriminates the records broadcast on the
  * model status feed in ROUTER mode. Matches the event names emitted by

tools/ui/src/lib/hooks/use-chat-form-model.svelte.ts

@@ -2,10 +2,10 @@
  * Model state for the chat form's action bar.
  *
  * Syncs the conversation's model into the global selection (a conversation
- * opened from history names the model that produced it), falls back to the
- * first loaded model on a router, and reactively tracks which modalities
- * (vision / audio / video) the active model supports - fetching model props
- * from the server on demand if they are not cached yet.
+ * opened from history names the model that produced it), falls back to a
+ * loading, then the first loaded, model on a router, and reactively tracks
+ * which modalities (vision / audio / video) the active model supports -
+ * fetching model props from the server on demand if they are not cached yet.
  */
 
 import { conversationsStore, modelsStore, serverStore } from '$lib/stores';
@@ -36,12 +36,16 @@ export function useChatFormModel() {
 		} else if (
 			isRouter &&
 			!modelsStore.selectedModelId &&
-			modelsStore.loadedModelIds.length > 0 &&
 			conversationsStore.activeMessages.length > 0 &&
 			!conversationModel
 		) {
 			lastSyncedConversationModel = null;
-			const first = modelsStore.models.find((m) => modelsStore.loadedModelIds.includes(m.model));
+			// a model that is loading wins over one that is already up; a model the
+			// selector doesn't list is never picked
+			const visible = modelsStore.getVisibleModels();
+			const first =
+				visible.find((m) => modelsStore.isModelLoading(m.model)) ??
+				visible.find((m) => modelsStore.loadedModelIds.includes(m.model));
 
 			if (first) modelsStore.selectModelById(first.id);
 		}

tools/ui/src/lib/hooks/use-models-selector.svelte.ts

@@ -167,8 +167,14 @@ export function useModelsSelector(opts: UseModelsSelectorOptions): UseModelsSele
 			});
 		}
 
-		// only the built-in server loads on request, and only in router mode
-		if (!onModelChange && isRouter && !modelsStore.isModelLoaded(option.model)) {
+		// only the built-in server loads on request, and only in router mode; a model
+		// that is already loading is only selected, not loaded a second time
+		if (
+			!onModelChange &&
+			isRouter &&
+			!modelsStore.isModelLoaded(option.model) &&
+			!modelsStore.isModelLoading(option.model)
+		) {
 			isLoadingModel = true;
 
 			modelsStore.status

tools/ui/src/lib/stores/chat/index.svelte.ts

@@ -337,7 +337,7 @@ class ChatStore implements ChatStreamHost, ChatFlowsHost {
 
 		const pending = this.consumePendingMessage(convId);
 
-		if (pending) await this.sendMessage(pending.content, pending.extras);
+		if (pending) await this.sendQueuedMessage(convId, pending);
 	}
 
 	consumePendingMessage(
@@ -416,6 +416,20 @@ class ChatStore implements ChatStreamHost, ChatFlowsHost {
 		return this.flows.editUserMessagePreserveResponses(messageId, newContent, newExtras);
 	}
 
+	/**
+	 * Send the message left queued for `convId` once that conversation is
+	 * active and idle (e.g. after navigating back to it).
+	 */
+	async flushPendingMessage(convId: string): Promise<void> {
+		if (conversationsStore.activeConversation?.id !== convId) return;
+
+		if (this.isChatLoadingInternal(convId) || agenticStore.isRunning(convId)) return;
+
+		const pending = this.consumePendingMessage(convId);
+
+		if (pending) await this.sendMessage(pending.content, pending.extras);
+	}
+
 	getAddFilesHandler(): ((files: File[]) => void) | null {
 		return this.addFilesHandler;
 	}
@@ -1141,7 +1155,7 @@ class ChatStore implements ChatStreamHost, ChatFlowsHost {
 					const pending = this.consumePendingMessage(convId);
 
 					if (pending) {
-						this.sendMessage(pending.content, pending.extras);
+						void this.sendQueuedMessage(convId, pending);
 					}
 
 					return;
@@ -1268,7 +1282,7 @@ class ChatStore implements ChatStreamHost, ChatFlowsHost {
 				const pending = agenticStore.consumePendingSteeringMessage(convId);
 
 				if (pending) {
-					await this.sendMessage(pending.content, pending.extras);
+					await this.sendQueuedMessage(convId, pending);
 				}
 
 				return;
@@ -1333,7 +1347,7 @@ class ChatStore implements ChatStreamHost, ChatFlowsHost {
 					const pending = this.consumePendingMessage(convId);
 
 					if (pending) {
-						await this.sendMessage(pending.content, pending.extras);
+						await this.sendQueuedMessage(convId, pending);
 					}
 				},
 				onCompletionId: streamCallbacks.onCompletionId,
@@ -1555,6 +1569,24 @@ class ChatStore implements ChatStreamHost, ChatFlowsHost {
 		}
 	}
 
+	/**
+	 * Re-send a message queued while `convId` was busy. sendMessage targets the
+	 * active conversation, so if the user navigated away the message stays
+	 * queued on its own conversation until flushPendingMessage picks it up.
+	 */
+	private async sendQueuedMessage(
+		convId: string,
+		msg: { content: string; extras?: DatabaseMessageExtra[] }
+	): Promise<void> {
+		if (conversationsStore.activeConversation?.id !== convId) {
+			this.pendingMessages.set(convId, msg);
+
+			return;
+		}
+
+		await this.sendMessage(msg.content, msg.extras);
+	}
+
 	/**
 	 * Streams a fresh summary of `source` into the compaction node and re-stamps
 	 * the reduced context afterwards. `rollback` restores the row when nothing

tools/ui/src/lib/stores/models/index.svelte.ts

@@ -12,10 +12,12 @@ import {
 	FAVORITE_MODELS_LOCALSTORAGE_KEY,
 	HIDDEN_MODELS_LOCALSTORAGE_KEY,
 	MODEL_GROUP_OPEN_LOCALSTORAGE_KEY,
+	MODEL_REMOVE_BLOCKED_FALLBACK,
+	MODEL_REMOVE_BLOCKED_REASONS,
 	RECENT_MODEL_LIMIT,
 	RECENT_MODELS_LOCALSTORAGE_KEY
 } from '$lib/constants';
-import { ReasoningEffort, ServerModelStatus } from '$lib/enums';
+import { ReasoningEffort, ServerModelSource, ServerModelStatus } from '$lib/enums';
 import { HuggingFaceService } from '$lib/services/huggingface.service';
 import { ModelsService } from '$lib/services/models.service';
 // direct imports between stores, not via the barrel, to avoid circular deps
@@ -23,6 +25,7 @@ import { conversationsStore } from '$lib/stores/conversations/index.svelte';
 import { type ModelPropsHost, ModelPropsManager } from '$lib/stores/models/props.svelte';
 import { type ModelStatusHost, ModelStatusManager } from '$lib/stores/models/status.svelte';
 import { serverStore } from '$lib/stores/server.svelte';
+import type { ApiModelDataEntry } from '$lib/types';
 import type { ParsedModelId } from '$lib/types/models';
 import { getConversationModel } from '$lib/utils/conversation-utils';
 import { repoOf } from '$lib/utils/model-names';
@@ -205,12 +208,22 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 	 * Prioritizes:
 	 * 1. Model from active conversation's last assistant response (if loaded)
 	 * 2. Model from active conversation's last assistant response (if not loaded)
-	 * 3. First loaded model (not from active conversation)
-	 * 4. A favorite model
-	 * 5. First available model
+	 * 3. A model that is loading
+	 * 4. First loaded model (not from active conversation)
+	 * 5. A favorite model
+	 * 6. First available model
+	 *
+	 * A chat without a model of its own also trades an idle selection for a
+	 * model that is loading.
 	 */
 	async ensureFirstModelSelected(): Promise<void> {
-		if (this.selectedModelName) return;
+		if (this.selectedModelName) {
+			const loadingModel = this.getLoadingModelToPreselect();
+
+			if (loadingModel) await this.selectModelById(loadingModel.id);
+
+			return;
+		}
 
 		const availableModels = this.getVisibleModels();
 
@@ -233,6 +246,15 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 			}
 		}
 
+		// a model that is loading is the one the user is waiting for
+		const loadingModel = availableModels.find((m) => this.isModelLoading(m.model));
+
+		if (loadingModel) {
+			await this.selectModelById(loadingModel.id);
+
+			return;
+		}
+
 		// Try a loaded model first
 		const loadedModel = availableModels.find((m) => this.isModelLoaded(m.model));
 
@@ -321,6 +343,20 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 		);
 	}
 
+	/**
+	 * The loading model a chat without a model of its own should switch to: none
+	 * while the current selection is loaded or loading itself.
+	 */
+	getLoadingModelToPreselect(): ModelOption | null {
+		if (this.getModelFromLastAssistantResponse()) return null;
+
+		const selected = this.selectedModelName;
+
+		if (selected && (this.isModelLoaded(selected) || this.isModelLoading(selected))) return null;
+
+		return this.getVisibleModels().find((m) => this.isModelLoading(m.model)) ?? null;
+	}
+
 	/**
 	 * Gets the model name from the last assistant message in the active conversation.
 	 * Used by both the chat page and settings page to maintain model consistency.
@@ -345,6 +381,16 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 		return (model?.status?.value as ServerModelStatus) ?? null;
 	}
 
+	/**
+	 * Models the selector lists: not flagged out of the UI (ui === false) and not
+	 * hidden by the user, so an automatic pick never lands on one it doesn't show.
+	 */
+	getVisibleModels(): ModelOption[] {
+		return this.models.filter(
+			(option) => this.props.getModelProps(option.model)?.ui !== false && !this.isHidden(option.id)
+		);
+	}
+
 	hasModel(modelName: string): boolean {
 		return this.models.some((model) => model.model === modelName);
 	}
@@ -367,11 +413,37 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 		return status === ServerModelStatus.LOADED || status === ServerModelStatus.SLEEPING;
 	}
 
+	/** Loading on the server, or a load this client started that has not finished yet. */
+	isModelLoading(modelId: string): boolean {
+		// an unload raises the same in-progress flag, but the model stays loaded until it ends
+		return (
+			this.getModelStatus(modelId) === ServerModelStatus.LOADING ||
+			(this._status.isOperationInProgress(modelId) && !this.isModelLoaded(modelId))
+		);
+	}
+
 	/** Loaded, or sleeping, and not mid-operation: what a row shows as running. */
 	isModelRunning(modelId: string): boolean {
 		return this.isModelLoaded(modelId) && !this._status.isOperationInProgress(modelId);
 	}
 
+	/**
+	 * Why the router would refuse to delete a model from disk, or null when it can.
+	 * Only a model from the HF cache is removable; a server that does not report
+	 * `can_remove` is let through and answers for itself.
+	 */
+	removeBlockedReason(modelId: string): string | null {
+		const model = this.routerModels.find((m) => m.id === modelId);
+
+		if (!model || model.can_remove !== false) return null;
+
+		if (model.source && model.source !== ServerModelSource.CACHE) {
+			return MODEL_REMOVE_BLOCKED_REASONS[model.source];
+		}
+
+		return MODEL_REMOVE_BLOCKED_FALLBACK;
+	}
+
 	/**
 	 * Select a model. `recordRecent` keeps automatic picks out of the recency list.
 	 */
@@ -422,13 +494,24 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 	}
 
 	/**
-	 * Auto-selects the model from the last assistant response if available and loaded.
+	 * Auto-selects the model from the last assistant response if available and loaded,
+	 * or a loading model when the conversation has no model yet.
 	 * Returns true if a model was selected, false otherwise.
 	 */
 	async selectModelFromLastAssistantResponse(): Promise<boolean> {
 		const lastModel = this.getModelFromLastAssistantResponse();
 
-		if (!lastModel || this.selectedModelName === lastModel) return false;
+		if (!lastModel) {
+			const loadingModel = this.getLoadingModelToPreselect();
+
+			if (!loadingModel) return false;
+
+			await this.selectModelById(loadingModel.id);
+
+			return true;
+		}
+
+		if (this.selectedModelName === lastModel) return false;
 
 		const matchingModel = this.models.find((option) => option.model === lastModel);
 
@@ -612,19 +695,12 @@ class ModelsStore implements ModelPropsHost, ModelStatusHost {
 				})
 		);
 	}
-
 	/** Fetch models in MODEL mode (single model, standard OpenAI-compatible). */
 	private async fetchModelModeInternal(): Promise<ModelOption[]> {
 		const response = await ModelsService.list();
 
 		return this.buildModelOptions(response);
 	}
-	/**
-	 * Filter to models visible in the UI (ui !== false).
-	 */
-	private getVisibleModels(): ModelOption[] {
-		return this.models.filter((option) => this.props.getModelProps(option.model)?.ui !== false);
-	}
 
 	private loadFavoritesFromStorage(): Set<string> {
 		try {

tools/ui/src/lib/types/api.d.ts

@@ -3,6 +3,7 @@ import type { ModelSidecar } from '$lib/constants';
 import type {
 	ContentPartType,
 	FileTypeAudio,
+	ServerModelSource,
 	ServerModelsSseEventType,
 	ServerModelStatus,
 	ServerRole
@@ -95,6 +96,10 @@ export interface ApiModelDataEntry {
 	path: string;
 	/** Current status of the model */
 	status: ApiModelStatus;
+	/** Where the router found the model: the HF cache, --models-dir or a preset */
+	source?: ServerModelSource;
+	/** Whether the router can delete the model's files (only cached models) */
+	can_remove?: boolean;
 	/** Alternative names that resolve to this model */
 	aliases?: string[];
 	/** Informational tags for this model */

tools/ui/src/routes/(chat)/chat/[id]/+page.svelte

@@ -106,6 +106,8 @@
 				chatStore.syncLoadingStateForChat(chatId);
 				// server probe (with localStorage fallback) and attach
 				await chatStore.discoverActiveStream(chatId);
+				// a message queued here may have finished its turn while another chat was open
+				await chatStore.flushPendingMessage(chatId);
 
 				if ((qParam !== null || modelParam !== null) && !urlParamsProcessed) {
 					await handleUrlParams();

tools/ui/tests/client/exec-shell-command-streaming.svelte.test.ts

@@ -81,6 +81,7 @@ describe('exec_shell_command streaming title', () => {
 
 	it('streams the partial command into the title with a cursor', async () => {
 		const { container } = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: messageWithArgs('{"command":"git sta'),
 			toolMessages: []
@@ -98,6 +99,7 @@ describe('exec_shell_command streaming title', () => {
 
 	it('drops the cursor once the args JSON is complete', async () => {
 		const screen = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: messageWithArgs('{"command":"git sta'),
 			toolMessages: []
@@ -108,6 +110,7 @@ describe('exec_shell_command streaming title', () => {
 		expect(hasCursor(screen.container)).toBe(true);
 
 		screen.rerender({
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: messageWithArgs('{"command":"git status"}'),
 			toolMessages: []
@@ -120,6 +123,7 @@ describe('exec_shell_command streaming title', () => {
 
 	it('switches the prompt from the conversation cwd to the persisted one', async () => {
 		const screen = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: messageWithArgs('{"command":"git status"}'),
 			toolMessages: []
@@ -131,6 +135,7 @@ describe('exec_shell_command streaming title', () => {
 
 		// Result message lands: its x-tool-cwd takes over the prompt.
 		screen.rerender({
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: messageWithArgs('{"command":"git status"}'),
 			toolMessages: [toolResultMessage('', '/other/dir')]
@@ -143,6 +148,7 @@ describe('exec_shell_command streaming title', () => {
 
 	it('shows the executing label only after the command finishes streaming', async () => {
 		const screen = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: messageWithArgs('{"command":"git sta'),
 			toolMessages: []
@@ -154,6 +160,7 @@ describe('exec_shell_command streaming title', () => {
 		expect(screen.container.textContent).not.toContain('Executing command...');
 
 		screen.rerender({
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: messageWithArgs('{"command":"git status"}'),
 			toolMessages: []
@@ -192,6 +199,7 @@ describe('exec_shell_command streaming title', () => {
 		conversationsStore.activeMessages = [messageWithArgs('')];
 
 		const screen = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: conversationsStore.activeMessages[0]!,
 			toolMessages: []
@@ -214,3 +222,57 @@ describe('exec_shell_command streaming title', () => {
 		}
 	});
 });
+
+describe('exec_shell_command interrupted call', () => {
+	beforeEach(() => {
+		settingsStore.updateMultipleConfig({
+			[SETTINGS_KEYS.ALWAYS_SHOW_TOOL_CALL_CONTENT]: true,
+			[SETTINGS_KEYS.SHOW_TOOL_CALL_IN_PROGRESS]: true
+		});
+	});
+
+	it('marks a call without a result as interrupted once streaming stops', async () => {
+		const { container } = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
+			isStreaming: false,
+			message: messageWithArgs('{"command":"git status"}'),
+			toolMessages: []
+		});
+
+		await tick();
+
+		expect(triggerText(container)).toContain('interrupted');
+		expect(container.textContent).toContain('Command interrupted');
+		expect(container.textContent).not.toContain('Executing command...');
+	});
+
+	it('marks a command stopped before any output as interrupted', async () => {
+		// exec_shell_command creates its result message before running, so a
+		// stop mid-run leaves an empty result rather than none.
+		const { container } = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
+			isStreaming: false,
+			message: messageWithArgs('{"command":"sleep 40"}'),
+			toolMessages: [toolResultMessage('', CWD)]
+		});
+
+		await tick();
+
+		expect(triggerText(container)).toContain('interrupted');
+		expect(container.textContent).toContain('Command interrupted');
+	});
+
+	it('does not keep spinning in an earlier message while a newer reply streams', async () => {
+		const { container } = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: false,
+			isStreaming: true,
+			message: messageWithArgs('{"command":"git status"}'),
+			toolMessages: []
+		});
+
+		await tick();
+
+		expect(triggerText(container)).toContain('interrupted');
+		expect(container.textContent).not.toContain('Executing command...');
+	});
+});

tools/ui/tests/client/tool-call-in-progress.svelte.test.ts

@@ -95,6 +95,7 @@ describe('Show tool call in progress', () => {
 		setSettings({ alwaysShow: false, inProgress: true });
 
 		const { container } = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: []
@@ -107,6 +108,7 @@ describe('Show tool call in progress', () => {
 
 	it('keeps the call collapsed while executing when the setting is off', async () => {
 		const { container } = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: []
@@ -123,6 +125,7 @@ describe('Show tool call in progress', () => {
 		setSettings({ alwaysShow: false, inProgress: true });
 
 		const screen = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: []
@@ -135,6 +138,7 @@ describe('Show tool call in progress', () => {
 
 		// Same block, now finished: it must close on its own, no toggle given.
 		screen.rerender({
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: [toolResultMessage()]
@@ -150,6 +154,7 @@ describe('Show tool call in progress', () => {
 		setSettings({ alwaysShow: false, inProgress: true });
 
 		const screen = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: []
@@ -161,6 +166,7 @@ describe('Show tool call in progress', () => {
 		await collapseBlock(screen);
 
 		screen.rerender({
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: [toolResultMessage()]
@@ -175,6 +181,7 @@ describe('Show tool call in progress', () => {
 		setSettings({ alwaysShow: true, inProgress: false });
 
 		const { container } = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: [toolResultMessage()]
@@ -189,6 +196,7 @@ describe('Show tool call in progress', () => {
 		setSettings({ alwaysShow: false, inProgress: true });
 
 		const screen = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: []
@@ -206,6 +214,7 @@ describe('Show tool call in progress', () => {
 
 	it('lets a manual expand win over the collapsed default', async () => {
 		const screen = render(ChatMessageAgenticContent, {
+			isLastAssistantMessage: true,
 			isStreaming: true,
 			message: assistantMessage(),
 			toolMessages: []

tools/ui/tests/stories/ModelsManager.stories.svelte

@@ -6,7 +6,7 @@
 		ModelsTableGroup
 	} from '$lib/components/app/models/ModelsManager/utils';
 	import { MODELS_TABLE_GROUP_LABELS } from '$lib/constants';
-	import { ModelGroupKind, ModelsTableGroupKind } from '$lib/enums';
+	import { ModelGroupKind, ModelsTableGroupKind, ServerModelStatus } from '$lib/enums';
 	import { modelsStore } from '$lib/stores/models/index.svelte';
 	import type { ModelOption } from '$lib/types/models';
 
@@ -60,7 +60,7 @@
 			object: 'model',
 			owned_by: 'llamacpp',
 			path: `/models/${qwen.model}`,
-			status: { value: 'loaded' }
+			status: { value: ServerModelStatus.LOADED }
 		}
 	];
 </script>

tools/ui/tests/unit/queued-message-isolation.test.ts

@@ -0,0 +1,81 @@
+// A message queued (steering or pending) while a conversation is busy belongs
+// to that conversation: when its turn ends while another chat is open, it must
+// stay queued there instead of being sent into the chat the user is viewing.
+
+import { beforeEach, describe, expect, it, vi } from 'vitest';
+
+vi.mock('$lib/services/database.service', () => ({
+	DatabaseService: {
+		getConversation: vi.fn(),
+		getConversationMessages: vi.fn(),
+		updateConversation: vi.fn(),
+		updateCurrentNode: vi.fn()
+	}
+}));
+
+import { DatabaseService } from '$lib/services/database.service';
+import { chatStore } from '$lib/stores/chat/index.svelte';
+import { conversationsStore } from '$lib/stores/conversations/index.svelte';
+
+type QueuedMessage = { content: string };
+type QueueAccess = { sendQueuedMessage(convId: string, msg: QueuedMessage): Promise<void> };
+
+const sendQueuedMessage = (convId: string, msg: QueuedMessage) =>
+	(chatStore as unknown as QueueAccess).sendQueuedMessage(convId, msg);
+
+async function openConversation(id: string) {
+	vi.mocked(DatabaseService.getConversation).mockResolvedValue({
+		currNode: null,
+		id,
+		lastModified: 1000,
+		name: id
+	});
+	vi.mocked(DatabaseService.getConversationMessages).mockResolvedValue([]);
+
+	expect(await conversationsStore.loadConversation(id)).toBe(true);
+}
+
+beforeEach(() => {
+	vi.restoreAllMocks();
+	chatStore.clearPendingMessage('conv-a');
+	chatStore.clearPendingMessage('conv-b');
+});
+
+describe('queued message isolation', () => {
+	it('keeps the message on its own conversation when another one is active', async () => {
+		const send = vi.spyOn(chatStore, 'sendMessage').mockResolvedValue();
+
+		await openConversation('conv-b');
+		await sendQueuedMessage('conv-a', { content: 'steer a' });
+
+		expect(send).not.toHaveBeenCalled();
+		expect(chatStore.getPendingMessageContent('conv-a')).toBe('steer a');
+		expect(chatStore.getPendingMessageContent('conv-b')).toBeNull();
+	});
+
+	it('sends right away when its conversation is still active', async () => {
+		const send = vi.spyOn(chatStore, 'sendMessage').mockResolvedValue();
+
+		await openConversation('conv-a');
+		await sendQueuedMessage('conv-a', { content: 'steer a' });
+
+		expect(send).toHaveBeenCalledWith('steer a', undefined);
+		expect(chatStore.hasPendingMessage('conv-a')).toBe(false);
+	});
+
+	it('flushes the left-behind message only once its conversation is active again', async () => {
+		const send = vi.spyOn(chatStore, 'sendMessage').mockResolvedValue();
+
+		chatStore.injectPendingMessage('conv-a', 'steer a');
+		await openConversation('conv-b');
+		await chatStore.flushPendingMessage('conv-a');
+
+		expect(send).not.toHaveBeenCalled();
+
+		await openConversation('conv-a');
+		await chatStore.flushPendingMessage('conv-a');
+
+		expect(send).toHaveBeenCalledWith('steer a', undefined);
+		expect(chatStore.hasPendingMessage('conv-a')).toBe(false);
+	});
+});