From 8e330954adb6e86c329c9d7e338f01f93ffe4b88 Mon Sep 17 00:00:00 2001
From: Xuan-Son Nguyen <son@huggingface.co>
Date: Sun, 13 Sep 2026 01:36:34 +0200
Subject: [PATCH 01/34] common: add LOG_JSON macro to log structured data
 (#28586)

* add LOG_JSON macro

* fit: add demo LOG_JSON
---
 common/arg.cpp |  2 +-
 common/fit.cpp | 41 +++++++++++++++++++++++++++
 common/log.cpp | 76 +++++++++++++++++++++++++++++++++++++++++---------
 common/log.h   | 19 ++++++++++++-
 4 files changed, 123 insertions(+), 15 deletions(-)

diff --git a/common/arg.cpp b/common/arg.cpp
index b1c0f2352..c4c4e143c 100644
--- a/common/arg.cpp
+++ b/common/arg.cpp
@@ -3875,7 +3875,7 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
         {"--no-log-jsonl"},
         "Log as JSONL (one JSON object per line) to stdout, this also disables colored logging (default: disabled)",
         [](common_params &, bool value) {
-            common_log_set_jsonl(common_log_main(), value);
+            common_log_set_jsonl(value);
         }
     ).set_env("LLAMA_ARG_LOG_JSONL"));
     add_opt(common_arg(
diff --git a/common/fit.cpp b/common/fit.cpp
index c601fe405..7a0300829 100644
--- a/common/fit.cpp
+++ b/common/fit.cpp
@@ -1,5 +1,6 @@
 #include "fit.h"
 
+#include "json.h"
 #include "log.h"
 
 #include "../src/llama-ext.h"
@@ -915,6 +916,9 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
 
     std::vector<std::array<std::string, 9>> table_data;
     table_data.reserve(devices.size());
+
+    // same data as the table below, for --log-jsonl consumers
+    common_json rows = common_json::array();
     const std::string template_header = "%s: | %s | %s   %s    %s   %s   %s   %s    %s |\n";
     const std::string template_gpu    = "%s: | %s | %s = %s + (%s = %s + %s + %s) + %s |\n";
     const std::string template_other  = "%s: | %s | %s   %s    %s = %s + %s + %s    %s |\n";
@@ -989,6 +993,19 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
             std::to_string(mb.context / MiB),
             std::to_string(mb.compute / MiB),
             std::to_string(unaccounted / static_cast<int64_t>(MiB))});
+
+        rows.push_back({
+            {"kind",        "device"},
+            {"name",        name},
+            {"description", desc},
+            {"total",       total / MiB},
+            {"free",        free / MiB},
+            {"self",        self / MiB},
+            {"model",       mb.model / MiB},
+            {"context",     mb.context / MiB},
+            {"compute",     mb.compute / MiB},
+            {"unaccounted", unaccounted / static_cast<int64_t>(MiB)},
+        });
     }
 
     // print memory breakdown for host:
@@ -1004,6 +1021,15 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
             std::to_string(mb_host.context / MiB),
             std::to_string(mb_host.compute / MiB),
             ""}); // unaccounted
+
+        rows.push_back({
+            {"kind",    "host"},
+            {"name",    "Host"},
+            {"self",    self / MiB},
+            {"model",   mb_host.model / MiB},
+            {"context", mb_host.context / MiB},
+            {"compute", mb_host.compute / MiB},
+        });
     }
 
     // print memory breakdown for all remaining buffer types:
@@ -1025,6 +1051,16 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
             std::to_string(mb.context / MiB),
             std::to_string(mb.compute / MiB),
             ""}); // unaccounted
+
+        rows.push_back({
+            {"kind",    "buffer_type"},
+            {"name",    name},
+            {"self",    self / MiB},
+            {"model",   mb.model / MiB},
+            {"context", mb.context / MiB},
+            {"compute", mb.compute / MiB},
+        });
+
         seen_buffer_types.insert(buft);
     }
 
@@ -1042,6 +1078,11 @@ void common_memory_breakdown_print(const struct llama_context * ctx) {
             __func__, td[1].c_str(), td[2].c_str(), td[3].c_str(), td[4].c_str(), td[5].c_str(),
             td[6].c_str(), td[7].c_str(), td[8].c_str());
     }
+
+    LOG_JSON("fit_memory_breakdown", common_json({
+        {"unit", "MiB"},
+        {"rows", rows},
+    }));
 }
 
 void common_fit_print(
diff --git a/common/log.cpp b/common/log.cpp
index 42951190c..0a9a4eb9e 100644
--- a/common/log.cpp
+++ b/common/log.cpp
@@ -37,6 +37,16 @@ void common_log_set_verbosity_thold(int verbosity) {
     common_log_verbosity_thold = verbosity;
 }
 
+static bool common_log_jsonl = false;
+
+bool common_log_get_jsonl(void) {
+    return common_log_jsonl;
+}
+
+void common_log_set_jsonl(bool jsonl) {
+    common_log_jsonl = jsonl;
+}
+
 static int64_t t_us() {
     return std::chrono::duration_cast<std::chrono::microseconds>(std::chrono::system_clock::now().time_since_epoch()).count();
 }
@@ -87,6 +97,7 @@ struct common_log_entry {
     bool is_end       { false }; // signals the worker thread to stop
     bool prefix       { false };
     bool jsonl        { false };
+    bool is_json      { false }; // msg already holds a serialized JSON object
 
     common_log_entry(size_t size = 256) : msg(size) { }
 
@@ -107,6 +118,12 @@ struct common_log_entry {
         }
 
         if (jsonl) {
+            if (is_json) {
+                fprintf(fcur, "%s\n", msg.data());
+                fflush(fcur);
+                return;
+            }
+
             common_json obj = {
                 {"type",  "log"},
                 {"time",  timestamp},
@@ -156,7 +173,6 @@ struct common_log {
         file       = nullptr;
         prefix     = false;
         timestamps = false;
-        jsonl      = false;
         running    = false;
         t_start    = t_us();
 
@@ -184,7 +200,6 @@ private:
 
     bool prefix;
     bool timestamps;
-    bool jsonl;
     bool running;
 
     int64_t t_start;
@@ -273,7 +288,8 @@ public:
         entry.is_end    = false;
         entry.level     = level;
         entry.prefix    = prefix;
-        entry.jsonl     = jsonl;
+        entry.jsonl     = common_log_jsonl;
+        entry.is_json   = false;
         entry.timestamp = 0;
         if (timestamps) {
             entry.timestamp = t_us() - t_start;
@@ -283,6 +299,42 @@ public:
         cv_new.notify_one();
     }
 
+    void add_json(const char * type, const common_json & obj) {
+        const common_json full = {
+            {"type", type},
+            {"data", obj},
+        };
+
+        const std::string text = full.dump_safe();
+
+        std::unique_lock<std::mutex> lock(mtx);
+
+        // block if the queue is full
+        cv_full.wait(lock, [this]() { return !running || !is_full(); });
+
+        if (!running) {
+            // discard messages while the worker thread is paused
+            return;
+        }
+
+        auto & entry = queue[tail];
+
+        if (entry.msg.size() < text.size() + 1) {
+            entry.msg.resize(text.size() + 1);
+        }
+        memcpy(entry.msg.data(), text.c_str(), text.size() + 1);
+
+        entry.is_end    = false;
+        entry.level     = GGML_LOG_LEVEL_NONE;
+        entry.prefix    = false;
+        entry.jsonl     = true;
+        entry.is_json   = true;
+        entry.timestamp = 0;
+
+        tail = (tail + 1) % queue.size();
+        cv_new.notify_one();
+    }
+
     void resume() {
         std::lock_guard<std::mutex> lock(mtx);
 
@@ -388,12 +440,6 @@ public:
 
         this->timestamps = timestamps;
     }
-
-    void set_jsonl(bool jsonl) {
-        std::lock_guard<std::mutex> lock(mtx);
-
-        this->jsonl = jsonl;
-    }
 };
 
 //
@@ -440,6 +486,14 @@ void common_log_add(struct common_log * log, enum ggml_log_level level, const ch
     va_end(args);
 }
 
+void common_log_add_json(struct common_log * log, const char * type, const common_json & obj) {
+    if (!common_log_jsonl) {
+        return;
+    }
+
+    log->add_json(type, obj);
+}
+
 void common_log_set_file(struct common_log * log, const char * file) {
     log->set_file(file);
 }
@@ -467,10 +521,6 @@ void common_log_set_timestamps(struct common_log * log, bool timestamps) {
     log->set_timestamps(timestamps);
 }
 
-void common_log_set_jsonl(struct common_log * log, bool jsonl) {
-    log->set_jsonl(jsonl);
-}
-
 void common_log_flush(struct common_log * log) {
     log->pause();
     log->resume();
diff --git a/common/log.h b/common/log.h
index 37f4de92b..e36b09463 100644
--- a/common/log.h
+++ b/common/log.h
@@ -43,6 +43,10 @@ int  common_log_get_verbosity_thold(void);
 
 void common_log_set_verbosity_thold(int verbosity); // not thread-safe
 
+bool common_log_get_jsonl(void);
+
+void common_log_set_jsonl(bool jsonl); // not thread-safe
+
 int common_log_get_verbosity(enum ggml_log_level level);
 
 void common_log_default_callback(enum ggml_log_level level, const char * text, void * user_data);
@@ -91,7 +95,6 @@ void common_log_set_file      (struct common_log * log, const char * file); // n
 void common_log_set_colors    (struct common_log * log, log_colors colors); // not thread-safe
 void common_log_set_prefix    (struct common_log * log, bool prefix);       // whether to output prefix to each log
 void common_log_set_timestamps(struct common_log * log, bool timestamps);   // whether to output timestamps in the prefix
-void common_log_set_jsonl     (struct common_log * log, bool jsonl);        // print each log as a JSON object on one line, not thread-safe
 void common_log_flush         (struct common_log * log);                    // flush all pending log messages
 
 // helper macros for logging
@@ -127,3 +130,17 @@ void common_log_flush         (struct common_log * log);                    // f
 #define LOG_WRNV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_WARN,  verbosity, __VA_ARGS__)
 #define LOG_ERRV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_ERROR, verbosity, __VA_ARGS__)
 #define LOG_CNTV(verbosity, ...) LOG_TMPL(GGML_LOG_LEVEL_CONT,  verbosity, __VA_ARGS__)
+
+class common_json; // defined in common/json.h
+
+// helper allows different types of json output
+// no-op if --log-jsonl is not set
+void common_log_add_json(struct common_log * log, const char * type, const common_json & data);
+
+// will only print if --log-jsonl is set
+#define LOG_JSON(type, data) \
+    do { \
+        if (common_log_get_jsonl()) { \
+            common_log_add_json(common_log_main(), type, data); \
+        } \
+    } while (0)

From 790cf51aabd61763486050dec7451d9147cb7c61 Mon Sep 17 00:00:00 2001
From: Aldehir Rojas <hello@alde.dev>
Date: Sat, 12 Sep 2026 19:08:52 -0500
Subject: [PATCH 02/34] chat : improve parsing of complex types in qwen3-coder
 (#28742)

* chat : improve schema support in qwen3 parser

* cont : clean up grammar a bit
---
 common/parsers/qwen3-coder.cpp | 31 ++++++++++++++++--
 tests/test-chat.cpp            | 59 ++++++++++++++++++++++++++++++++++
 2 files changed, 87 insertions(+), 3 deletions(-)

diff --git a/common/parsers/qwen3-coder.cpp b/common/parsers/qwen3-coder.cpp
index dfc744084..7938a2027 100644
--- a/common/parsers/qwen3-coder.cpp
+++ b/common/parsers/qwen3-coder.cpp
@@ -104,9 +104,34 @@ common_chat_params common_chat_params_init_qwen3_coder(const common_chat_templat
 
                     auto arg_open = p.tool_arg_open("<parameter=" + p.tool_arg_name(p.literal(param.name)) + ">\n");
 
-                    auto arg_value = param.schema->may_be_string() ?
-                        arg_string :
-                        p.tool_arg_json_value(p.schema(p.json(), rule_name + "-schema", doc, *param.schema)) + arg_close;
+                    auto types = param.schema->value_types();
+
+                    auto arg_value = p.eps();
+                    if (!types.has(common_chat_schema::TYPE_STRING)) {
+                        arg_value = p.tool_arg_json_value(p.schema(p.json(), rule_name + "-schema", doc, *param.schema)) + arg_close;
+                    } else if (types.is_only(common_chat_schema::TYPE_STRING)) {
+                        arg_value = arg_string;
+                    } else {
+                        // The string alternative accepts any text, so the grammar only keeps the raw string
+                        // rule. The parser still tries the JSON alternatives first to type the value.
+                        auto json_value = p.choice();
+                        if (types.has(common_chat_schema::TYPE_OBJECT)) {
+                            json_value |= p.json_object();
+                        }
+                        if (types.has(common_chat_schema::TYPE_ARRAY)) {
+                            json_value |= p.json_array();
+                        }
+                        if (types.has(common_chat_schema::TYPE_NUMBER) || types.has(common_chat_schema::TYPE_INTEGER)) {
+                            json_value |= p.json_number();
+                        }
+                        if (types.has(common_chat_schema::TYPE_BOOLEAN)) {
+                            json_value |= p.json_bool();
+                        }
+                        if (types.has(common_chat_schema::TYPE_NULL)) {
+                            json_value |= p.json_null();
+                        }
+                        arg_value = p.gbnf(p.atomic(p.tool_arg_json_value(json_value) + arg_close) | arg_string, "xml-arg-string");
+                    }
 
                     auto arg_rule = p.rule(rule_name, p.tool_arg(arg_open + arg_value));
 
diff --git a/tests/test-chat.cpp b/tests/test-chat.cpp
index 1aef83f43..30a7237e3 100644
--- a/tests/test-chat.cpp
+++ b/tests/test-chat.cpp
@@ -846,6 +846,25 @@ static common_chat_tool nullable_int_tool{
     })",
 };
 
+static common_chat_tool string_union_tool{
+    /* .name = */ "set_union",
+    /* .description = */ "Set values whose types are unions with string",
+    /* .parameters = */ R"({
+        "type": "object",
+        "properties": {
+            "value": {
+                "type": ["string", "object"],
+                "description": "A string or object value"
+            },
+            "amount": {
+                "type": ["string", "integer"],
+                "description": "A string or integer value"
+            }
+        },
+        "required": ["value", "amount"]
+    })",
+};
+
 static common_chat_tool enum_no_type_tool{
     /* .name = */ "set_unit",
     /* .description = */ "Set a temperature unit",
@@ -3805,6 +3824,46 @@ static void test_template_output_peg_parsers(bool detailed_debug) {
             })
             .run();
 
+        // nullable string given null - parses as JSON null, not the string "null"
+        tst.test(
+               "<tool_call>\n"
+               "<function=set_nullable_str>\n"
+               "<parameter=name>\nnull\n</parameter>\n"
+               "</function>\n"
+               "</tool_call>")
+            .tools({ nullable_string_tool })
+            .expect_tool_calls({
+                { "set_nullable_str", R"({"name": null})", {} },
+            })
+            .run();
+
+        // unions with string - JSON values of the other types are typed, everything else is a string
+        tst.test(
+               "<tool_call>\n"
+               "<function=set_union>\n"
+               "<parameter=value>\n{\"a\": 1}\n</parameter>\n"
+               "<parameter=amount>\n2 dollars\n</parameter>\n"
+               "</function>\n"
+               "</tool_call>")
+            .tools({ string_union_tool })
+            .expect_tool_calls({
+                { "set_union", R"({"value": {"a": 1}, "amount": "2 dollars"})", {} },
+            })
+            .run();
+
+        tst.test(
+               "<tool_call>\n"
+               "<function=set_union>\n"
+               "<parameter=value>\n{not valid json\n</parameter>\n"
+               "<parameter=amount>\n42\n</parameter>\n"
+               "</function>\n"
+               "</tool_call>")
+            .tools({ string_union_tool })
+            .expect_tool_calls({
+                { "set_union", R"({"value": "{not valid json", "amount": 42})", {} },
+            })
+            .run();
+
         // enum without explicit type key - should infer string from enum values
         tst.test(
                "<tool_call>\n"

From 56b9eb280a67796379d8625729fb03d72c70789d Mon Sep 17 00:00:00 2001
From: Hongqiang Wang <wangh@qti.qualcomm.com>
Date: Sat, 12 Sep 2026 21:33:23 -0700
Subject: [PATCH 03/34] opencl: apply the noshuffle row-alignment rule to q4_K,
 q5_K and q8_0, not just q6_K (#28575)

---
 ggml/src/ggml-opencl/ggml-opencl.cpp | 17 ++++++++++++++---
 1 file changed, 14 insertions(+), 3 deletions(-)

diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index c107281a2..39c592e88 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -8302,9 +8302,20 @@ inline bool use_adreno_kernels(const ggml_backend_opencl_context *backend_ctx, c
     bool threashold_ok = tensor->ne[0] >= threshold_ne0 && tensor->ne[1] >= threshold_ne1 &&
             tensor->ne[2] == 1 && tensor->ne[3] == 1;
 
-    // q6_K adreno kernels requires ne1 is multiple of 128
-    if (tensor->type == GGML_TYPE_Q6_K) {
-        return threashold_ok && tensor->ne[1] % 128 == 0;
+    // The noshuffle layout packs 2 rows per 32-bit texel and the GEMV reads it at an
+    // ne1/2 texel stride with an exact-cover dispatch, so it is only addressable when
+    // ne1 is a multiple of 64; an unaligned ne1 truncates the stride and the weight is
+    // read misaligned. That is a property of the layout, not of one quant -- q4_K, q5_K
+    // and q8_0 read the same packing as q6_K. The bound is 64, not 128: a q8_0 attention
+    // weight of ne1 = 2880 is a multiple of 64 but not 128 and is correct.
+    switch (tensor->type) {
+        case GGML_TYPE_Q4_K:
+        case GGML_TYPE_Q5_K:
+        case GGML_TYPE_Q6_K:
+        case GGML_TYPE_Q8_0:
+            return threashold_ok && tensor->ne[1] % 64 == 0;
+        default:
+            break;
     }
     return threashold_ok;
 }

From f1e44dcc11d8802d107bd7331a3d3fd3e6f57b93 Mon Sep 17 00:00:00 2001
From: Jeff Bolz <jbolz@nvidia.com>
Date: Sun, 13 Sep 2026 01:18:19 -0500
Subject: [PATCH 04/34] vulkan: workaround NV queuesubmit driver bug (#28830)

There is a driver bug where two queues on the same VkDevice simultaneously
submitting can break some internal synchronization. Until it's fixed, add a
mutex around queuesubmit.
---
 ggml/src/ggml-vulkan/ggml-vulkan.cpp | 18 +++++++++++++++++-
 1 file changed, 17 insertions(+), 1 deletion(-)

diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
index b28fdc9bb..0dfa44dbf 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
@@ -333,6 +333,7 @@ static void ggml_vk_print_device_lost_info(const vk_device& device);
 struct vk_queue_handle {
     vk::Queue queue;
     vk_device_ref device;
+    std::mutex * device_submit_mutex = nullptr;
     virtual void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) = 0;
     virtual void lock()   {}   // no-op by default (internally synchronized case)
     virtual void unlock() {}
@@ -342,6 +343,11 @@ struct vk_queue_handle {
 struct vk_queue_handle_synchronized : vk_queue_handle {
     std::mutex mutex;
     void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) override {
+        // Workaround for NVIDIA driver bug
+        std::unique_lock<std::mutex> device_guard;
+        if (device_submit_mutex) {
+            device_guard = std::unique_lock<std::mutex>(*device_submit_mutex);
+        }
         std::lock_guard<std::mutex> guard(mutex);
         try {
             queue.submit(submits, fence);
@@ -356,9 +362,14 @@ struct vk_queue_handle_synchronized : vk_queue_handle {
     void unlock() override { mutex.unlock(); }
 };
 
+// Driver guarantees internal synchronization via VK_KHR_internally_synchronized_queues
 struct vk_queue_handle_unsynchronized : vk_queue_handle {
     void submit(vk::ArrayProxy<const vk::SubmitInfo> submits, vk::Fence fence) override {
-        // Driver guarantees internal synchronization via VK_KHR_internally_synchronized_queues
+        // Workaround for NVIDIA driver bug
+        std::unique_lock<std::mutex> device_guard;
+        if (device_submit_mutex) {
+            device_guard = std::unique_lock<std::mutex>(*device_submit_mutex);
+        }
         try {
             queue.submit(submits, fence);
         } catch (vk::DeviceLostError &) {
@@ -835,6 +846,7 @@ static bool ggml_vk_lightning_indexer_k_type_supported(ggml_type type) {
 
 struct vk_device_struct {
     std::recursive_mutex mutex;
+    std::mutex queue_submit_mutex;
     mutable std::shared_mutex pinned_memory_mutex;
 
     // Guards compile_pending, all_pipelines, and the dynamic pipeline maps
@@ -3520,6 +3532,10 @@ static std::unique_ptr<vk_queue> ggml_vk_create_queue(vk_device& device, uint32_
 
     h->queue = device->device.getQueue2(queue_info2);
     h->device = device;
+    // Avoid concurrent submissions on NVIDIA due to driver bug.
+    if (device->vendor_id == VK_VENDOR_ID_NVIDIA) {
+        h->device_submit_mutex = &device->queue_submit_mutex;
+    }
     q->handle = h;
 
     q->cmd_pool.init(device, q.get());

From 002a12ad25503a93501b2e188c360029830a241a Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Sun, 13 Sep 2026 09:18:28 +0300
Subject: [PATCH 05/34] ci : cap test-backend-ops parallel jobs at 2 and add a
 3600s timeout (#28833)

- Clamp the -j parallelism to min(nproc, 2) so a single-core runner
  uses -j 1 and multi-core runners use at most -j 2, instead of
  unconditionally using $(nproc).
- Add a 3600s timeout to both test-backend-ops runs (the high-perf CPU
  path and the default path) so a hung test cannot stall CI indefinitely.
- Note a TODO to reduce the timeout to 1800s in the future.

Assisted-by: pi:llama.cpp/Qwen3.8-27B
---
 ci/run.sh | 11 ++++++++---
 1 file changed, 8 insertions(+), 3 deletions(-)

diff --git a/ci/run.sh b/ci/run.sh
index 1ceb19fd5..e510b8c45 100755
--- a/ci/run.sh
+++ b/ci/run.sh
@@ -775,7 +775,11 @@ function gg_run_test_backend_ops {
 
     set -e
 
-    local args_extra="-j $(nproc)"
+    local n_jobs=$(nproc)
+    if [ "${n_jobs}" -gt 2 ]; then
+        n_jobs=2
+    fi
+    local args_extra="-j ${n_jobs}"
 
     # TODO: fix multi-threaded for ROCm
     #       https://github.com/ggml-org/llama.cpp/actions/runs/34576278519/job/103297889044?pr=28740#step:3:4865
@@ -789,10 +793,11 @@ function gg_run_test_backend_ops {
         args_extra=""
     fi
 
+    # TODO: reduce the test-backend-ops timeout to 1800s
     if [ ! -z ${GG_BUILD_HIGH_PERF} ]; then
-        (time ./bin/test-backend-ops ${args_extra} -b CPU) 2>&1 | tee -a $OUT/${ci}-test-backend-ops.log
+        (time timeout 3600 ./bin/test-backend-ops ${args_extra} -b CPU) 2>&1 | tee -a $OUT/${ci}-test-backend-ops.log
     else
-        (time ./bin/test-backend-ops ${args_extra}       ) 2>&1 | tee -a $OUT/${ci}-test-backend-ops.log
+        (time timeout 3600 ./bin/test-backend-ops ${args_extra}       ) 2>&1 | tee -a $OUT/${ci}-test-backend-ops.log
     fi
 
     set +e

From 37b3a9e0ccba261d1cc245a971deae0b18c201ab Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Sigbj=C3=B8rn=20Skj=C3=A6ret?=
 <sigbjorn.skjaeret@huggingface.co>
Date: Sun, 13 Sep 2026 09:41:27 +0200
Subject: [PATCH 06/34] ci : remove leftover command (#28839)

---
 .github/workflows/server-sanitize.yml | 2 --
 1 file changed, 2 deletions(-)

diff --git a/.github/workflows/server-sanitize.yml b/.github/workflows/server-sanitize.yml
index 11237cf51..43746e91e 100644
--- a/.github/workflows/server-sanitize.yml
+++ b/.github/workflows/server-sanitize.yml
@@ -116,7 +116,6 @@ jobs:
         run: |
           source .venv/bin/activate
           cd tools/server/tests
-          export ${{ matrix.extra_args }}
           PYTEST_WORKERS=1 ./tests.sh
 
       - name: Slow tests
@@ -125,5 +124,4 @@ jobs:
         run: |
           source .venv/bin/activate
           cd tools/server/tests
-          export ${{ matrix.extra_args }}
           PYTEST_WORKERS=1 SLOW_TESTS=1 ./tests.sh

From 4a89937354190cef5a97baf8eeb17336105eb72d Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Sun, 13 Sep 2026 13:05:28 +0300
Subject: [PATCH 07/34] tests : reduce FA test sizes (#28842)

---
 tests/test-backend-ops.cpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index b63b3773e..d02297cf5 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -10678,9 +10678,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_flash_attn_ext(256, 256, 2, {16, 1}, 16384,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0));
 
     // MLA shape: the V cache is a sub-view of the K cache, with quantized KV
-    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1},  113,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
-    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 1024,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
-    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {20, 1}, 1024,  64, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
+    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {8, 1},  113,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
+    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {8, 1}, 1024,   1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
+    test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {8, 1}, 1024,  64, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q8_0, {0, 1, 2, 3}, true, true));
 
     // Sparse mask hint: supported decode/prefill layouts and dense fallbacks.
     test_cases.emplace_back(new test_flash_attn_ext(512, 512, 1, { 8, 1}, 4096, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, false,  512));

From bc52a12b38941b0a690ade65fbc5749715224e30 Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Sun, 13 Sep 2026 18:13:32 +0300
Subject: [PATCH 08/34] pi : prefer PI_MODEL_NAME env var for model disclosure
 (#28853)

Assisted-by: pi:llama.cpp/Qwen3.8-27B
---
 .pi/gg/SYSTEM.md | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/.pi/gg/SYSTEM.md b/.pi/gg/SYSTEM.md
index 47883081c..369b87bcd 100644
--- a/.pi/gg/SYSTEM.md
+++ b/.pi/gg/SYSTEM.md
@@ -6,6 +6,7 @@ General:
 - PR and commit titles format: `<module> : <title>`. Lookup recents for examples
 - Don't try to build or run the code unless you are explicitly asked to do so
 - Use the `gh` CLI tool when querying PRs, issues, or other GitHub resources
+- When [MODEL] is needed, first try to get it from the `PI_MODEL_NAME` env var before asking the user
 
 Coding:
 - When in doubt, always refer to the CONTRIBUTING.md file of the project
@@ -20,7 +21,7 @@ Pull requests (PRs):
 - Don't explicitly wrap lines in the PR description (each paragraph and bullet is a single line)
 - When creating a pull request, look for the repository's PR template and follow it
 - For the AI usage disclosure section, write "YES. pi:llama.cpp/[MODEL]"
-- Ask the user to tell you what model was used and write it in place of [MODEL]
+- If `PI_MODEL_NAME` env var is not set, ask the user to tell you what model was used and write it in place of [MODEL]
 - Always create the pull requests in draft mode
 
 Commits:

From c95f8e47b8d9796659ed513957b5077888ad9acb Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Sun, 13 Sep 2026 18:16:45 +0300
Subject: [PATCH 09/34] ci : run editorconfig and code-style checks on
 ubuntu-slim (#28854)

Move the EditorConfig Checker and Code Style Checker workflows from the
`[self-hosted, fast]` runners to `ubuntu-slim`, which is an established
runner label in the repo.

Assisted-by: pi:llama.cpp/Qwen3.8-27B
---
 .github/workflows/code-style.yml   | 2 +-
 .github/workflows/editorconfig.yml | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/.github/workflows/code-style.yml b/.github/workflows/code-style.yml
index 50b598b84..c88396c0a 100644
--- a/.github/workflows/code-style.yml
+++ b/.github/workflows/code-style.yml
@@ -15,7 +15,7 @@ concurrency:
 
 jobs:
   model-naming:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim
     steps:
       - uses: actions/checkout@v6
       - name: Check model naming conventions
diff --git a/.github/workflows/editorconfig.yml b/.github/workflows/editorconfig.yml
index 59159cd41..53f6a0ccf 100644
--- a/.github/workflows/editorconfig.yml
+++ b/.github/workflows/editorconfig.yml
@@ -15,7 +15,7 @@ concurrency:
 
 jobs:
   editorconfig:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim
     steps:
       - uses: actions/checkout@v6
       - uses: editorconfig-checker/action-editorconfig-checker@840e866d93b8e032123c23bac69dece044d4d84c # v2.2.0

From b6b003d2cb29647d968302eb2db8da6f66303b3e Mon Sep 17 00:00:00 2001
From: Neo Zhang <zhang.jianyu@outlook.com>
Date: Sun, 13 Sep 2026 23:31:34 +0800
Subject: [PATCH 10/34] sycl : Fix get mem error (#28227)

* fix for unsupport zes API

* optimize the code

* adjust the log level

* rm unused head files

* Update docs/backend/SYCL.md

Co-authored-by: Titaniumtown <titaniumtown@proton.me>

* fix the error to detect level zero SDK/dev package, stop build after detect the error

* update the message

* fix the build error when missed to install level zero dev package

* rm GGML_SYCL_DEV_DEBUG, mv read env vars in all entry functions

---------

Co-authored-by: Neo Zhang Jianyu <jianyu.zhang@intel.com>
Co-authored-by: Titaniumtown <titaniumtown@proton.me>
Co-authored-by: Neo Zhang <NA>
---
 docs/backend/SYCL.md              |  5 +-
 ggml/src/ggml-sycl/CMakeLists.txt | 14 +++--
 ggml/src/ggml-sycl/base.hpp       |  7 +++
 ggml/src/ggml-sycl/ggml-sycl.cpp  | 82 +++++++++++++++++++-------
 ggml/src/ggml-sycl/mem.cpp        | 95 ++++++++++++++-----------------
 5 files changed, 123 insertions(+), 80 deletions(-)

diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md
index 4a640e442..91b209741 100644
--- a/docs/backend/SYCL.md
+++ b/docs/backend/SYCL.md
@@ -790,14 +790,15 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm
 
 | Name              | Value            | Function                                                                                                                  |
 |-------------------|------------------|---------------------------------------------------------------------------------------------------------------------------|
-| GGML_SYCL_DEBUG   | 0 (default) or 1 | Enable log function by macro: GGML_SYCL_DEBUG                                                                             |
+| GGML_SYCL_DEBUG   | 0 (default) or 1 | Enable log function: GGML_SYCL_DEBUG() for common debug. |
+| GGML_SYCL_DEV_DEBUG   | 0 (default) or 1 | Enable log function: GGML_SYCL_DEV_DEBUG() for developmental purposes by replacing GGML_SYCL_DEBUG() in special codes. Restore to GGML_SYCL_DEBUG() before committing code.|
 | GGML_SYCL_DEV2DEV_MEMCPY | 0 (default), 1, 2 | Choose the method of dev2dev memory copy.<br>Value: <br>*  0: SYCL API (default), only support dGPUs.<br>* 1: L0 API -- Better performance, only support dGPUs, found to lead to abnormal crash in some case. <br>* 2: Host Forward -- Most stable method for all cases (including iGPU + dGPU*N), but with lower performance (-2% to -5%).<br>SYCL & L0 API are easy to be impacted by Intel GPU driver issue. When you meet the garbled output or crash issues in multiple GPUs case, try with this debug flag to work around or check the issue.|
 | GGML_SYCL_ENABLE_FLASH_ATTN | 1 (default) or 0| Enable Flash-Attention. It can reduce memory usage. The performance impact depends on the LLM.|
 | GGML_SYCL_ENABLE_OPT | 0 or 1 (default)| Enable optimize features for Intel GPUs. (Recommended to 0 for Intel devices older than Gen 10) |
 | GGML_SYCL_ENABLE_GRAPH | 0 (default) or 1 | Enable running computations through SYCL Graphs feature. Disabled by default because SYCL Graph is still on development, no better performance. |
 | GGML_SYCL_ENABLE_HOST_PINNED_MEM | 0 or 1 (default) | Enable host pinned memory to speed up copy data from host to device. When disable it, host memory will common malloc() on CPU. Disable it when use `--load-model mlock`.|
 | GGML_SYCL_HOST_PINNED_MEM_2G | 0 (default) or 1 | Limit the max memory allocation to be no more than 2GB when enable host pinned memory. USM allocations above 2 GiB take the relaxed/large-allocation path, which serializes H2D copies with compute and prevents copy/compute overlap. It will impact the startup time. Need more test. Depend on `GGML_SYCL_ENABLE_HOST_PINNED_MEM=1`.|
-| GGML_SYCL_GET_MEM_API | 0 (default) or 1  | Set to get memory info (free, total) by Level Zero or SYCL API:<br>0 - Level Zero API: support more GPUs, only run on Level Zero running time. When there is an error, fallback to call SYCL API. Depend on GGML_SYCL_SUPPORT_LEVEL_ZERO_API.<br>1 - SYCL API: legacy, support more running time, it can't get the free size of some GPUs (like Arc770). In such case, return total size for free size.|
+| GGML_SYCL_GET_MEM_API | 0 (default) or 1  | Set to get memory info (free, total) by Level Zero or SYCL API:<br>0 - Level Zero API: support more GPUs, only run on Level Zero running time. When there is an error, fallback to call SYCL API. Depend on GGML_SYCL_SUPPORT_LEVEL_ZERO_API.<br>1 - SYCL API: legacy, support more running time, it can't get the free size of some GPUs (like Arc770). In such case, return the free size as value of total size.|
 | GGML_SYCL_USE_LEVEL_ZERO_API | 1 (default) or 0 | Use Level Zero API for device memory allocation instead of SYCL. Reduces system RAM usage on Intel dGPUs by avoiding DMA-buf/TTM host memory staging. Requires GGML_SYCL_SUPPORT_LEVEL_ZERO_API=ON at build time. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).|
 | GGML_SYCL_ENABLE_DNN | 0 or 1 (default)| Enable running computations through oneDNN and always use oneMKL. |
 | GGML_SYCL_FA_ONEDNN | 1 (default) or 0 | Enable the oneDNN fused SDPA (flash-attention) path on supported GPUs. Set to 0 to always use the native SYCL flash-attention kernel. |
diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt
index a8d9c0d80..d2196f74d 100644
--- a/ggml/src/ggml-sycl/CMakeLists.txt
+++ b/ggml/src/ggml-sycl/CMakeLists.txt
@@ -110,15 +110,21 @@ if (GGML_SYCL_SUPPORT_LEVEL_ZERO_API)
     # Link against Level Zero loader for direct device memory allocation.
     # Avoids sycl::malloc_device triggering DMA-buf/TTM system RAM staging
     # in the xe kernel driver during multi-GPU inference.
-    find_path(LEVEL_ZERO_INCLUDE_DIR level_zero/ze_api.h HINTS ${ONEAPI_ROOT}/include ${LEVEL_ZERO_V1_SDK_PATH}/include)
+    find_path(LEVEL_ZERO_DEV_INCLUDE_DIR level_zero/ze_api.h HINTS ${ONEAPI_ROOT}/include ${LEVEL_ZERO_V1_SDK_PATH}/include)
     find_library(ZE_LOADER_LIB ze_loader HINTS ${ONEAPI_ROOT}/lib ${LEVEL_ZERO_V1_SDK_LIB_PATH} ENV LD_LIBRARY_PATH)
-    if(ZE_LOADER_LIB AND LEVEL_ZERO_INCLUDE_DIR)
+    if(ZE_LOADER_LIB AND LEVEL_ZERO_DEV_INCLUDE_DIR)
         target_link_libraries(ggml-sycl PRIVATE ${ZE_LOADER_LIB})
         target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_SUPPORT_LEVEL_ZERO_API)
         message(STATUS "Level Zero loader found: ${ZE_LOADER_LIB}")
-        message(STATUS "Level Zero headers found: ${LEVEL_ZERO_INCLUDE_DIR}")
+        message(STATUS "Level Zero development headers found: ${LEVEL_ZERO_DEV_INCLUDE_DIR}")
     else()
-        message(WARNING "Level Zero loader or headers not found, Level Zero support disabled")
+        message(WARNING "Level Zero loader or development headers not found, "
+                        "Level Zero API support disabled. "
+                        "Please install the Level Zero SDK/development package "
+                        "to support Level Zero API features. "
+                        "Level Zero API is not mandatory for SYCL backend, "
+                        "but it is required by the special features for better "
+                        "function & performance on Intel GPUs.")
     endif()
 endif()
 
diff --git a/ggml/src/ggml-sycl/base.hpp b/ggml/src/ggml-sycl/base.hpp
index 3afd57ccb..fe96c4ab8 100644
--- a/ggml/src/ggml-sycl/base.hpp
+++ b/ggml/src/ggml-sycl/base.hpp
@@ -17,6 +17,7 @@
 #include <cstdio>
 
 extern int g_ggml_sycl_debug;
+extern int g_ggml_sycl_dev_debug;
 
 #if defined(__clang__) && __has_builtin(__builtin_expect)
 // Hint the optimizer to pipeline the more likely following instruction in branches
@@ -33,4 +34,10 @@ extern int g_ggml_sycl_debug;
             fprintf(stderr, __VA_ARGS__); \
     } while (0)
 
+#define GGML_SYCL_DEV_DEBUG(...)          \
+    do {                                  \
+        if (UNLIKELY(g_ggml_sycl_dev_debug))  \
+            fprintf(stderr, __VA_ARGS__); \
+    } while (0)
+
 #endif  // GGML_SYCL_BASE_HPP
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index f225682f2..1eb82ad5a 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -91,6 +91,7 @@
 
 static bool g_sycl_loaded = false;
 int g_ggml_sycl_debug = 0;
+int g_ggml_sycl_dev_debug = 0;
 int g_ggml_sycl_enable_optimize = 1;
 int g_ggml_sycl_enable_graph = 0;
 int g_ggml_sycl_enable_dnn = 1;
@@ -113,8 +114,8 @@ int g_ggml_sycl_enable_host_pinned_mem = 1;
 int g_ggml_sycl_host_pinned_mem_2g = 0;
 int g_ggml_sycl_get_mem_api = MEMORY_API_TYPE_LEVEL_ZERO;
 
-
 static ggml_sycl_device_info ggml_sycl_init() {
+    GGML_SYCL_DEBUG("[SYCL] call ggml_sycl_init\n");
     ggml_sycl_device_info info = {};
 
     // Do not hard crash when there exists no SYCL devices.
@@ -205,12 +206,9 @@ static ggml_sycl_device_info ggml_sycl_init() {
     }
 
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
-    // Large buffers can be allocated before ggml_check_sycl() initializes other
-    // g_ggml_sycl_enable_* globals, so initialize this one as early as we can.
+    //update g_ggml_sycl_use_level_zero_api according to the device support
     g_ggml_sycl_use_level_zero_api =
-        info.ext_oneapi_level_zero && ggml_sycl_get_env("GGML_SYCL_USE_LEVEL_ZERO_API", 1);
-#else
-    g_ggml_sycl_use_level_zero_api = 0;
+        info.ext_oneapi_level_zero && g_ggml_sycl_use_level_zero_api;
 #endif
 
     return info;
@@ -314,23 +312,40 @@ static const char* dev2dev_int2str(int dev2dev) {
 * It's the first internal function to be called by them in SYCL backend.
 * This function is used to do initialize work for the SYCL backend and set the global variables.
 */
+#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+static ze_result_t init_zes() {
+    ze_result_t res = zesInit(0);
+    if (res != ZE_RESULT_SUCCESS) {
+        GGML_SYCL_DEBUG("Warning: [%s] zesInit failed with code %d. Sysman free-memory query be unavailable.\n",
+                            __func__, (int) res);
+    }
+    return res;
+}
+
+ze_result_t get_zes_init_res() {
+    static ze_result_t zes_init_res = init_zes();
+    GGML_SYCL_DEBUG("[SYCL] call %s: zesInit result: %d\n", __func__, (int) zes_init_res);
+    return zes_init_res;
+}
+#endif
+
 void initialize_sycl_begining() {
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
-    ze_result_t zes_init = zesInit(0);
-    if (zes_init != ZE_RESULT_SUCCESS) {
-        std::cerr << "Warning: zesInit failed [ggml_check_sycl] with code " << static_cast<int>(zes_init)
-                            << ". Sysman free-memory query may be unavailable.\n";
-    }
+    //must be called in initialization stage, before any other Level Zero API calls
+    GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__);
+    get_zes_init_res();
 #endif
 }
 
 static void ggml_check_sycl() try {
+    GGML_SYCL_DEBUG("[SYCL] ggml_check_sycl()\n");
     static bool initialized = false;
 
     if (!initialized) {
         initialize_sycl_begining();
 
         g_ggml_sycl_debug = ggml_sycl_get_env("GGML_SYCL_DEBUG", 0);
+        g_ggml_sycl_dev_debug = ggml_sycl_get_env("GGML_SYCL_DEV_DEBUG", 0);
         g_ggml_sycl_enable_optimize = ggml_sycl_get_env("GGML_SYCL_ENABLE_OPT", 1);
         g_ggml_sycl_enable_graph = ggml_sycl_get_env("GGML_SYCL_ENABLE_GRAPH", 0);
         g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1);
@@ -344,9 +359,13 @@ static void ggml_check_sycl() try {
         g_ggml_sycl_enable_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_ESIMD", 1);
         g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
 
+#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+        g_ggml_sycl_use_level_zero_api = ggml_sycl_get_env("GGML_SYCL_USE_LEVEL_ZERO_API", 1);
+#else
+        g_ggml_sycl_use_level_zero_api = 0;
+#endif
         g_ggml_sycl_dev2dev_memcpy = ggml_sycl_get_env("GGML_SYCL_DEV2DEV_MEMCPY", DEV2DEV_MEMCPY_SYCL);
         g_ggml_sycl_get_mem_api = ggml_sycl_get_env("GGML_SYCL_GET_MEM_API", MEMORY_API_TYPE_LEVEL_ZERO);
-
         if (g_ggml_sycl_use_level_zero_api == 0) {
             g_ggml_sycl_dev2dev_memcpy = DEV2DEV_MEMCPY_SYCL;
             g_ggml_sycl_get_mem_api = MEMORY_API_TYPE_SYCL;
@@ -405,6 +424,7 @@ static void ggml_check_sycl() try {
 
         GGML_LOG_INFO("Running with Environment Variables:\n");
         GGML_LOG_INFO("  GGML_SYCL_DEBUG: %d\n", g_ggml_sycl_debug);
+        GGML_LOG_INFO("  GGML_SYCL_DEV_DEBUG: %d\n", g_ggml_sycl_dev_debug);
 
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
         GGML_LOG_INFO("  GGML_SYCL_DEV2DEV_MEMCPY: %d (%s)\n", g_ggml_sycl_dev2dev_memcpy, dev2dev_int2str(g_ggml_sycl_dev2dev_memcpy));
@@ -945,6 +965,7 @@ inline void * aligned_malloc_host(size_t alignment, size_t size) {
 static ggml_backend_buffer_t
 ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft,
                                            size_t size) try {
+    GGML_SYCL_DEBUG("[SYCL] call %s: size=%zu\n", __func__, size);
     ggml_check_sycl();
 
     ggml_backend_sycl_buffer_type_context * buft_ctx = (ggml_backend_sycl_buffer_type_context *)buft->context;
@@ -1464,10 +1485,11 @@ static ggml_backend_buffer_type_i ggml_backend_sycl_split_buffer_type_interface
 };
 
 ggml_backend_buffer_type_t ggml_backend_sycl_split_buffer_type(const float * tensor_split) {
+    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_split_buffer_type\n");
+
     static std::mutex mutex;
     std::lock_guard<std::mutex> lock(mutex);
 
-    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_split_buffer_type\n");
     ggml_check_sycl();
     // FIXME: this is not thread safe
     static std::map<std::array<float, GGML_SYCL_MAX_DEVICES>, struct ggml_backend_buffer_type> buft_map;
@@ -1520,6 +1542,7 @@ static const char * ggml_backend_sycl_host_buffer_type_name(ggml_backend_buffer_
 
 //host pinned memory
 static void * ggml_backend_sycl_host_malloc(size_t size) {
+    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_host_malloc\n");
     void * ptr = nullptr;
     try {
         ggml_check_sycl();
@@ -5341,8 +5364,8 @@ catch (sycl::exception const &exc) {
 }
 
 static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tensor * dst) try {
+    GGML_SYCL_DEBUG("[SYCL] ggml_sycl_compute_forward: dst=%s, op=%s\n", dst->name, ggml_op_name(dst->op));
     if (!g_sycl_loaded) return false;
-    initialize_sycl_begining();
 
     if (dst->src[0] != nullptr && ggml_backend_buffer_is_sycl_split(dst->src[0]->buffer)) {
         ggml_sycl_set_peer_access(dst->src[1]->ne[1], ctx.device);
@@ -5725,11 +5748,27 @@ catch (sycl::exception const &exc) {
   std::exit(1);
 }
 
+bool sycl_get_mem_info(int device, size_t * free, size_t * total) {
+    GGML_SYCL_DEBUG("[SYCL] [%s] g_ggml_sycl_get_mem_api=%d\n",
+        __func__, g_ggml_sycl_get_mem_api);
+
+    MemoryAPIType mem_api_type = MemoryAPIType::MEMORY_API_TYPE_SYCL;
+
+#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+    mem_api_type = get_zes_init_res() == ZE_RESULT_SUCCESS ?
+        (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL;
+#else
+    mem_api_type = MemoryAPIType::MEMORY_API_TYPE_SYCL;
+#endif
+    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device),
+        *free, *total, mem_api_type);
+    GGML_SYCL_DEBUG("[SYCL] [%s] total = %zu free = %zu\n", __func__, *total, *free);
+    return res;
+}
+
 void ggml_backend_sycl_get_device_memory(int device, size_t * free, size_t * total) try {
     GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_get_device_memory\n");
-    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device), *free, *total,
-                               (MemoryAPIType) g_ggml_sycl_get_mem_api);
-    if (!res) {
+    if (!sycl_get_mem_info(device, free, total)) {
         GGML_ABORT("[%s] failed to get device memory size", __func__);
     }
     ggml_sycl_memtrace_report_device("device memory query", device, *free, *total);
@@ -6177,12 +6216,12 @@ static const char * ggml_backend_sycl_device_get_description(ggml_backend_dev_t
 }
 
 static void ggml_backend_sycl_device_get_memory(ggml_backend_dev_t dev, size_t * free, size_t * total) {
+    GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__);
     ggml_backend_sycl_device_context * ctx = (ggml_backend_sycl_device_context *) dev->context;
-    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(ctx->device), *free, *total,
-                               (MemoryAPIType) g_ggml_sycl_get_mem_api);
-    if (!res) {
+    if (!sycl_get_mem_info(ctx->device, free, total)) {
         GGML_ABORT("[%s] failed to get device memory size", __func__);
     }
+    GGML_SYCL_DEBUG("[SYCL] call %s total %zu free %zu\n", __func__, *total, *free);
     ggml_sycl_memtrace_report_device("device memory query (dev)", ctx->device, *free, *total);
 }
 
@@ -7061,6 +7100,7 @@ static const ggml_backend_reg_i ggml_backend_sycl_reg_interface = {
 // backend registry
 
 ggml_backend_reg_t ggml_backend_sycl_reg() {
+    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_reg\n");
     static ggml_backend_reg reg;
     static bool initialized = false;
 
@@ -7068,7 +7108,7 @@ ggml_backend_reg_t ggml_backend_sycl_reg() {
         static std::mutex mutex;
         std::lock_guard<std::mutex> lock(mutex);
         if (!initialized) {
-            initialize_sycl_begining();
+            ggml_check_sycl();
             ggml_backend_sycl_reg_context * ctx = new ggml_backend_sycl_reg_context;
             const int min_batch_size = getenv("GGML_OP_OFFLOAD_MIN_BATCH") ? atoi(getenv("GGML_OP_OFFLOAD_MIN_BATCH")) : 32;
 
diff --git a/ggml/src/ggml-sycl/mem.cpp b/ggml/src/ggml-sycl/mem.cpp
index 5ec466420..ad5bfe0ff 100644
--- a/ggml/src/ggml-sycl/mem.cpp
+++ b/ggml/src/ggml-sycl/mem.cpp
@@ -6,13 +6,13 @@
 #include <level_zero/zes_api.h>
 #endif
 
-#include "base.hpp"
-#include "mem.hpp"
-
 #include <cstdint>
 #include <iostream>
 #include <vector>
 
+#include "base.hpp"
+#include "mem.hpp"
+
 const char * mem_api_int2str(int mem_api) {
     if (mem_api == MEMORY_API_TYPE_SYCL) {
         return "SYCL API";
@@ -24,7 +24,12 @@ const char * mem_api_int2str(int mem_api) {
 }
 
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+/*
+* Depend on to call zesInit(0) before any other Level Zero API calls, otherwise the Level Zero API calls may fail.
+*/
 bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
+    GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using Level Zero API.\n", __func__);
+
     free_bytes  = 0;
     total_bytes = 0;
 
@@ -37,41 +42,28 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot
 #endif
 
     try {
-        ze_result_t zes_init = zesInit(0);
-        if (zes_init != ZE_RESULT_SUCCESS) {
-            std::cerr << "Warning: zesInit failed with code " << static_cast<int>(zes_init)
-                      << ". Sysman free-memory query may be unavailable.\n";
-        }
 
         if (dev.get_platform().get_backend() != kL0Backend) {
-            GGML_SYCL_DEBUG("Device backend is not Level Zero; falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Device backend is not Level Zero.\n");
             return false;
         }
 
         ze_device_handle_t ze_dev = sycl::get_native<kL0Backend>(dev);
         if (ze_dev == nullptr) {
-            GGML_SYCL_DEBUG("Level Zero device handle is null; falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Level Zero device handle is null.\n");
             return false;
         }
 
         ze_result_t r = zesDeviceEnumMemoryModules(ze_dev, &module_count, nullptr);
         if (r != ZE_RESULT_SUCCESS || module_count == 0) {
-            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules. Falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules.\n");
             return false;
         }
 
         std::vector<zes_mem_handle_t> modules(module_count);
         r = zesDeviceEnumMemoryModules(ze_dev, &module_count, modules.data());
         if (r != ZE_RESULT_SUCCESS || module_count == 0) {
-            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules. Falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules.\n");
             return false;
         }
 
@@ -90,73 +82,70 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot
         }
 
         if (total_bytes == 0) {
-            GGML_SYCL_DEBUG("Level Zero memory query returned zero total bytes. Falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Level Zero memory query returned zero total bytes.\n");
             return false;
         }
-        return true;
+        return total_bytes >= free_bytes;
+
     } catch (const sycl::exception & e) {
         GGML_SYCL_DEBUG("Level Zero memory query failed: %s\n", e.what());
-        total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-        free_bytes  = total_bytes;
         return false;
     }
 }
 #endif
 
 bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
-    GGML_SYCL_DEBUG("[%s]Querying free memory using SYCL API.\n", __func__);
+    GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using SYCL API.\n", __func__);
     total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
 
 #if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105)
     if (dev.has(sycl::aspect::ext_intel_free_memory)) {
         try {
-            GGML_SYCL_DEBUG("Querying free memory using SYCL aspect::ext_intel_free_memory.");
+            GGML_SYCL_DEBUG("Querying free memory using SYCL aspect::ext_intel_free_memory.\n");
             free_bytes = dev.get_info<sycl::ext::intel::info::device::free_memory>();
             return true;
         } catch (const sycl::exception &) {
             GGML_SYCL_DEBUG(
-                "Failed to query free memory using SYCL aspect::ext_intel_free_memory. Using total memory as free "
-                "memory.");
-            free_bytes = total_bytes;
+                "Failed to query free memory using SYCL aspect::ext_intel_free_memory.\n");
             return false;
         }
     } else {
         GGML_SYCL_DEBUG(
-            "Device does not support SYCL aspect::ext_intel_free_memory. Using total memory as free memory.");
-        free_bytes = total_bytes;
+            "Device does not support SYCL aspect::ext_intel_free_memory.\n");
     }
 #else
-    GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105. Using total memory as free memory.");
-    free_bytes = total_bytes;
+    GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105.\n");
 #endif
-    return true;
+    return false;
 }
 
 bool get_memory_size(sycl::device dev, size_t & free_bytes, size_t & total_bytes, MemoryAPIType api_type) {
-    const auto name       = dev.get_info<sycl::info::device::name>();
-    const auto vendor     = dev.get_info<sycl::info::device::vendor>();
-    const auto global_mem = dev.get_info<sycl::info::device::global_mem_size>();
 
-    GGML_SYCL_DEBUG("[%s]GPU Name:          %s\n", __func__, name.c_str());
-    GGML_SYCL_DEBUG("[%s]GPU Vendor:        %s\n", __func__, vendor.c_str());
-    GGML_SYCL_DEBUG("[%s]GPU Global Memory: %zu bytes\n", __func__, static_cast<size_t>(global_mem));
+    GGML_SYCL_DEBUG("[%s]GPU Name:          %s\n", __func__,
+        dev.get_info<sycl::info::device::name>().c_str());
+    GGML_SYCL_DEBUG("[%s]GPU Vendor:        %s\n", __func__,
+        dev.get_info<sycl::info::device::vendor>().c_str());
 
     if (api_type == MEMORY_API_TYPE_LEVEL_ZERO) {
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
-        GGML_SYCL_DEBUG("[%s]Querying free memory using Level Zero API.\n", __func__);
-        if (!query_free_memory_by_ze(dev, free_bytes, total_bytes)) {
-            //fallback to SYCL API if Level Zero API fails
-            GGML_SYCL_DEBUG("[%s]Falling back to SYCL API for memory query.\n", __func__);
-            return get_memory_size_by_sycl_api(dev, free_bytes, total_bytes);
+        GGML_SYCL_DEBUG("[%s] Querying free memory using Level Zero API.\n", __func__);
+        if (query_free_memory_by_ze(dev, free_bytes, total_bytes)) {
+            return true;
         }
-        return true;
-#else
-        GGML_SYCL_DEBUG("[%s]Level Zero API support is not enabled. Please enable it to use this feature.\n", __func__);
-        return false;
+        //fallback to SYCL API if Level Zero API fails
+        GGML_SYCL_DEBUG("[%s] Falling back to SYCL API for memory query.\n", __func__);
 #endif
-    } else {  //MEMORY_API_TYPE_SYCL
-        return get_memory_size_by_sycl_api(dev, free_bytes, total_bytes);
     }
+
+    //MEMORY_API_TYPE_SYCL
+    if(get_memory_size_by_sycl_api(dev, free_bytes, total_bytes)){
+        return true;
+    }
+
+    //Todo, fallback to other methods to get free memory size, such as using OS-specific APIs (e.g., /proc/meminfo on Linux, GlobalMemoryStatusEx on Windows, etc.)
+    GGML_SYCL_DEBUG(
+        "[%s] Can't get free mem size by Level Zero and SYCL API. Using total memory as free memory.\n", __func__);
+    free_bytes = total_bytes;
+
+    return true;
 }

From 243a3082d437383f778c5671191a5c5880563855 Mon Sep 17 00:00:00 2001
From: Michael Taylor <162068037+mctylr-gh@users.noreply.github.com>
Date: Sun, 13 Sep 2026 13:50:46 -0300
Subject: [PATCH 11/34] tests : fix typo in test-quant-type-selection for
 nemotron 3 nano (#28835)

Corrects a typo in `tests/test-quant-type-selection` for the
Nvidia Nemotron 3 Nano 30B A3B model, which was referred to as
*nvidia-nemotron-nano-3-30b-a3b*.

The error made the test skip that test case, rather than failing
the test.

[no release]
---
 ...o-3-30b-a3b.schema => nvidia-nemotron-3-nano-30b-a3b.schema} | 0
 tests/test-quant-type-selection.cpp                             | 2 +-
 2 files changed, 1 insertion(+), 1 deletion(-)
 rename tests/snapshots/{nemotron-nano-3-30b-a3b.schema => nvidia-nemotron-3-nano-30b-a3b.schema} (100%)

diff --git a/tests/snapshots/nemotron-nano-3-30b-a3b.schema b/tests/snapshots/nvidia-nemotron-3-nano-30b-a3b.schema
similarity index 100%
rename from tests/snapshots/nemotron-nano-3-30b-a3b.schema
rename to tests/snapshots/nvidia-nemotron-3-nano-30b-a3b.schema
diff --git a/tests/test-quant-type-selection.cpp b/tests/test-quant-type-selection.cpp
index 9a5f5e53e..1696ec164 100644
--- a/tests/test-quant-type-selection.cpp
+++ b/tests/test-quant-type-selection.cpp
@@ -221,7 +221,7 @@ static const remote_model_spec model_specs[] = {
     { "ggml-org/Step-3.5-Flash-GGUF",                 "Q4_K"   },
     { "ggml-org/Qwen3-Coder-Next-GGUF",               "Q8_0"   },
     { "ggml-org/Qwen3-14B-GGUF",                      "Q8_0"   },
-    { "ggml-org/NVIDIA-Nemotron-Nano-3-30B-A3B-GGUF", "Q8_0"   },
+    { "ggml-org/NVIDIA-Nemotron-3-Nano-30B-A3B-GGUF", "Q8_0"   },
     { "ggml-org/gpt-oss-120b-GGUF",                   "mxfp4"  },
     { "ggml-org/gemma-3-4b-it-GGUF",                  "Q8_0"   },
     { "bartowski/Meta-Llama-3.1-70B-Instruct-GGUF",   "Q4_K_M" },

From 6978052985cb094da528c829b9f57858ca111025 Mon Sep 17 00:00:00 2001
From: Bernard Ladenthin <bernard.ladenthin@gmail.com>
Date: Sun, 13 Sep 2026 19:18:53 +0200
Subject: [PATCH 12/34] ggml-cpu(s390x): guard VXE-only repack helpers (#28775)

---
 ggml/src/ggml-cpu/arch/s390/repack.cpp | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/ggml/src/ggml-cpu/arch/s390/repack.cpp b/ggml/src/ggml-cpu/arch/s390/repack.cpp
index 3990a6b04..abf3433ad 100644
--- a/ggml/src/ggml-cpu/arch/s390/repack.cpp
+++ b/ggml/src/ggml-cpu/arch/s390/repack.cpp
@@ -70,6 +70,7 @@ void ggml_quantize_mat_q8_0_4x4(const float * GGML_RESTRICT x, void * GGML_RESTR
 #endif
 }
 
+#if defined(__VXE__) || defined(__VXE2__)
 static inline int16x8_t vxe_dot_acc(const int8x16_t v_x, const int8x16_t v_y, const int16x8_t v_acc) {
     return vec_meadd(v_x, v_y, vec_moadd(v_x, v_y, v_acc));
 }
@@ -84,6 +85,7 @@ static inline int32x4_t vxe_fold(const int16x8_t v_sumi) {
     const int16x8_t v_ones = vec_splats((int16_t)1);
     return vec_add(vec_mule(v_sumi, v_ones), vec_mulo(v_sumi, v_ones));
 }
+#endif
 
 void ggml_gemv_q4_0_4x4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int nr, int nc) {
     const int qk = QK8_0;

From e49d2c27605ec3c5b299b9583fa8dc0442739e18 Mon Sep 17 00:00:00 2001
From: Yaniss Amazouz <yaniss91600@gmail.com>
Date: Sun, 13 Sep 2026 20:20:50 +0300
Subject: [PATCH 13/34] models : guard the expert FFN size fallback in
 nemotron-h against a zero divisor (#28779)

The NextN/MTP tail loop derives the expert FFN size as n_ff/n_expert_used
when expert_feed_forward_length gives nothing for the layer. Both values come
from per-layer arrays that legitimately hold 0 on layers that are not MoE, so
a checkpoint whose predict layers hold 0 in both divides by zero and dies with
SIGFPE at load time, with no error message. Report the malformed metadata
instead.
---
 src/models/nemotron-h.cpp | 10 ++++++++--
 1 file changed, 8 insertions(+), 2 deletions(-)

diff --git a/src/models/nemotron-h.cpp b/src/models/nemotron-h.cpp
index d2c48f125..ff8784d18 100644
--- a/src/models/nemotron-h.cpp
+++ b/src/models/nemotron-h.cpp
@@ -145,8 +145,14 @@ void llama_model_nemotron_h::load_arch_tensors(llama_model_loader & ml) {
         const int64_t n_head_i       = hparams.n_head(i);
         const int64_t n_embd_k_gqa_i = hparams.n_embd_k_gqa(i);
         const int64_t n_embd_v_gqa_i = hparams.n_embd_v_gqa(i);
-        const int64_t n_ff_exp       = hparams.n_ff_exp(i) ? (int64_t)hparams.n_ff_exp(i) : n_ff / (int64_t)hparams.n_expert_used(i);
-        const int64_t n_ff_shexp     = hparams.n_ff_shexp;
+        const int64_t n_expert_used_i = hparams.n_expert_used(i);
+        const int64_t n_ff_exp_i      = hparams.n_ff_exp(i);
+        if (n_ff_exp_i == 0 && n_expert_used_i == 0) {
+            throw std::runtime_error(format("%s: layer %d declares neither expert_feed_forward_length nor expert_used_count, "
+                                            "cannot determine the expert FFN size", __func__, i));
+        }
+        const int64_t n_ff_exp   = n_ff_exp_i ? n_ff_exp_i : n_ff / n_expert_used_i;
+        const int64_t n_ff_shexp = hparams.n_ff_shexp;
 
         // NextN input-fusion tensors
         layer.nextn.enorm            = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM,            "weight", i), {n_embd}, mtp_flags);

From 5f436dddb440a288ee5611d7d1eca564a6aca9f4 Mon Sep 17 00:00:00 2001
From: fairydreaming <166155368+fairydreaming@users.noreply.github.com>
Date: Sun, 13 Sep 2026 19:24:11 +0200
Subject: [PATCH 14/34] tests : exclude HY_V4 from WebGPU test-llama-archs
 tests (#28855)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Co-authored-by: Stanisław Szymczyk <sszymczy@gmail.com>
---
 tests/test-llama-archs.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp
index 3496f72e4..018ff1f42 100644
--- a/tests/test-llama-archs.cpp
+++ b/tests/test-llama-archs.cpp
@@ -563,7 +563,8 @@ static bool arch_supported(const llm_arch arch) {
     }
     // FIXME: these hit scheduler/view-backed-output issues with WebGPU on CI.
 #ifdef GGML_USE_WEBGPU
-    if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_DOTS3NOTE || arch == LLM_ARCH_QWEN4EXP) {
+    if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_DOTS3NOTE || arch == LLM_ARCH_QWEN4EXP ||
+            arch == LLM_ARCH_HY_V4) {
         return false;
     }
 #endif // GGML_USE_WEBGPU

From 7a16a6ce326f69752aafaf11468b3103331e26d9 Mon Sep 17 00:00:00 2001
From: Clint Herron <hanclinto@gmail.com>
Date: Sun, 13 Sep 2026 17:56:46 -0400
Subject: [PATCH 15/34] grammar : coalesce find + insert into a single insert
 and adjust move/copy mechanics (#26885)

1) Combine two consecutive lookups (find + insert) into a single insert-attempt/lookup routine so that we don't per
form two O(log(n)) lookup operations in a row anymore -- we only need to do it once and then see if the insert succeeded.
2) Instead of copying every potential stack (expensive) and then moving it (cheap) to new_stacks when it's a final output state, we switch the order so that we move every potential stack (cheap), and then only copy it (expensive) to new stacks when it's a final output state. There are a LOT of intermediate states that get generated, and unless they become final output states, then all of these expensive intermediate copies are wasted.

Before: lookup -> lookup/insert + copy -> optional move to output
New: lookup/insert + move -> optional copy to output
---
 src/llama-grammar.cpp | 11 ++++++-----
 1 file changed, 6 insertions(+), 5 deletions(-)

diff --git a/src/llama-grammar.cpp b/src/llama-grammar.cpp
index 6aa03c766..deffec8c0 100644
--- a/src/llama-grammar.cpp
+++ b/src/llama-grammar.cpp
@@ -871,17 +871,18 @@ static void llama_grammar_advance_stack(
     std::set<llama_grammar_stack, decltype(stack_cmp)> seen(stack_cmp);
 
     while (!todo.empty()) {
-        llama_grammar_stack curr_stack = std::move(todo.back());
+        llama_grammar_stack curr_stack_candidate = std::move(todo.back());
         todo.pop_back();
 
-        if (seen.find( curr_stack) != seen.end()) {
+        auto [curr_stack_it, inserted] = seen.insert(std::move(curr_stack_candidate));
+        if (!inserted) {
             continue;
         }
-        seen.insert(curr_stack);
+        const llama_grammar_stack & curr_stack = *curr_stack_it;
 
         if (curr_stack.empty()) {
             if (std::find(new_stacks.begin(), new_stacks.end(), curr_stack) == new_stacks.end()) {
-                new_stacks.emplace_back(std::move(curr_stack));
+                new_stacks.emplace_back(curr_stack);
             }
             continue;
         }
@@ -924,7 +925,7 @@ static void llama_grammar_advance_stack(
         case LLAMA_GRETYPE_TOKEN_NOT:
             if (std::find(new_stacks.begin(), new_stacks.end(), curr_stack) == new_stacks.end()) {
                 // only add the stack if it's not a duplicate of one we already have
-                new_stacks.emplace_back(std::move(curr_stack));
+                new_stacks.emplace_back(curr_stack);
             }
             break;
         default:

From ad6c66839af3c5646fba8c6c2e2087a1e4e38948 Mon Sep 17 00:00:00 2001
From: thelittlefireman <5165783+thelittlefireman@users.noreply.github.com>
Date: Mon, 14 Sep 2026 00:05:10 +0200
Subject: [PATCH 16/34] ggml-cuda: fallback to F32 on device without BF16
 hardware acceleration (#28846)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

* ggml-cuda: fallback to F32 on device without BF16 hardware acceleration: (Nvidia >= AMPERE, AMD >= RDNA3 or = CDNA)

* apply logic to NVIDIA as well

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
---
 ggml/src/ggml-cuda/common.cuh   |  6 ++++++
 ggml/src/ggml-cuda/ggml-cuda.cu | 12 ++++++++++--
 2 files changed, 16 insertions(+), 2 deletions(-)

diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
index 7d14ce906..2e78ae4fa 100644
--- a/ggml/src/ggml-cuda/common.cuh
+++ b/ggml/src/ggml-cuda/common.cuh
@@ -329,6 +329,12 @@ static bool fp16_mma_hardware_available(const int cc) {
         (GGML_CUDA_CC_IS_MTHREADS(cc) && cc >= GGML_CUDA_CC_QY2);
 }
 
+// To be used for feature selection of external libraries, e.g. cuBLAS.
+static bool fast_bf16_hardware_available(const int cc) {
+        return (GGML_CUDA_CC_IS_AMD(cc) && (cc >= GGML_CUDA_CC_RDNA3 || GGML_CUDA_CC_IS_CDNA(cc)))
+            || (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE);
+}
+
 static bool bf16_mma_hardware_available(const int cc) {
     return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE) ||
         GGML_CUDA_CC_IS_CDNA(cc) || cc >= GGML_CUDA_CC_RDNA3 ||
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 5ae3b8d22..790553888 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -1620,11 +1620,19 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
 }
 
 static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    const int cc = ggml_cuda_info().devices[ctx.device].cc;
     ggml_type compute_type = src0->type;
     if (ggml_is_quantized(compute_type)) {
-        compute_type = fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
-    } else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc)) {
+        compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
+    } else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc)) {
         compute_type = GGML_TYPE_F32;
+    } else if (compute_type == GGML_TYPE_BF16 && !fast_bf16_hardware_available(cc)) {
+        if (GGML_CUDA_CC_IS_AMD(cc) && src1->ne[1] > 32) {
+            compute_type = GGML_TYPE_F32;
+        }
+        if (GGML_CUDA_CC_IS_NVIDIA(cc) && src1->ne[1] > (cc >= GGML_CUDA_CC_VOLTA ? 8 : 128)) {
+            compute_type = GGML_TYPE_F32;
+        }
     }
     if (dst->op_params[0] == GGML_PREC_F32) {
         compute_type = GGML_TYPE_F32;

From 093a2f86c3e37c54fa3e1f9efb17b304f3433abd Mon Sep 17 00:00:00 2001
From: Daniel Bevenius <daniel.bevenius@gmail.com>
Date: Mon, 14 Sep 2026 05:24:05 +0200
Subject: [PATCH 17/34] common : move llama_n_rs_seq to before llama_decode
 (#28749)

This commit moves the llama_n_rs_seq function call to before the
llama_decode call and returns directly if the check is true, removing
the setting of res and the goto statement.

The motivation for this change is to avoid the llama_decode call if it
is not needed.
---
 common/common.cpp | 11 +++++------
 1 file changed, 5 insertions(+), 6 deletions(-)

diff --git a/common/common.cpp b/common/common.cpp
index d162a3880..d8319cd9a 100644
--- a/common/common.cpp
+++ b/common/common.cpp
@@ -1586,6 +1586,11 @@ common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) {
         return COMMON_CONTEXT_SEQ_RM_TYPE_NO;
     }
 
+    if (llama_n_rs_seq(ctx) > 0) {
+        COM_TRC("%s", "the context supports bounded partial sequence removal\n");
+        return COMMON_CONTEXT_SEQ_RM_TYPE_RS;
+    }
+
     common_context_seq_rm_type res = COMMON_CONTEXT_SEQ_RM_TYPE_PART;
 
     llama_memory_clear(mem, true);
@@ -1602,12 +1607,6 @@ common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) {
         goto done;
     }
 
-    if (llama_n_rs_seq(ctx) > 0) {
-        COM_TRC("%s", "the context supports bounded partial sequence removal\n");
-        res = COMMON_CONTEXT_SEQ_RM_TYPE_RS;
-        goto done;
-    }
-
     // try to remove the last tokens
     if (!llama_memory_seq_rm(mem, 0, 1, -1)) {
         COM_TRC("%s", "the context does not support partial sequence removal\n");

From 661643e43079a4ee6faab4c1895291767b67ea8d Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?=C5=81ukasz=20=C5=9Alusarczyk?=
 <lukasz.slusarczyk@intel.com>
Date: Mon, 14 Sep 2026 08:24:06 +0200
Subject: [PATCH 18/34] sycl : fix oneDNN scratchpad breaking the pool free
 order (#28704)

---
 ggml/src/ggml-sycl/common.hpp | 19 -------------------
 ggml/src/ggml-sycl/gemm.hpp   |  6 ++++--
 2 files changed, 4 insertions(+), 21 deletions(-)

diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
index 355dd442b..dc6cdd3df 100644
--- a/ggml/src/ggml-sycl/common.hpp
+++ b/ggml/src/ggml-sycl/common.hpp
@@ -401,29 +401,10 @@ struct ggml_backend_sycl_context {
     dnnl::stream stream_dnnl() {
         return stream_dnnl(device, 0);
     }
-    dnnl::memory get_scratchpad_mem(const dnnl::memory::desc & scratchpad_md,
-                                    const dnnl::engine & eng, const queue_ptr q) {
-        ggml_sycl_pool_alloc<uint8_t> * pool;
-        auto it = scratchpad_map.find(q);
-        if (it == scratchpad_map.end()) {
-            scratchpad_map[q] = std::make_unique<ggml_sycl_pool_alloc<uint8_t>>(this->pool());
-            pool = scratchpad_map[q].get();
-        } else {
-            pool = it->second.get();
-        }
-
-        size_t scratchpad_size = scratchpad_md.get_size();
-        if (scratchpad_size > pool->actual_size) {
-            pool->realloc(scratchpad_size);
-        }
-        void * mem_ptr = pool->get();
-        return dnnl::memory(scratchpad_md, eng, mem_ptr);
-    }
 #endif
 
     // pool
     std::unique_ptr<ggml_sycl_pool> pools[GGML_SYCL_MAX_DEVICES];
-    std::unordered_map<sycl::queue *, std::unique_ptr<ggml_sycl_pool_alloc<uint8_t>>> scratchpad_map;
 
     std::unique_ptr<ggml_sycl_fattn_kv_buffers> fattn_bufs[GGML_SYCL_MAX_DEVICES];
 
diff --git a/ggml/src/ggml-sycl/gemm.hpp b/ggml/src/ggml-sycl/gemm.hpp
index c202da110..81bc5c2e6 100644
--- a/ggml/src/ggml-sycl/gemm.hpp
+++ b/ggml/src/ggml-sycl/gemm.hpp
@@ -66,8 +66,10 @@ public:
         auto matmul_pd = dnnl::matmul::primitive_desc(eng, a_in_md, b_in_md, c_md, primitive_attr);
         auto c_mem = dnnl::memory(matmul_pd.dst_desc(), eng, c);
 
-        auto scratchpad_md = matmul_pd.scratchpad_desc();
-        auto scratchpad_mem = ctx.get_scratchpad_mem(scratchpad_md, eng, q);
+        const auto scratchpad_md = matmul_pd.scratchpad_desc();
+        ggml_sycl_pool_alloc<uint8_t> scratchpad(ctx.pool());
+        void * scratchpad_ptr = scratchpad_md.get_size() > 0 ? scratchpad.alloc(scratchpad_md.get_size()) : nullptr;
+        auto scratchpad_mem = dnnl::memory(scratchpad_md, eng, scratchpad_ptr);
 
         auto matmul_prim = dnnl::matmul(matmul_pd);
 

From 15d8f2d592622107d5ff3931997ef405e1c56a48 Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Mon, 14 Sep 2026 11:50:43 +0300
Subject: [PATCH 19/34] ci : remove gg_sum summary logic (#28857)

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
---
 ci/run.sh | 183 +++++++-----------------------------------------------
 1 file changed, 23 insertions(+), 160 deletions(-)

diff --git a/ci/run.sh b/ci/run.sh
index e510b8c45..a9f92a065 100755
--- a/ci/run.sh
+++ b/ci/run.sh
@@ -58,8 +58,6 @@ if [ -n "${GG_BUILD_ROCM}" ] && [ -n "${GITHUB_RUN_ID}" ]; then
 fi
 
 rm -f $OUT/*.log
-rm -f $OUT/*.exit
-rm -f $OUT/*.md
 
 sd=`dirname $0`
 cd $sd/../
@@ -211,10 +209,6 @@ function gg_wget {
     cd $cwd
 }
 
-function gg_printf {
-    printf -- "$@" >> $OUT/README.md
-}
-
 function gg_run {
     ci=$1
 
@@ -223,13 +217,10 @@ function gg_run {
 
     gg_run_$ci | tee $OUT/$ci.log
     cur=$?
-    echo "$cur" > $OUT/$ci.exit
 
     set +x
     set +o pipefail
 
-    gg_sum_$ci
-
     ret=$((ret | cur))
 }
 
@@ -255,17 +246,6 @@ function gg_run_ctest_debug {
     set +e
 }
 
-function gg_sum_ctest_debug {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs ctest in debug mode\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-ctest.log)"
-    gg_printf '```\n'
-    gg_printf '\n'
-}
-
 # ctest_release
 
 function gg_run_ctest_release {
@@ -290,16 +270,6 @@ function gg_run_ctest_release {
     set +e
 }
 
-function gg_sum_ctest_release {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs ctest in release mode\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-ctest.log)"
-    gg_printf '```\n'
-}
-
 # test_llama_archs_tensor_split
 
 function gg_run_test_llama_archs_tensor_split {
@@ -324,16 +294,6 @@ function gg_run_test_llama_archs_tensor_split {
     set +e
 }
 
-function gg_sum_test_llama_archs_tensor_split {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs test-llama-archs with 1 to 4 devices\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}.log)"
-    gg_printf '```\n'
-}
-
 # test_llama_archs_models
 
 function gg_run_test_llama_archs_models {
@@ -353,16 +313,6 @@ function gg_run_test_llama_archs_models {
     set +e
 }
 
-function gg_sum_test_llama_archs_models {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Generates the dummy models used by the model-dependent tests\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}.log)"
-    gg_printf '```\n'
-}
-
 # test_scripts
 
 function gg_run_test_scripts {
@@ -376,17 +326,6 @@ function gg_run_test_scripts {
     set +e
 }
 
-function gg_sum_test_scripts {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs test scripts\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-scripts.log)"
-    gg_printf '```\n'
-    gg_printf '\n'
-}
-
 function gg_get_model {
     #local gguf_0="$MNT/models/qwen3/0.6B/ggml-model-f16.gguf"
     local gguf_0="$MNT/models/qwen3/0.6B/ggml-model-q4_0.gguf"
@@ -430,26 +369,6 @@ function gg_run_ctest_with_model_release {
     cd ..
 }
 
-function gg_sum_ctest_with_model_debug {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs ctest with model files in debug mode\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-ctest.log)"
-    gg_printf '```\n'
-}
-
-function gg_sum_ctest_with_model_release {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs ctest with model files in release mode\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-ctest.log)"
-    gg_printf '```\n'
-}
-
 # qwen3_0_6b
 
 function gg_run_qwen3_0_6b {
@@ -554,50 +473,24 @@ function gg_run_qwen3_0_6b {
         return 0
     }
 
-    check_ppl "f16"  "$(cat $OUT/${ci}-tg-f16.log  | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
+    check_ppl "f16"  "$(cat $OUT/${ci}-tg-f16.log  | grep "^\[1\]")"
     if [ -z ${GG_BUILD_NO_BF16} ]; then
-        check_ppl "bf16" "$(cat $OUT/${ci}-tg-bf16.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
+        check_ppl "bf16" "$(cat $OUT/${ci}-tg-bf16.log | grep "^\[1\]")"
     fi
-    check_ppl "q8_0" "$(cat $OUT/${ci}-tg-q8_0.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q4_0" "$(cat $OUT/${ci}-tg-q4_0.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q4_1" "$(cat $OUT/${ci}-tg-q4_1.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q5_0" "$(cat $OUT/${ci}-tg-q5_0.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q5_1" "$(cat $OUT/${ci}-tg-q5_1.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-   #check_ppl "q2_k" "$(cat $OUT/${ci}-tg-q2_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log # note: ppl > 20.0 for this quant and model
-    check_ppl "q3_k" "$(cat $OUT/${ci}-tg-q3_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q4_k" "$(cat $OUT/${ci}-tg-q4_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q5_k" "$(cat $OUT/${ci}-tg-q5_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q6_k" "$(cat $OUT/${ci}-tg-q6_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-
-    cat $OUT/${ci}-imatrix.log | grep "Final" >> $OUT/${ci}-imatrix-sum.log
+    check_ppl "q8_0" "$(cat $OUT/${ci}-tg-q8_0.log | grep "^\[1\]")"
+    check_ppl "q4_0" "$(cat $OUT/${ci}-tg-q4_0.log | grep "^\[1\]")"
+    check_ppl "q4_1" "$(cat $OUT/${ci}-tg-q4_1.log | grep "^\[1\]")"
+    check_ppl "q5_0" "$(cat $OUT/${ci}-tg-q5_0.log | grep "^\[1\]")"
+    check_ppl "q5_1" "$(cat $OUT/${ci}-tg-q5_1.log | grep "^\[1\]")"
+   #check_ppl "q2_k" "$(cat $OUT/${ci}-tg-q2_k.log | grep "^\[1\]")" # note: ppl > 20.0 for this quant and model
+    check_ppl "q3_k" "$(cat $OUT/${ci}-tg-q3_k.log | grep "^\[1\]")"
+    check_ppl "q4_k" "$(cat $OUT/${ci}-tg-q4_k.log | grep "^\[1\]")"
+    check_ppl "q5_k" "$(cat $OUT/${ci}-tg-q5_k.log | grep "^\[1\]")"
+    check_ppl "q6_k" "$(cat $OUT/${ci}-tg-q6_k.log | grep "^\[1\]")"
 
     set +e
 }
 
-function gg_sum_qwen3_0_6b {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Qwen3 0.6B:\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '- perplexity:\n%s\n' "$(cat $OUT/${ci}-ppl.log)"
-    gg_printf '- imatrix:\n```\n%s\n```\n' "$(cat $OUT/${ci}-imatrix-sum.log)"
-    gg_printf '- f16:\n```\n%s\n```\n'  "$(cat $OUT/${ci}-tg-f16.log)"
-    if [ -z ${GG_BUILD_NO_BF16} ]; then
-        gg_printf '- bf16:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-bf16.log)"
-    fi
-    gg_printf '- q8_0:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q8_0.log)"
-    gg_printf '- q4_0:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q4_0.log)"
-    gg_printf '- q4_1:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q4_1.log)"
-    gg_printf '- q5_0:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q5_0.log)"
-    gg_printf '- q5_1:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q5_1.log)"
-    gg_printf '- q2_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q2_k.log)"
-    gg_printf '- q3_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q3_k.log)"
-    gg_printf '- q4_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q4_k.log)"
-    gg_printf '- q5_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q5_k.log)"
-    gg_printf '- q6_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q6_k.log)"
-    gg_printf '- save-load-state: \n```\n%s\n```\n' "$(cat $OUT/${ci}-save-load-state.log)"
-}
-
 # bge-small
 
 function gg_run_embd_bge_small {
@@ -639,15 +532,6 @@ function gg_run_embd_bge_small {
     set +e
 }
 
-function gg_sum_embd_bge_small {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'BGE Small (BERT):\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '- f16: \n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-f16.log)"
-    gg_printf '- q8_0:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q8_0.log)"
-}
-
 # rerank_tiny
 
 function gg_run_rerank_tiny {
@@ -704,66 +588,58 @@ function gg_run_rerank_tiny {
     set +e
 }
 
-function gg_sum_rerank_tiny {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Rerank Tiny (Jina):\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '- f16: \n```\n%s\n```\n' "$(cat $OUT/${ci}-rk-f16.log)"
-}
-
 function gg_check_build_requirements {
     if ! command -v git &> /dev/null; then
-        gg_printf 'git not found, please install\n'
+        echo 'git not found, please install'
         exit 1
     fi
 
     if ! command -v git-lfs &> /dev/null; then
-        gg_printf 'git-lfs not found, please install\n'
+        echo 'git-lfs not found, please install'
         exit 1
     fi
 
     if ! git config --get filter.lfs.clean &> /dev/null; then
-        gg_printf 'git-lfs not initialized, please run `git lfs install`\n'
+        echo 'git-lfs not initialized, please run `git lfs install`'
         exit 1
     fi
 
     if ! command -v wget &> /dev/null; then
-        gg_printf 'wget not found, please install\n'
+        echo 'wget not found, please install'
         exit 1
     fi
 
     if ! command -v python3 &> /dev/null; then
-        gg_printf 'python3 not found, please install\n'
+        echo 'python3 not found, please install'
         exit 1
     fi
 
     if ! command -v pip3 &> /dev/null; then
-        gg_printf 'pip3 not found, please install\n'
+        echo 'pip3 not found, please install'
         exit 1
     fi
 
     if ! python3 -m ensurepip --help &> /dev/null; then
-        gg_printf 'ensurepip not found, please install python3-venv package\n'
+        echo 'ensurepip not found, please install python3-venv package'
         exit 1
     fi
 
     if ! command -v cmake &> /dev/null; then
-        gg_printf 'cmake not found, please install\n'
+        echo 'cmake not found, please install'
         exit 1
     fi
 
     if ! command -v ccache &> /dev/null; then
-        gg_printf 'ccache not found, please consider installing for faster builds\n'
+        echo 'ccache not found, please consider installing for faster builds'
     fi
 
     if ! command -v ctest &> /dev/null; then
-        gg_printf 'ctest not found, please install\n'
+        echo 'ctest not found, please install'
         exit 1
     fi
 
     if ! command -v unzip &> /dev/null; then
-        gg_printf 'unzip not found, please install\n'
+        echo 'unzip not found, please install'
         exit 1
     fi
 }
@@ -803,17 +679,6 @@ function gg_run_test_backend_ops {
     set +e
 }
 
-function gg_sum_test_backend_ops {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs test-backend-ops\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-test-backend-ops.log)"
-    gg_printf '```\n'
-    gg_printf '\n'
-}
-
 ## main
 
 export LLAMA_ARG_LOG_PREFIX=1
@@ -861,6 +726,4 @@ if [ -z ${GG_BUILD_LOW_PERF} ]; then
     test $ret -eq 0 && gg_run ctest_with_model_release
 fi
 
-cat $OUT/README.md
-
 exit $ret

From 89fe24240548456477870b2a627cd8021fea1e39 Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Mon, 14 Sep 2026 11:51:06 +0300
Subject: [PATCH 20/34] ci : trigger self-hosted CI on changes to ci/run.sh
 (#28859)

The workflow's push/pull_request path filters did not include the
ci/run.sh script that all of its jobs execute, so changes to it never
re-triggered the self-hosted CI.

Assisted-by: pi:llama.cpp/Qwen3.8-27B
---
 .github/workflows/build-self-hosted.yml | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/.github/workflows/build-self-hosted.yml b/.github/workflows/build-self-hosted.yml
index c7f992540..2e05988bf 100644
--- a/.github/workflows/build-self-hosted.yml
+++ b/.github/workflows/build-self-hosted.yml
@@ -7,6 +7,7 @@ on:
       - master
     paths: [
       '.github/workflows/build-self-hosted.yml',
+      'ci/run.sh',
       '**/CMakeLists.txt',
       '**/.cmake',
       '**/*.h',
@@ -27,6 +28,7 @@ on:
     types: [opened, synchronize, reopened]
     paths: [
       '.github/workflows/build-self-hosted.yml',
+      'ci/run.sh',
       '**/CMakeLists.txt',
       '**/.cmake',
       '**/*.h',

From 2f539596c6e9a977e91b6bc6344650422c6bc3b0 Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Mon, 14 Sep 2026 13:03:41 +0300
Subject: [PATCH 21/34] ggml-cpu : disable PCH and fix CACHE_LINE_SIZE
 ambiguity to fix heap corruption (#28882)

Disable the ggml-cpu precompiled header and remove the
std::hardware_destructive_interference_size branch from CACHE_LINE_SIZE.

The PCH force-includes ggml-impl.h before ops.h, which pulls in <new>
via <array>/<vector> and defines __cpp_lib_hardware_interference_size.
This makes the C++ kernels use CACHE_LINE_SIZE = 256 (hardware
destructive interference size) while the C work-buffer sizing code in
ggml-cpu.c always uses the fallback 64. The mismatch undersizes the
rope work buffer by (CACHE_LINE_SIZE/4 - 16) * n_threads * 4 bytes,
causing a heap-buffer-overflow that corrupts the heap and later crashes
in ggml_compute_forward_rope_flt.

Disabling the ggml-cpu PCH restores the natural include order so
ops.h is processed before <new>, keeping CACHE_LINE_SIZE consistent.
Removing the std::hardware_destructive_interference_size branch makes
the value deterministic and include-order independent.

ref: https://github.com/ggml-org/llama.cpp/issues/28858

Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp
---
 ggml/src/ggml-cpu/CMakeLists.txt |  6 ------
 ggml/src/ggml-cpu/ops.h          | 17 ++++-------------
 2 files changed, 4 insertions(+), 19 deletions(-)

diff --git a/ggml/src/ggml-cpu/CMakeLists.txt b/ggml/src/ggml-cpu/CMakeLists.txt
index 83088e147..1c7338eea 100644
--- a/ggml/src/ggml-cpu/CMakeLists.txt
+++ b/ggml/src/ggml-cpu/CMakeLists.txt
@@ -675,12 +675,6 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
     target_compile_options(${GGML_CPU_NAME} PRIVATE ${ARCH_FLAGS})
     target_compile_definitions(${GGML_CPU_NAME} PRIVATE ${ARCH_DEFINITIONS})
 
-    if (CMAKE_C_COMPILER_ID STREQUAL "GNU" AND NOT GGML_SYSTEM_ARCH STREQUAL "x86")
-        message(STATUS "Skipping PCH for ${GGML_CPU_NAME}: GCC PCH is only enabled for x86 (arch: ${GGML_SYSTEM_ARCH})")
-    else()
-        target_precompile_headers(${GGML_CPU_NAME} PRIVATE ggml-impl.h)
-    endif()
-
     if (EMSCRIPTEN)
         set_target_properties(${GGML_CPU_NAME} PROPERTIES COMPILE_FLAGS "-msimd128")
     endif()
diff --git a/ggml/src/ggml-cpu/ops.h b/ggml/src/ggml-cpu/ops.h
index ce2b3e870..2728b08b6 100644
--- a/ggml/src/ggml-cpu/ops.h
+++ b/ggml/src/ggml-cpu/ops.h
@@ -5,10 +5,10 @@
 //
 // cache line
 //
-
-#if defined(__cpp_lib_hardware_interference_size)
-#define CACHE_LINE_SIZE std::hardware_destructive_interference_size
-#else
+// TODO: rework CACHE_LINE_SIZE so std::hardware_destructive_interference_size
+// can be used consistently between C and C++ TUs; the previous macro form
+// diverged based on include order and undersized the work buffer.
+// ref: https://github.com/ggml-org/llama.cpp/pull/28882
 #if defined(__POWER9_VECTOR__)
 #define CACHE_LINE_SIZE 128
 #elif defined(__VXE__) || defined(__VXE2__)
@@ -16,17 +16,8 @@
 #else
 #define CACHE_LINE_SIZE 64
 #endif
-#endif
 
-// -Winterference-size was introduced in GCC 12
-#if defined(__cplusplus) && defined(__GNUC__) && !defined(__clang__) && __GNUC__ >= 12
-#pragma GCC diagnostic push
-#pragma GCC diagnostic ignored "-Winterference-size"
-#endif
 static const size_t CACHE_LINE_SIZE_F32 = CACHE_LINE_SIZE/sizeof(float);
-#if defined(__cplusplus) && defined(__GNUC__) && !defined(__clang__) && __GNUC__ >= 12
-#pragma GCC diagnostic pop
-#endif
 
 // Work buffer size for im2col operations in CONV2D
 #define GGML_IM2COL_WORK_SIZE (16 * 1024 * 1024)

From 21f6b0d22c6c87efd42d22485cd8b13f61ad5736 Mon Sep 17 00:00:00 2001
From: cwriter <silvan.niederer@bluewin.ch>
Date: Mon, 14 Sep 2026 13:02:44 +0200
Subject: [PATCH 22/34] sycl: rfc: Use radix select for top_k (#28670)

* sycl: GPU-resident TOP_K for large k, parallelised over the device

The SYCL backend refused GGML_OP_TOP_K above k = 32 and let it fall back to
the CPU, a backend round-trip per call. The limit was not conservatism: the
scan-merge kernels keep (split_block + 1) * k candidate (value, index) pairs
in SLM, so at k = 128 a work-group already needs 132 KB and cannot launch.
qwen4exp's sparse-attention indexer asks for k = 2048 in 12 layers on every
token, so this fired at every context length.

Add a radix select for large k. The k-th largest is found by four
most-significant-first passes over an order-preserving unsigned key: histogram
the digit over the candidate set, walk the buckets from the top, and recurse
into the one where the running count reaches what is still needed. SLM holds
the histogram rather than candidates, so the footprint is independent of k.
A final pass emits every column beating the pivot plus exactly as many
pivot-equal columns as are still missing, so duplicate keys still yield
exactly k distinct indices. Output order is not required and is not paid for:
ggml-cpu/ops.cpp swaps its first two outputs to say so.

The key folds -0.0 onto +0.0 so its equivalence classes match the reference
comparator, under which the two tie. NaN has no defined order in the reference
(its comparator is not a strict weak order there); here +NaN keys above +inf
and -NaN below -inf, which at least makes the result deterministic.

One work-group per row leaves the device idle whenever a graph has fewer rows
than it has cores, which at batch size 1 means one work-group full stop:
qwen4exp tops-k a tensor of shape [n_kv, n_tokens/n_stream, n_stream], so
token generation gives nrows == 1, and the backend sampler reshapes logits to
a single row as well. Measured, ne=[200000,1] and ne=[200000,16] cost 358.0 us
and 363.4 us -- sixteen rows for 1.5% more wall-clock.

So also spread a row over several groups when there are too few rows to cover
the device. Per-pass state moves to global memory and each digit pass becomes
its own launch, since a work-group barrier can no longer span the row. Groups
accumulate in SLM and contribute 256 global atomics each, keeping global
traffic per-group rather than per-element, and the last group of a row -- the
one whose fetch_add returns G-1 -- performs that pass's scan, holding the
launch count at one per digit plus one emit. The group count comes from the
device and is floor-divided by nrows, so a row count that already covers the
device is left whole and pays nothing. Below 64K columns the single-group
kernel finishes inside the cost of the extra launches and stays in charge.

Reading the row's prefix/mask/need through a device-scope atomic_ref costs
more than the sweep it guards: those loads are uncached, so passes 2-4 ran at
49 us against 12 us for pass 1. One lane reads them into SLM and the group
takes them from there -- 208 us -> 44.6 us at ne=[131072,1], k=2048.

The block size now takes the device's max_work_group_size instead of a cap of
512. The cap was never a floor, so a device reporting 512 is unaffected; one
allowing 1024 was being given half its width.

Finally, put the scan-merge gate where the two paths actually cross. That
kernel's cost climbs with k while the radix select's does not; measured over
widths from 2 to 200K columns and row counts from 1 to 8192, radix is ahead
everywhere from k = 8 up and behind at k <= 2, where scan-merge's smaller
fixed cost wins. The short-row corner (ncols=2, nrows=65536, as in bailingmoe2
group selection) is exactly where radix loses at low k, and the gate keeps it
on scan-merge.

Op-level against the CPU-fallback path this replaces, and against the
single-group radix select for the split: 4.98x at ne=[131072,1] k=2048,
6.65x at ne=[151936,1] k=40, 13.35x at k=20, 118x at ne=[65000,16] k=32.
No measured shape regressed. End to end on 3x Arc Pro B60 with
Qwen3.8-Flash-Next UD-IQ4_XS, llama-bench tg64, the parallelisation is worth
5.91 -> 6.05 t/s at d=131072 and a wash at shallower depths. Perplexity over
wikitext-2 is unchanged within noise at both 512 and 81920 context.

test-backend-ops: 525/525 TOP_K (previously every k > 32 case was refused),
880/880 MUL_MAT_ID. Perf coverage added for k > 32 at large widths and for the
short-row corner, neither of which was exercised before.

* move topk-select to topk-radix.{cpp|hpp}

---------

Co-authored-by: cwriter <cwriter@localhost>
---
 ggml/src/ggml-sycl/backend.hpp    |   1 +
 ggml/src/ggml-sycl/ggml-sycl.cpp  |  10 +-
 ggml/src/ggml-sycl/topk-radix.cpp | 531 ++++++++++++++++++++++++++++++
 ggml/src/ggml-sycl/topk-radix.hpp |  24 ++
 tests/test-backend-ops.cpp        |  24 ++
 5 files changed, 587 insertions(+), 3 deletions(-)
 create mode 100644 ggml/src/ggml-sycl/topk-radix.cpp
 create mode 100644 ggml/src/ggml-sycl/topk-radix.hpp

diff --git a/ggml/src/ggml-sycl/backend.hpp b/ggml/src/ggml-sycl/backend.hpp
index 51ab6f930..ab80a2a3b 100644
--- a/ggml/src/ggml-sycl/backend.hpp
+++ b/ggml/src/ggml-sycl/backend.hpp
@@ -44,6 +44,7 @@
 #include "ssm_conv.hpp"
 #include "softmax.hpp"
 #include "topk-moe.hpp"
+#include "topk-radix.hpp"
 #include "tsembd.hpp"
 #include "upscale.hpp"
 #include "wkv.hpp"
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 1eb82ad5a..686a4c76e 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -3121,10 +3121,14 @@ static void ggml_sycl_op_top_k(ggml_backend_sycl_context & ctx, ggml_tensor * ds
     const int64_t ncols = src0->ne[0];
     const int64_t nrows = ggml_nrows(src0);
 
-    GGML_ASSERT(k > 0 && k <= 32);
+    GGML_ASSERT(k > 0);
     GGML_ASSERT(k <= ncols);
 
-    top_k_f32_sycl(ctx, src0_dd, dst_dd, ncols, nrows, k, main_stream);
+    if (k <= SYCL_TOP_K_SCAN_MERGE_MAX_K) {
+        top_k_f32_sycl(ctx, src0_dd, dst_dd, ncols, nrows, k, main_stream);
+    } else {
+        ggml_sycl_top_k_radix(ctx, src0_dd, dst_dd, ncols, nrows, k, main_stream);
+    }
 }
 
 inline void ggml_sycl_op_argmax(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
@@ -6644,7 +6648,7 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
                 op->type == GGML_TYPE_I32 &&
                 src0->type == GGML_TYPE_F32 &&
                 ggml_is_contiguous(src0) &&
-                k > 0 && k <= 32;
+                k > 0 && k <= src0->ne[0];
         }
         case GGML_OP_POOL_2D:
         case GGML_OP_POOL_1D:
diff --git a/ggml/src/ggml-sycl/topk-radix.cpp b/ggml/src/ggml-sycl/topk-radix.cpp
new file mode 100644
index 000000000..8cd0bd2f5
--- /dev/null
+++ b/ggml/src/ggml-sycl/topk-radix.cpp
@@ -0,0 +1,531 @@
+#include "topk-radix.hpp"
+
+#include "common.hpp"
+
+#include <algorithm>
+
+// Large-k top-k by radix select on an order-preserving unsigned key.
+//
+// The k-th largest key of a row is found by four most-significant-first passes over its
+// 8-bit digits: histogram the digit over the candidate set, walk the buckets from the
+// top, and recurse into the bucket where the running count reaches what is still
+// needed. Everything strictly above that bucket is in the top-k. A final pass emits
+// every column whose key beats the pivot, then exactly as many pivot-equal columns as
+// are still missing, so duplicate keys yield exactly k distinct indices.
+//
+// SLM holds only the histogram, so unlike the scan-merge kernels the cost does not grow
+// with k. One work-group owns a row and runs every pass, so a top-k is one launch and
+// needs no pool scratch. The row is re-read once per pass rather than compacted, which
+// keeps the candidate set implicit: (key & mask) == prefix.
+//
+// The output is the set of winning indices in no particular order, which is what the
+// reference op provides (it swaps its first two outputs to say so) and what
+// test-backend-ops compares.
+
+static constexpr int SYCL_TOP_K_RADIX_BITS       = 8;
+static constexpr int SYCL_TOP_K_RADIX_BUCKETS    = 1 << SYCL_TOP_K_RADIX_BITS;
+// Private histogram copies, interleaved per bucket so neighbouring lanes hit
+// neighbouring banks. Lanes of one instruction spread over the copies, which is what
+// bounds the atomic serialisation on tie-heavy rows.
+static constexpr int SYCL_TOP_K_RADIX_HIST_COPIES = 8;
+static constexpr int SYCL_TOP_K_RADIX_HIST_SIZE   = SYCL_TOP_K_RADIX_BUCKETS * SYCL_TOP_K_RADIX_HIST_COPIES;
+// Past the histogram: pivot digit, pivot bucket count, remaining need, then the two
+// emit counters.
+static constexpr int SYCL_TOP_K_RADIX_SLM_WORDS   = SYCL_TOP_K_RADIX_HIST_SIZE + 5;
+
+// Larger float <=> larger key. The reference comparator is a plain float '>', under which
+// -0.0 and +0.0 tie, so -0.0 is folded onto +0.0 first. NaN has no defined order in the
+// reference (its comparator is not a strict weak order on NaN); here a positive NaN keys
+// above +inf and a negative NaN below -inf, which at least makes the result deterministic.
+static inline uint32_t top_k_radix_key(float f) {
+    uint32_t u = sycl::bit_cast<uint32_t>(f);
+    if (u == 0x80000000u) {
+        u = 0u;
+    }
+    return (u & 0x80000000u) ? ~u : (u | 0x80000000u);
+}
+
+static void top_k_radix_select_f32(
+    const float *   src,
+    int32_t *       dst_idx,
+    const int       ncols,
+    const int       k,
+    uint32_t *      slm,
+    const sycl::nd_item<1> & item_ct1
+) {
+    using local_atomic = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed,
+                                          sycl::memory_scope::work_group,
+                                          sycl::access::address_space::local_space>;
+
+    const int tid        = item_ct1.get_local_id(0);
+    const int block_size = item_ct1.get_local_range(0);
+
+    uint32_t * hist     = slm;
+    uint32_t * s_digit  = slm + SYCL_TOP_K_RADIX_HIST_SIZE;
+    uint32_t * s_bucket = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 1;
+    uint32_t * s_need   = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 2;
+    uint32_t * s_cnt_gt = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 3;
+    uint32_t * s_cnt_eq = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 4;
+
+    if (tid == 0) {
+        *s_cnt_gt = 0;
+        *s_cnt_eq = 0;
+    }
+
+    const int copy = tid & (SYCL_TOP_K_RADIX_HIST_COPIES - 1);
+
+    uint32_t prefix = 0;   // digits fixed so far, in place
+    uint32_t mask   = 0;   // which bits of prefix are fixed
+    uint32_t need   = (uint32_t) k;
+
+    for (int shift = 32 - SYCL_TOP_K_RADIX_BITS; shift >= 0; shift -= SYCL_TOP_K_RADIX_BITS) {
+        for (int i = tid; i < SYCL_TOP_K_RADIX_HIST_SIZE; i += block_size) {
+            hist[i] = 0;
+        }
+        item_ct1.barrier(sycl::access::fence_space::local_space);
+
+        for (int col = tid; col < ncols; col += block_size) {
+            const uint32_t key = top_k_radix_key(src[col]);
+            if ((key & mask) == prefix) {
+                const uint32_t bucket = (key >> shift) & (SYCL_TOP_K_RADIX_BUCKETS - 1);
+                local_atomic(hist[bucket * SYCL_TOP_K_RADIX_HIST_COPIES + copy]).fetch_add(1u);
+            }
+        }
+        item_ct1.barrier(sycl::access::fence_space::local_space);
+
+        // Lane t takes bucket 255 - t, so an inclusive scan over lanes counts from the top
+        // bucket downward. The pivot is the unique bucket whose cumulative count first
+        // reaches need; the previous cumulative count is what the higher buckets contribute.
+        uint32_t cnt = 0;
+        if (tid < SYCL_TOP_K_RADIX_BUCKETS) {
+            const uint32_t * h = hist + (SYCL_TOP_K_RADIX_BUCKETS - 1 - tid) * SYCL_TOP_K_RADIX_HIST_COPIES;
+            for (int c = 0; c < SYCL_TOP_K_RADIX_HIST_COPIES; c++) {
+                cnt += h[c];
+            }
+        }
+        const uint32_t incl = sycl::inclusive_scan_over_group(item_ct1.get_group(), cnt, sycl::plus<uint32_t>());
+
+        if (tid < SYCL_TOP_K_RADIX_BUCKETS && incl >= need && incl - cnt < need) {
+            *s_digit = (uint32_t) (SYCL_TOP_K_RADIX_BUCKETS - 1 - tid);
+            *s_bucket = cnt;
+            *s_need   = need - (incl - cnt);
+        }
+        item_ct1.barrier(sycl::access::fence_space::local_space);
+
+        const uint32_t digit      = *s_digit;
+        const uint32_t bucket_cnt = *s_bucket;
+        need   = *s_need;
+        prefix |= digit << shift;
+        mask   |= (uint32_t) (SYCL_TOP_K_RADIX_BUCKETS - 1) << shift;
+
+        // Every candidate in the pivot bucket is wanted: the remaining digits cannot
+        // change the answer, and the masked emit below is exact as it stands.
+        if (bucket_cnt == need) {
+            break;
+        }
+        // The next pass rewrites hist and s_*; the reads above must land first.
+        item_ct1.barrier(sycl::access::fence_space::local_space);
+    }
+
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    // Exactly k - need columns have (key & mask) > prefix; the first need of the pivot-equal
+    // columns fill the tail. Both counters live in SLM since the whole row is this group.
+    const uint32_t base_eq = (uint32_t) k - need;
+
+    for (int col = tid; col < ncols; col += block_size) {
+        const uint32_t kp = top_k_radix_key(src[col]) & mask;
+        if (kp > prefix) {
+            const uint32_t pos = local_atomic(*s_cnt_gt).fetch_add(1u);
+            dst_idx[pos] = col;
+        } else if (kp == prefix) {
+            const uint32_t pos = local_atomic(*s_cnt_eq).fetch_add(1u);
+            if (pos < need) {
+                dst_idx[base_eq + pos] = col;
+            }
+        }
+    }
+}
+
+static void top_k_radix_f32_sycl(
+    ggml_backend_sycl_context & ctx,
+    const float * src,
+    int32_t * dst_indices,
+    const int64_t ncols,
+    const int64_t nrows,
+    const int k,
+    dpct::queue_ptr main_stream
+) {
+    GGML_ASSERT(ncols <= INT32_MAX);
+
+    // One group per row; every pass is a strided sweep of the row, so lanes in flight is the
+    // only lever, and the device's own limit is the answer -- there is nothing here that
+    // wants a smaller group. Must still cover the 256 buckets for the scan step.
+    const int block_size = ggml_sycl_info().max_work_group_sizes[ctx.device];
+    GGML_ASSERT(block_size >= SYCL_TOP_K_RADIX_BUCKETS);
+
+    const sycl::range<1> block_dims(block_size);
+    const sycl::range<1> grid_dims(nrows);
+
+    main_stream->submit([&](sycl::handler &cgh) {
+        sycl::local_accessor<uint32_t, 1> slm(sycl::range<1>(SYCL_TOP_K_RADIX_SLM_WORDS), cgh);
+
+        cgh.parallel_for(
+            sycl::nd_range<1>(grid_dims * block_dims, block_dims),
+            [=](sycl::nd_item<1> item_ct1) {
+                const int row = item_ct1.get_group(0);
+
+                top_k_radix_select_f32(
+                    src + (int64_t) row * ncols, dst_indices + (int64_t) row * k,
+                    (int) ncols, k,
+                    slm.get_multi_ptr<sycl::access::decorated::no>().get(),
+                    item_ct1);
+            });
+    });
+}
+
+// One work-group owns a whole row above, which leaves the device idle whenever a graph
+// has fewer rows than it has cores -- the common case at batch size 1, where the
+// sparse-attention indexer and the backend sampler both top-k a single row. The kernels
+// below spread one row over several groups instead.
+//
+// A digit pass now needs the whole row's histogram before any group can pick the pivot,
+// so the per-pass state moves to global memory and the passes become separate launches:
+// a work-group barrier no longer spans the row. Each group still accumulates into SLM
+// and contributes 256 global atomics at the end, so global traffic is per-group, not
+// per-element. The last group to finish a pass (the one whose fetch_add returns G - 1)
+// does the scan for the row and clears the histogram for the next pass, which keeps the
+// launch count at one per digit rather than two.
+//
+// Running all four digits unconditionally costs nothing in correctness: once a bucket
+// holds exactly the elements still needed, later digits only extend the prefix, and the
+// count of columns above that longer prefix grows by exactly as much as `need` shrinks.
+// The emit below therefore stays exact whatever pass the answer settled on.
+
+static constexpr int SYCL_TOP_K_RADIX_ROW_DONE   = SYCL_TOP_K_RADIX_BUCKETS + 0;
+static constexpr int SYCL_TOP_K_RADIX_ROW_PREFIX = SYCL_TOP_K_RADIX_BUCKETS + 1;
+static constexpr int SYCL_TOP_K_RADIX_ROW_MASK   = SYCL_TOP_K_RADIX_BUCKETS + 2;
+static constexpr int SYCL_TOP_K_RADIX_ROW_NEED   = SYCL_TOP_K_RADIX_BUCKETS + 3;
+static constexpr int SYCL_TOP_K_RADIX_ROW_CNT_GT = SYCL_TOP_K_RADIX_BUCKETS + 4;
+static constexpr int SYCL_TOP_K_RADIX_ROW_CNT_EQ = SYCL_TOP_K_RADIX_BUCKETS + 5;
+static constexpr int SYCL_TOP_K_RADIX_ROW_WORDS  = SYCL_TOP_K_RADIX_BUCKETS + 6;
+
+// How wide the split goes is a property of the device, not of the model: enough groups to
+// cover the cores, and no more. Past that the extra groups add histogram traffic without
+// adding bandwidth (measured on this device: 20 and 40 groups tie, 60 and 160 lose).
+//
+// nsm is max_compute_units / 16, i.e. it counts an Xe core as 16 EUs. That is a core's
+// width on Xe-HPG, but an Xe2 core is 8 XVEs wide, so on Battlemage the field reads half
+// the cores actually present (10 for a 20-core B60). The measured curve is flat from one
+// group per core to two and only falls off at three, so a factor of two covers the device
+// on Xe2 and lands in the flat region on Xe-HPG. It is the one number here that a correct
+// core count would remove; it was tuned on Xe2 and has not been measured on Xe-HPG.
+static constexpr int SYCL_TOP_K_RADIX_GROUPS_PER_NSM = 2;
+// Splitting trades one kernel for five. Below the width at which the single-group kernel
+// runs longer than those four extra launches, it wins on its own; measured break-even on
+// this device sits just under 64K columns.
+static constexpr int SYCL_TOP_K_RADIX_MIN_SPLIT_COLS = 65536;
+// A partition thinner than this cannot keep a group's sweep busy.
+static constexpr int SYCL_TOP_K_RADIX_MIN_PART_COLS  = 4096;
+
+static int top_k_radix_split_groups(const int device, const int64_t ncols, const int64_t nrows) {
+    const int64_t target = (int64_t) SYCL_TOP_K_RADIX_GROUPS_PER_NSM * ggml_sycl_info().devices[device].nsm;
+
+    // One group per row already, so a graph with rows enough to cover the device gains
+    // nothing from splitting and would only pay the extra launches.
+    if (ncols < SYCL_TOP_K_RADIX_MIN_SPLIT_COLS || nrows >= target) {
+        return 1;
+    }
+
+    const int64_t by_rows = target / nrows;   // floor: never overshoot a row that is nearly covered
+    const int64_t by_cols = ncols / SYCL_TOP_K_RADIX_MIN_PART_COLS;
+
+    return (int) std::max<int64_t>(1, std::min(by_rows, by_cols));
+}
+
+using top_k_radix_gatomic = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed,
+                                             sycl::memory_scope::device,
+                                             sycl::access::address_space::global_space>;
+
+static void top_k_radix_split_pass_f32(
+    const float *   src,
+    uint32_t *      state,
+    const int       ncols,
+    const int       k,
+    const int       shift,
+    const bool      first,
+    const int       part,
+    const int       nparts,
+    uint32_t *      slm,
+    const sycl::nd_item<1> & item_ct1
+) {
+    using local_atomic = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed,
+                                          sycl::memory_scope::work_group,
+                                          sycl::access::address_space::local_space>;
+
+    const int tid        = item_ct1.get_local_id(0);
+    const int block_size = item_ct1.get_local_range(0);
+
+    uint32_t * hist   = slm;
+    uint32_t * s_last = slm + SYCL_TOP_K_RADIX_HIST_SIZE;
+    uint32_t * s_row  = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 1;   // prefix, mask, need
+
+    // The previous launch is the barrier that publishes these, so a plain load is enough.
+    // One lane reads them and the group takes them from SLM: a device-scope atomic load
+    // is uncached here, and having every work-item issue three of them off the same
+    // address costs more than the whole sweep below.
+    if (tid == 0) {
+        s_row[0] = first ? 0u : state[SYCL_TOP_K_RADIX_ROW_PREFIX];
+        s_row[1] = first ? 0u : state[SYCL_TOP_K_RADIX_ROW_MASK];
+        s_row[2] = first ? (uint32_t) k : state[SYCL_TOP_K_RADIX_ROW_NEED];
+    }
+
+    for (int i = tid; i < SYCL_TOP_K_RADIX_HIST_SIZE; i += block_size) {
+        hist[i] = 0;
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    const uint32_t prefix = s_row[0];
+    const uint32_t mask   = s_row[1];
+    const uint32_t need   = s_row[2];
+
+    const int copy  = tid & (SYCL_TOP_K_RADIX_HIST_COPIES - 1);
+    const int chunk = (ncols + nparts - 1) / nparts;
+    const int col0  = part * chunk;
+    const int col1  = std::min(ncols, col0 + chunk);
+
+    for (int col = col0 + tid; col < col1; col += block_size) {
+        const uint32_t key = top_k_radix_key(src[col]);
+        if ((key & mask) == prefix) {
+            const uint32_t bucket = (key >> shift) & (SYCL_TOP_K_RADIX_BUCKETS - 1);
+            local_atomic(hist[bucket * SYCL_TOP_K_RADIX_HIST_COPIES + copy]).fetch_add(1u);
+        }
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    // One global atomic per bucket per group, not per element.
+    for (int b = tid; b < SYCL_TOP_K_RADIX_BUCKETS; b += block_size) {
+        uint32_t sum = 0;
+        for (int c = 0; c < SYCL_TOP_K_RADIX_HIST_COPIES; c++) {
+            sum += hist[b * SYCL_TOP_K_RADIX_HIST_COPIES + c];
+        }
+        if (sum) {
+            top_k_radix_gatomic(state[b]).fetch_add(sum);
+        }
+    }
+
+    // Publish this group's bins, then claim the scan if this group is the row's last.
+    // The group-wide barrier flushes the atomics above; only the claiming lane needs the
+    // release, so the device-scope fence is paid once per group rather than per work-item.
+    item_ct1.barrier(sycl::access::fence_space::global_and_local);
+    if (tid == 0) {
+        sycl::atomic_fence(sycl::memory_order::release, sycl::memory_scope::device);
+        sycl::atomic_ref<uint32_t, sycl::memory_order::acq_rel, sycl::memory_scope::device,
+                         sycl::access::address_space::global_space> done(state[SYCL_TOP_K_RADIX_ROW_DONE]);
+        *s_last = (done.fetch_add(1u) == (uint32_t) (nparts - 1)) ? 1u : 0u;
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    if (*s_last == 0u) {
+        return;
+    }
+    sycl::atomic_fence(sycl::memory_order::acquire, sycl::memory_scope::device);
+
+    // Lane t takes bucket 255 - t, so an inclusive scan counts down from the top bucket.
+    uint32_t cnt = 0;
+    if (tid < SYCL_TOP_K_RADIX_BUCKETS) {
+        cnt = top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_BUCKETS - 1 - tid]).load();
+    }
+    const uint32_t incl = sycl::inclusive_scan_over_group(item_ct1.get_group(), cnt, sycl::plus<uint32_t>());
+
+    if (tid < SYCL_TOP_K_RADIX_BUCKETS && incl >= need && incl - cnt < need) {
+        const uint32_t digit = (uint32_t) (SYCL_TOP_K_RADIX_BUCKETS - 1 - tid);
+        top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_PREFIX]).store(prefix | (digit << shift));
+        top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_MASK]).store(
+            mask | ((uint32_t) (SYCL_TOP_K_RADIX_BUCKETS - 1) << shift));
+        top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_NEED]).store(need - (incl - cnt));
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    // Clear for the next pass; the next launch is the barrier that orders this.
+    for (int b = tid; b < SYCL_TOP_K_RADIX_BUCKETS; b += block_size) {
+        top_k_radix_gatomic(state[b]).store(0u);
+    }
+    if (tid == 0) {
+        top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_DONE]).store(0u);
+    }
+}
+
+static void top_k_radix_split_emit_f32(
+    const float *   src,
+    int32_t *       dst_idx,
+    uint32_t *      state,
+    const int       ncols,
+    const int       k,
+    const int       part,
+    const int       nparts,
+    uint32_t *      slm,
+    const sycl::nd_item<1> & item_ct1
+) {
+    using local_atomic = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed,
+                                          sycl::memory_scope::work_group,
+                                          sycl::access::address_space::local_space>;
+
+    const int tid        = item_ct1.get_local_id(0);
+    const int block_size = item_ct1.get_local_range(0);
+
+    uint32_t * s_gt      = slm;
+    uint32_t * s_eq      = slm + 1;
+    uint32_t * s_base_gt = slm + 2;
+    uint32_t * s_base_eq = slm + 3;
+
+    uint32_t * s_row = slm + 4;   // prefix, mask, need
+
+    if (tid == 0) {
+        *s_gt = 0;
+        *s_eq = 0;
+        s_row[0] = state[SYCL_TOP_K_RADIX_ROW_PREFIX];
+        s_row[1] = state[SYCL_TOP_K_RADIX_ROW_MASK];
+        s_row[2] = state[SYCL_TOP_K_RADIX_ROW_NEED];
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    const uint32_t prefix = s_row[0];
+    const uint32_t mask   = s_row[1];
+    const uint32_t need   = s_row[2];
+
+    // Exactly k - need columns beat the pivot; the first need pivot-equal ones fill the tail.
+    const uint32_t base_eq = (uint32_t) k - need;
+
+    const int chunk = (ncols + nparts - 1) / nparts;
+    const int col0  = part * chunk;
+    const int col1  = std::min(ncols, col0 + chunk);
+
+    // Counting first and reserving one range per group keeps the row's two counters out of
+    // the inner loop: a per-element global atomic on a single address serialises the whole
+    // emit, and at k in the thousands that alone outweighs every read the kernel does.
+    for (int col = col0 + tid; col < col1; col += block_size) {
+        const uint32_t kp = top_k_radix_key(src[col]) & mask;
+        if (kp > prefix) {
+            local_atomic(*s_gt).fetch_add(1u);
+        } else if (kp == prefix) {
+            local_atomic(*s_eq).fetch_add(1u);
+        }
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    if (tid == 0) {
+        const uint32_t n_gt = *s_gt;
+        const uint32_t n_eq = *s_eq;
+        *s_base_gt = n_gt ? top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_CNT_GT]).fetch_add(n_gt) : 0u;
+        *s_base_eq = n_eq ? top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_CNT_EQ]).fetch_add(n_eq) : 0u;
+        *s_gt = 0;
+        *s_eq = 0;
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    const uint32_t base_gt_g = *s_base_gt;
+    const uint32_t base_eq_g = *s_base_eq;
+
+    for (int col = col0 + tid; col < col1; col += block_size) {
+        const uint32_t kp = top_k_radix_key(src[col]) & mask;
+        if (kp > prefix) {
+            dst_idx[base_gt_g + local_atomic(*s_gt).fetch_add(1u)] = col;
+        } else if (kp == prefix) {
+            const uint32_t pos = base_eq_g + local_atomic(*s_eq).fetch_add(1u);
+            if (pos < need) {
+                dst_idx[base_eq + pos] = col;
+            }
+        }
+    }
+}
+
+static void top_k_radix_split_f32_sycl(
+    ggml_backend_sycl_context & ctx,
+    const float * src,
+    int32_t * dst_indices,
+    const int64_t ncols,
+    const int64_t nrows,
+    const int k,
+    const int nparts,
+    dpct::queue_ptr main_stream
+) {
+    GGML_ASSERT(ncols <= INT32_MAX);
+    GGML_ASSERT(nparts > 1);
+
+    const int block_size = ggml_sycl_info().max_work_group_sizes[ctx.device];
+    GGML_ASSERT(block_size >= SYCL_TOP_K_RADIX_BUCKETS);
+
+    const size_t state_words = (size_t) nrows * SYCL_TOP_K_RADIX_ROW_WORDS;
+    ggml_sycl_pool_alloc<uint32_t> state_alloc(ctx.pool(), state_words);
+    uint32_t * state = state_alloc.get();
+
+    // Zero histogram, done counter and both emit counters. prefix/mask/need are seeded by
+    // the first pass, which ignores the stored values.
+    // The queue is in-order, so the passes below are already ordered after this fill.
+    SYCL_CHECK(CHECK_TRY_ERROR(main_stream->memset(state, 0, state_words * sizeof(uint32_t))));
+
+    const sycl::range<1> block_dims(block_size);
+    const sycl::range<1> grid_dims(nrows * nparts);
+
+    bool first = true;
+    for (int shift = 32 - SYCL_TOP_K_RADIX_BITS; shift >= 0; shift -= SYCL_TOP_K_RADIX_BITS) {
+        const bool is_first = first;
+        first = false;
+        main_stream->submit([&](sycl::handler &cgh) {
+            sycl::local_accessor<uint32_t, 1> slm(sycl::range<1>(SYCL_TOP_K_RADIX_HIST_SIZE + 4), cgh);
+
+            cgh.parallel_for(
+                sycl::nd_range<1>(grid_dims * block_dims, block_dims),
+                [=](sycl::nd_item<1> item_ct1) {
+                    const int g    = item_ct1.get_group(0);
+                    const int row  = g / nparts;
+                    const int part = g % nparts;
+
+                    top_k_radix_split_pass_f32(
+                        src + (int64_t) row * ncols,
+                        state + (int64_t) row * SYCL_TOP_K_RADIX_ROW_WORDS,
+                        (int) ncols, k, shift, is_first, part, nparts,
+                        slm.get_multi_ptr<sycl::access::decorated::no>().get(),
+                        item_ct1);
+                });
+        });
+    }
+
+    main_stream->submit([&](sycl::handler &cgh) {
+        sycl::local_accessor<uint32_t, 1> slm(sycl::range<1>(8), cgh);
+
+        cgh.parallel_for(
+            sycl::nd_range<1>(grid_dims * block_dims, block_dims),
+            [=](sycl::nd_item<1> item_ct1) {
+                const int g    = item_ct1.get_group(0);
+                const int row  = g / nparts;
+                const int part = g % nparts;
+
+                top_k_radix_split_emit_f32(
+                    src + (int64_t) row * ncols,
+                    dst_indices + (int64_t) row * k,
+                    state + (int64_t) row * SYCL_TOP_K_RADIX_ROW_WORDS,
+                    (int) ncols, k, part, nparts,
+                    slm.get_multi_ptr<sycl::access::decorated::no>().get(),
+                    item_ct1);
+            });
+    });
+}
+
+void ggml_sycl_top_k_radix(
+    ggml_backend_sycl_context & ctx,
+    const float *   src,
+    int32_t *       dst_indices,
+    const int64_t   ncols,
+    const int64_t   nrows,
+    const int       k,
+    dpct::queue_ptr main_stream
+) {
+    const int nparts = top_k_radix_split_groups(ctx.device, ncols, nrows);
+    if (nparts > 1) {
+        top_k_radix_split_f32_sycl(ctx, src, dst_indices, ncols, nrows, k, nparts, main_stream);
+    } else {
+        top_k_radix_f32_sycl(ctx, src, dst_indices, ncols, nrows, k, main_stream);
+    }
+}
diff --git a/ggml/src/ggml-sycl/topk-radix.hpp b/ggml/src/ggml-sycl/topk-radix.hpp
new file mode 100644
index 000000000..db479607e
--- /dev/null
+++ b/ggml/src/ggml-sycl/topk-radix.hpp
@@ -0,0 +1,24 @@
+#pragma once
+
+#include "common.hpp"
+
+// The legacy implementation uses SLM to implement sorting and top_k selection.
+// SLM is limited to 128KB on Xe, which limits how much can be sorted to k<32.
+// After a k=8, the radix selection becomes beneficial for most cases, because
+// scan-merge has (block + 1) * k pairs of (value, index). Given normal sorting of nlog(n),
+// radix-select becomes beneficial quite early. This sets it to 8 - however, the other parameters
+// (columns and rows) may also be a driving factor.
+// We select the legacy implementation for k below this constant because the overhead of radix select
+// exceeds the benefit for very small problems
+constexpr int SYCL_TOP_K_SCAN_MERGE_MAX_K = 8;
+
+// Top-k of every row of src, k indices per row into dst_indices, in no particular order.
+// Picks between the one-group-per-row and the split-row kernel from the shape and the device.
+void ggml_sycl_top_k_radix(
+    ggml_backend_sycl_context & ctx,
+    const float *   src,
+    int32_t *       dst_indices,
+    const int64_t   ncols,
+    const int64_t   nrows,
+    const int       k,
+    dpct::queue_ptr main_stream);
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index d02297cf5..f650e0123 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -11298,6 +11298,30 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
         }
     }
 
+    // qwen4exp sparse-attention indexer: nrows = n_tokens/n_stream, so tg gives nrows==1.
+    // Sweep nrows to expose how much of the device a single row leaves idle.
+    for (auto cols : {8192, 32768, 131072}) {
+        for (auto nrows : {1, 2, 4, 8, 16, 32}) {
+            test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {cols, nrows, 1, 1}, 2048));
+        }
+    }
+    // backend sampler: one row of the vocab (llama-sampler.cpp top_k)
+    for (auto k : {20, 40}) {
+        test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {151936, 1, 1, 1}, k));
+    }
+
+    // short rows, many of them: MoE routing and group selection. The opposite corner from
+    // the indexer, and the one where a work-group per row is the wasteful choice.
+    for (auto cols : {2, 16, 128, 1024}) {
+        for (auto nrows : {1024, 8192}) {
+            for (auto k : {1, 2, 8, 16, 32}) {
+                if (k <= cols) {
+                    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {cols, nrows, 1, 1}, k));
+                }
+            }
+        }
+    }
+
     for (auto nrows : {1, 4, 8, 16}) {
         for (auto cols : {128, 1024, 4096, 8192, 16384, 32768, 65536, 131072, 200000, 2000000}) {
             test_cases.emplace_back(new test_cumsum(GGML_TYPE_F32, {cols, nrows, 1, 1}));

From 3d10bcd19785c7b70626d7ded4a2276ef92bc850 Mon Sep 17 00:00:00 2001
From: Alex <59368173+AlexGabbia@users.noreply.github.com>
Date: Mon, 14 Sep 2026 13:04:05 +0200
Subject: [PATCH 23/34] llama: add Maple 20B-A1B ternary MoE architecture (CPU)
 (#27000)

* gguf-py: add Maple tensor constants

Add MODEL_ARCH.MAPLE, its "maple" name, and the tensor list for the
Maple 20B-A1B ternary MoE architecture: token embeddings, output,
attention with Q/K RMS norms, and per-expert FFN tensors.

* convert: add Maple HF->GGUF converter

Register MapleForCausalLM in the HF architecture map and add the
converter for the Maple 20B-A1B ternary MoE model: 24 layers, 256
experts with 8 active, sliding-window attention (SWA-512) interleaved
with global attention at a 3:1 ratio, partial rotary factor 0.5, and
per-expert weight stacking into merged 3D tensors.

* llama: add Maple architecture (20B-A1B ternary MoE)

Add the Maple 20B-A1B ternary MoE architecture: 24 layers, 256
experts with 8 active, sliding-window attention (SWA-512) interleaved
with global attention at a 3:1 ratio, and ternary TQ1_0/TQ2_0
quantization support.

- register LLM_ARCH_MAPLE between MAMBA2 and JAMBA
- implement llama_model_maple: Q/K RMS norms after projection (GEMMA4
  style), rope applied only on SWA layers (nope_on_global_attention),
  ISWA KV cache, and MoE FFN with swiglu gate clamp at +7 (DEEPSEEK4
  style)
- mark MAPLE as unsupported by the model saver (roundtrip skipped)

* tests: mark Maple as MoE-mandatory

Maple is always-MoE: the model throws when n_expert == 0, so the
test harness must only run the MoE config for LLM_ARCH_MAPLE.

* maple: apply review feedback (n_ff_exp_arr, get_arr, rope params)

- load_arch_hparams: use n_ff_exp_arr + n_ff_exp() accessor (upstream
  changed these from a scalar member during the rebase)
- sliding_window_pattern: get_arr, the pattern is mandatory for this arch
- partial_rotary_factor: read only from rope_parameters (base.py mirrors
  the top-level key automatically)
- document why TOKEN_EMBD/OUTPUT are forced to F16 (they are the two
  dense tensors in Maple, and the reference GGUFs ship them as F16)
- add @ModelBase.example("deepgrove/maple-preview")

* tests: add Maple to the SWA pattern array list

get_arr for maple.attention.sliding_window_pattern requires an array, but
the harness only emitted a per-layer array for the arches in its list, so
test-llama-archs -a maple failed to load the model.

Assisted-by: DeepSeek Harness

* maple: move swiglu_clamp_exp to the converter

The loader prefilled 7.0 and read the key optionally. The converter now
writes it and the loader reads it as required, because llama-graph.cpp
skips the clamp when the limit is 0 and an optional read would silently
run unclamped. The test harness provides the key for the same reason.

Also drops tensor_force_quant: base.py already forces FFN_GATE_INP to F32
and TOKEN_EMBD/OUTPUT to F16 for ternary file types.

Assisted-by: DeepSeek Harness

* convert: fix the LazyBase func signature in the Maple converter

ty flagged the stack() closure: it takes no argument, while LazyBase is
annotated with func: Callable[[Any], Any]. Pass the tensor list through
args instead of closing over it, the same way kimi_k3 does, so the
callable shape matches.

Assisted-by: DeepSeek Harness
---
 conversion/__init__.py     |   1 +
 conversion/maple.py        |  87 +++++++++++++++++++++
 gguf-py/gguf/constants.py  |  19 +++++
 src/llama-arch.cpp         |   1 +
 src/llama-arch.h           |   1 +
 src/llama-graph.cpp        |   2 +-
 src/llama-model-saver.cpp  |   1 +
 src/llama-model.cpp        |   3 +
 src/models/maple.cpp       | 150 +++++++++++++++++++++++++++++++++++++
 src/models/models.h        |  13 ++++
 tests/test-llama-archs.cpp |   9 ++-
 11 files changed, 285 insertions(+), 2 deletions(-)
 create mode 100644 conversion/maple.py
 create mode 100644 src/models/maple.cpp

diff --git a/conversion/__init__.py b/conversion/__init__.py
index 4d58bcd10..9c5d98438 100644
--- a/conversion/__init__.py
+++ b/conversion/__init__.py
@@ -168,6 +168,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
     "Mamba2ForCausalLM": "mamba",
     "MambaForCausalLM": "mamba",
     "MambaLMHeadModel": "mamba",
+    "MapleForCausalLM": "maple",
     "MellumForCausalLM": "mellum",
     "MiMoV2FlashForCausalLM": "mimo",
     "MiMoV2ForCausalLM": "mimo",
diff --git a/conversion/maple.py b/conversion/maple.py
new file mode 100644
index 000000000..fb0e87804
--- /dev/null
+++ b/conversion/maple.py
@@ -0,0 +1,87 @@
+from __future__ import annotations
+
+from typing import Iterable, TYPE_CHECKING, cast
+
+import torch
+
+if TYPE_CHECKING:
+    from torch import Tensor
+
+from .base import LazyTorchTensor, ModelBase, TextModel, gguf
+
+
+@ModelBase.register("MapleForCausalLM")
+@ModelBase.example("deepgrove/maple-preview")
+class MapleModel(TextModel):
+    model_arch = gguf.MODEL_ARCH.MAPLE
+
+    def set_gguf_parameters(self):
+        super().set_gguf_parameters()
+        hparams = self.hparams
+
+        assert hparams["hidden_act"] == "silu"
+        assert hparams.get("num_shared_experts", 0) == 0
+        assert hparams.get("norm_topk_prob", True)
+        assert hparams.get("nope_on_global_attention", False)
+
+        head_dim = hparams.get("head_dim", hparams["hidden_size"] // hparams["num_attention_heads"])
+        partial_rotary_factor = self.rope_parameters.get("partial_rotary_factor", 1.0)
+
+        self.gguf_writer.add_vocab_size(hparams["vocab_size"])
+        self.gguf_writer.add_rope_dimension_count(int(head_dim * partial_rotary_factor))
+        self.gguf_writer.add_sliding_window(hparams["sliding_window"])
+        self.gguf_writer.add_sliding_window_pattern([layer_type == "sliding_attention" for layer_type in hparams["layer_types"]])
+        self.gguf_writer.add_expert_feed_forward_length(hparams["moe_intermediate_size"])
+        # the reference clamps the MoE SwiGLU gate/up at 7.0 (modeling_maple.py)
+        self.gguf_writer.add_swiglu_clamp_exp([7.0] * self.block_count)
+
+    _experts: list[dict[str, Tensor]] | None = None
+
+    @staticmethod
+    def _stack_experts(tensors: list[Tensor]) -> Tensor:
+        shape = (len(tensors), *tensors[0].shape)
+        dtype = tensors[0].dtype
+        meta = LazyTorchTensor.meta_with_dtype_and_shape(dtype, shape)
+
+        # tensors goes through args, not the closure, so that `func` matches
+        # LazyBase's single-argument shape
+        def stack(ts: list[Tensor]) -> Tensor:
+            result = torch.empty(shape, dtype=dtype)
+            for expert_id, tensor in enumerate(ts):
+                result[expert_id].copy_(LazyTorchTensor.to_eager(tensor))
+            ts.clear()
+            return result
+
+        return cast(torch.Tensor, LazyTorchTensor(meta=meta, args=(tensors,), func=stack))
+
+    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        if "mlp.experts" in name:
+            n_experts = self.hparams["num_experts"]
+            assert bid is not None
+
+            if self._experts is None:
+                self._experts = [{} for _ in range(self.block_count)]
+
+            self._experts[bid][name] = data_torch
+
+            if len(self._experts[bid]) >= n_experts * 3:
+                for weight_name in ("down_proj", "gate_proj", "up_proj"):
+                    tensors = []
+
+                    for expert_id in range(n_experts):
+                        expert_name = f"model.layers.{bid}.mlp.experts.{expert_id}.{weight_name}.weight"
+                        tensors.append(self._experts[bid].pop(expert_name))
+
+                    merged_name = f"model.layers.{bid}.mlp.experts.{weight_name}.weight"
+                    yield from super().modify_tensors(self._stack_experts(tensors), merged_name, bid)
+            return
+
+        yield from super().modify_tensors(data_torch, name, bid)
+
+    def prepare_tensors(self):
+        super().prepare_tensors()
+
+        if self._experts is not None:
+            experts = [name for layer in self._experts for name in layer]
+            if experts:
+                raise ValueError(f"Unprocessed experts: {experts}")
diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py
index d3a639f37..e54ee5a0f 100644
--- a/gguf-py/gguf/constants.py
+++ b/gguf-py/gguf/constants.py
@@ -541,6 +541,7 @@ class MODEL_ARCH(IntEnum):
     ARWKV7           = auto()
     MAMBA            = auto()
     MAMBA2           = auto()
+    MAPLE            = auto()
     JAMBA            = auto()
     XVERSE           = auto()
     COMMAND_R        = auto()
@@ -1295,6 +1296,7 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
     MODEL_ARCH.ARWKV7:           "arwkv7",
     MODEL_ARCH.MAMBA:            "mamba",
     MODEL_ARCH.MAMBA2:           "mamba2",
+    MODEL_ARCH.MAPLE:            "maple",
     MODEL_ARCH.JAMBA:            "jamba",
     MODEL_ARCH.XVERSE:           "xverse",
     MODEL_ARCH.COMMAND_R:        "command-r",
@@ -3487,6 +3489,23 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
         MODEL_TENSOR.SSM_NORM,
         MODEL_TENSOR.SSM_OUT,
     ],
+    MODEL_ARCH.MAPLE: [
+        MODEL_TENSOR.TOKEN_EMBD,
+        MODEL_TENSOR.OUTPUT_NORM,
+        MODEL_TENSOR.OUTPUT,
+        MODEL_TENSOR.ATTN_NORM,
+        MODEL_TENSOR.ATTN_Q,
+        MODEL_TENSOR.ATTN_Q_NORM,
+        MODEL_TENSOR.ATTN_K,
+        MODEL_TENSOR.ATTN_K_NORM,
+        MODEL_TENSOR.ATTN_V,
+        MODEL_TENSOR.ATTN_OUT,
+        MODEL_TENSOR.FFN_NORM,
+        MODEL_TENSOR.FFN_GATE_INP,
+        MODEL_TENSOR.FFN_GATE_EXP,
+        MODEL_TENSOR.FFN_DOWN_EXP,
+        MODEL_TENSOR.FFN_UP_EXP,
+    ],
     MODEL_ARCH.JAMBA: [
         MODEL_TENSOR.TOKEN_EMBD,
         MODEL_TENSOR.OUTPUT_NORM,
diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp
index b5efb7206..0fac27efc 100644
--- a/src/llama-arch.cpp
+++ b/src/llama-arch.cpp
@@ -62,6 +62,7 @@ static const std::map<llm_arch, const char *> LLM_ARCH_NAMES = {
     { LLM_ARCH_STARCODER2,       "starcoder2"       },
     { LLM_ARCH_MAMBA,            "mamba"            },
     { LLM_ARCH_MAMBA2,           "mamba2"           },
+    { LLM_ARCH_MAPLE,            "maple"            },
     { LLM_ARCH_JAMBA,            "jamba"            },
     { LLM_ARCH_FALCON_H1,        "falcon-h1"        },
     { LLM_ARCH_XVERSE,           "xverse"           },
diff --git a/src/llama-arch.h b/src/llama-arch.h
index f1d173a57..6e67f5d65 100644
--- a/src/llama-arch.h
+++ b/src/llama-arch.h
@@ -67,6 +67,7 @@ enum llm_arch {
     LLM_ARCH_STARCODER2,
     LLM_ARCH_MAMBA,
     LLM_ARCH_MAMBA2,
+    LLM_ARCH_MAPLE,
     LLM_ARCH_JAMBA,
     LLM_ARCH_FALCON_H1,
     LLM_ARCH_XVERSE,
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index 5855393ef..fd4290cf0 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -2225,7 +2225,7 @@ ggml_tensor * llm_graph_context::build_moe_ffn(
                     const float limit = hparams.swiglu_clamp_exp[il];
                     constexpr float eps = 1e-6f;
                     if (limit > eps) {
-                        if (arch == LLM_ARCH_DEEPSEEK4 || (arch == LLM_ARCH_DFLASH && hparams.dsv4_hc_mult > 0) || arch == LLM_ARCH_HY_V4) {
+                        if (arch == LLM_ARCH_MAPLE || arch == LLM_ARCH_DEEPSEEK4 || (arch == LLM_ARCH_DFLASH && hparams.dsv4_hc_mult > 0) || arch == LLM_ARCH_HY_V4) {
                             cur = ggml_swiglu_clamp(ctx0, cur, up, limit);
                         } else {
                             up = ggml_clamp(ctx0, up, -limit, limit);
diff --git a/src/llama-model-saver.cpp b/src/llama-model-saver.cpp
index 66f8bdec3..59a8ff84f 100644
--- a/src/llama-model-saver.cpp
+++ b/src/llama-model-saver.cpp
@@ -33,6 +33,7 @@ bool llama_model_saver_supports_arch(llm_arch arch) {
         case LLM_ARCH_LAGUNA:
         case LLM_ARCH_GRANITE_SWA:
         case LLM_ARCH_DOTS3NOTE: // TODO: need to handle SWA pattern and MLA+SWA config
+        case LLM_ARCH_MAPLE:
             return false;
         default:
             return true;
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index f9e9a8bcb..3b2536283 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -162,6 +162,8 @@ static llama_model * llama_model_mapping(llm_arch arch, const llama_model_params
             return new llama_model_mamba(params);
         case LLM_ARCH_MAMBA2:
             return new llama_model_mamba2(params);
+        case LLM_ARCH_MAPLE:
+            return new llama_model_maple(params);
         case LLM_ARCH_JAMBA:
             return new llama_model_jamba(params);
         case LLM_ARCH_XVERSE:
@@ -3019,6 +3021,7 @@ llama_rope_type llama_model_rope_type(const llama_model * model) {
         case LLM_ARCH_SPARK2_5:
         case LLM_ARCH_TALKIE:
         case LLM_ARCH_MELLUM:
+        case LLM_ARCH_MAPLE:
             return LLAMA_ROPE_TYPE_NEOX;
 
         case LLM_ARCH_DFLASH:
diff --git a/src/models/maple.cpp b/src/models/maple.cpp
new file mode 100644
index 000000000..7604b7dfe
--- /dev/null
+++ b/src/models/maple.cpp
@@ -0,0 +1,150 @@
+#include "models.h"
+
+void llama_model_maple::load_arch_hparams(llama_model_loader & ml) {
+    hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
+
+    ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
+    ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW,    hparams.n_swa);
+    ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
+
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
+
+    hparams.rope_freq_base_train_swa  = hparams.rope_freq_base_train;
+    hparams.rope_freq_scale_train_swa = hparams.rope_freq_scale_train;
+    ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
+
+    ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_EXP, hparams.swiglu_clamp_exp, hparams.n_layer_all);
+
+    switch (hparams.n_layer()) {
+        case 24: type = LLM_TYPE_20B; break;
+        default: type = LLM_TYPE_UNKNOWN;
+    }
+}
+
+void llama_model_maple::load_arch_tensors(llama_model_loader &) {
+    LLAMA_LOAD_LOCALS;
+
+    const int64_t n_ff_exp = hparams.n_ff_exp();
+    const int64_t head_dim = hparams.n_embd_head_k();
+
+    tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
+
+    output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0);
+    output      = create_tensor(tn(LLM_TENSOR_OUTPUT,      "weight"), {n_embd, n_vocab}, 0);
+
+    if (n_expert == 0) {
+        throw std::runtime_error("n_expert must be > 0 for Maple");
+    }
+    if (n_expert_used == 0) {
+        throw std::runtime_error("n_expert_used must be > 0 for Maple");
+    }
+
+    for (int i = 0; i < n_layer; ++i) {
+        auto & layer = layers[i];
+
+        layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
+
+        create_tensor_qkv(layer, i, n_embd, n_head * head_dim, n_head_kv * head_dim, n_head_kv * head_dim, 0);
+        layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_head * head_dim, n_embd}, 0);
+
+        layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), {head_dim}, 0);
+        layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {head_dim}, 0);
+        layer.ffn_norm    = create_tensor(tn(LLM_TENSOR_FFN_NORM,    "weight", i), {n_embd}, 0);
+
+        layer.ffn_gate_inp  = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP,  "weight", i), {n_embd, n_expert}, 0);
+        layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
+        layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
+        layer.ffn_up_exps   = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS,   "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
+    }
+}
+
+std::unique_ptr<llm_graph_context> llama_model_maple::build_arch_graph(const llm_graph_params & params) const {
+    return std::make_unique<graph>(*this, params);
+}
+
+llama_model_maple::graph::graph(const llama_model & model, const llm_graph_params & params) :
+    llm_graph_context(params) {
+    const int64_t n_embd_head = hparams.n_embd_head_k();
+
+    GGML_ASSERT(n_embd_head == hparams.n_embd_head_v());
+
+    ggml_tensor * inpL = build_inp_embd(model.tok_embd);
+    ggml_tensor * inp_pos = build_inp_pos();
+    auto * inp_attn = build_attn_inp_kv_iswa();
+    ggml_tensor * inp_out_ids = build_inp_out_ids();
+
+    for (int il = 0; il < n_layer; ++il) {
+        ggml_tensor * inpSA = inpL;
+
+        ggml_tensor * cur = build_norm(inpL, model.layers[il].attn_norm, nullptr, LLM_NORM_RMS, il);
+        cb(cur, "attn_norm", il);
+
+        {
+            auto [Qcur, Kcur, Vcur] = build_qkv(model.layers[il], cur, n_embd_head, n_head, n_head_kv, il);
+
+            Qcur = build_norm(Qcur, model.layers[il].attn_q_norm, nullptr, LLM_NORM_RMS, il);
+            Kcur = build_norm(Kcur, model.layers[il].attn_k_norm, nullptr, LLM_NORM_RMS, il);
+            cb(Qcur, "Qcur_normed", il);
+            cb(Kcur, "Kcur_normed", il);
+
+            if (hparams.is_swa(il)) {
+                const int64_t n_rot_l = hparams.n_rot(il);
+                const float freq_base_l = model.get_rope_freq_base(cparams, il);
+                const float freq_scale_l = model.get_rope_freq_scale(cparams, il);
+
+                Qcur = ggml_rope_ext(ctx0, Qcur, inp_pos, nullptr, n_rot_l, rope_type, n_ctx_orig, freq_base_l,
+                                     freq_scale_l, ext_factor, attn_factor, beta_fast, beta_slow);
+                Kcur = ggml_rope_ext(ctx0, Kcur, inp_pos, nullptr, n_rot_l, rope_type, n_ctx_orig, freq_base_l,
+                                     freq_scale_l, ext_factor, attn_factor, beta_fast, beta_slow);
+            }
+            cb(Qcur, "Qcur", il);
+            cb(Kcur, "Kcur", il);
+            cb(Vcur, "Vcur", il);
+
+            cur = build_attn(inp_attn,
+                    model.layers[il].wo, nullptr, model.layers[il].wo_s,
+                    Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f / sqrtf(float(n_embd_head)), il);
+            cb(cur, "attn_out", il);
+        }
+
+        if (il == n_layer - 1 && inp_out_ids) {
+            cur   = ggml_get_rows(ctx0, cur, inp_out_ids);
+            inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
+        }
+
+        ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
+        cb(ffn_inp, "ffn_inp", il);
+
+        cur = build_norm(ffn_inp, model.layers[il].ffn_norm, nullptr, LLM_NORM_RMS, il);
+        cb(cur, "ffn_norm", il);
+
+        cur = build_moe_ffn(cur,
+                model.layers[il].ffn_gate_inp,
+                model.layers[il].ffn_up_exps,
+                model.layers[il].ffn_gate_exps,
+                model.layers[il].ffn_down_exps,
+                nullptr,
+                n_expert, n_expert_used,
+                LLM_FFN_SILU, true,
+                1.0f,
+                LLAMA_EXPERT_GATING_FUNC_TYPE_SOFTMAX,
+                il);
+        cb(cur, "ffn_moe_out", il);
+
+        cur = ggml_add(ctx0, cur, ffn_inp);
+        cur = build_cvec(cur, il);
+        cb(cur, "l_out", il);
+
+        inpL = cur;
+    }
+
+    ggml_tensor * cur = build_norm(inpL, model.output_norm, nullptr, LLM_NORM_RMS, -1);
+    cb(cur, "result_norm", -1);
+    res->t_embd = cur;
+
+    cur = build_lora_mm(model.output, cur, model.output_s);
+    cb(cur, "result_output", -1);
+    res->t_logits = cur;
+
+    ggml_build_forward_expand(gf, cur);
+}
diff --git a/src/models/models.h b/src/models/models.h
index 87195fddd..da519dcfd 100644
--- a/src/models/models.h
+++ b/src/models/models.h
@@ -945,6 +945,19 @@ struct llama_model_mamba2 : public llama_model_base {
 };
 
 
+struct llama_model_maple : public llama_model_base {
+    llama_model_maple(const struct llama_model_params & params) : llama_model_base(params) {}
+    void load_arch_hparams(llama_model_loader & ml) override;
+    void load_arch_tensors(llama_model_loader & ml) override;
+
+    struct graph : public llm_graph_context {
+        graph(const llama_model & model, const llm_graph_params & params);
+    };
+
+    std::unique_ptr<llm_graph_context> build_arch_graph(const llm_graph_params & params) const override;
+};
+
+
 struct llama_model_jamba : public llama_model_base {
     llama_model_jamba(const struct llama_model_params & params) : llama_model_base(params) {}
     void load_arch_hparams(llama_model_loader & ml) override;
diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp
index 018ff1f42..90a6a7162 100644
--- a/tests/test-llama-archs.cpp
+++ b/tests/test-llama-archs.cpp
@@ -239,7 +239,8 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
         // SWA pattern: every 5th layer is full attention (matches E2B layer_types)
         ms.add_kv(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, uint32_t(5));
     } 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_GRANITE_SWA || arch == LLM_ARCH_DOTS3NOTE) {
+            arch == LLM_ARCH_MUSE_GLIMMER || arch == LLM_ARCH_GRANITE_SWA || arch == LLM_ARCH_DOTS3NOTE ||
+            arch == LLM_ARCH_MAPLE) {
         std::vector<uint32_t> pattern;
         pattern.reserve(n_layer);
         for (uint32_t il = 0; il < n_layer; il++) {
@@ -323,6 +324,11 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
         ms.add_kv(LLM_KV_EXPERT_WEIGHTS_SCALE,                  1.0f);
         ms.add_kv(LLM_KV_EXPERT_WEIGHTS_NORM,                   true);
     }
+
+    if (arch == LLM_ARCH_MAPLE) {
+        ms.add_kv(LLM_KV_SWIGLU_CLAMP_EXP, 7.0f);
+    }
+
     ms.add_kv(LLM_KV_TOKENIZER_MODEL,         "no_vocab");
     // ms.add_kv(LLM_KV_DENSE_2_FEAT_OUT,     n_embd);
     // ms.add_kv(LLM_KV_DENSE_3_FEAT_IN,      n_embd);
@@ -505,6 +511,7 @@ static bool moe_mandatory(const llm_arch arch) {
         case LLM_ARCH_MISTRAL4:
         case LLM_ARCH_MELLUM:
         case LLM_ARCH_LAGUNA:
+        case LLM_ARCH_MAPLE:
             return true;
         default:
             return false;

From be2c6d7d1ff08b3059d8bf755623c77768f8482b Mon Sep 17 00:00:00 2001
From: Aaron Teo <aaron.teo1@ibm.com>
Date: Mon, 14 Sep 2026 19:04:58 +0800
Subject: [PATCH 24/34] tests(s390x): add non-vxe build to tests (#28776)

* tests: add non-vxe build to tests

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

ggml-cpu: add unused macro to fix ci

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

Revert "ggml-cpu: temporarily add #28775 patch until its merged"

This reverts commit d4645257b6b7e65c47b1b46baec3eb46a3f40968.

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

* ggml-cpu: revert back to upstream/master

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>

---------

Signed-off-by: Aaron Teo <aaron.teo1@ibm.com>
---
 .github/workflows/build-ibm.yml      | 10 ++++++++--
 ggml/src/ggml-cpu/arch/s390/quants.c |  1 +
 2 files changed, 9 insertions(+), 2 deletions(-)

diff --git a/.github/workflows/build-ibm.yml b/.github/workflows/build-ibm.yml
index d2e4f3cda..355487e97 100644
--- a/.github/workflows/build-ibm.yml
+++ b/.github/workflows/build-ibm.yml
@@ -34,10 +34,15 @@ env:
   LLAMA_ARG_LOG_TIMESTAMPS: 1
 
 jobs:
-
   ubuntu-24-s390x:
+    name: ubuntu-24-s390x (VXE ${{ matrix.vxe }})
     runs-on: ubuntu-24.04-s390x
 
+    strategy:
+      fail-fast: false
+      matrix:
+        vxe: ["ON", "OFF"]  # `-DGGML_VXE=ON/OFF`
+
     steps:
       - name: Clone
         id: checkout
@@ -77,7 +82,8 @@ jobs:
         run: |
           cmake -B build \
             -DLLAMA_FATAL_WARNINGS=ON \
-            -DGGML_RPC=ON
+            -DGGML_RPC=ON \
+            -DGGML_VXE=${{ matrix.vxe }}
           time cmake --build build --config Release -j $(nproc)
 
       - name: Test
diff --git a/ggml/src/ggml-cpu/arch/s390/quants.c b/ggml/src/ggml-cpu/arch/s390/quants.c
index 70f2882d8..52344828e 100644
--- a/ggml/src/ggml-cpu/arch/s390/quants.c
+++ b/ggml/src/ggml-cpu/arch/s390/quants.c
@@ -417,6 +417,7 @@ void ggml_vec_dot_mxfp4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const vo
     sumf = vec_hsum_f32x4(v_acc);
     *s = sumf;
 #else
+    UNUSED(nb);
     UNUSED(x);
     UNUSED(y);
     UNUSED(ib);

From 1aca1f9fcd238dda6ba81e787620d70ea9a10d5b Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Sigbj=C3=B8rn=20Skj=C3=A6ret?=
 <sigbjorn.skjaeret@huggingface.co>
Date: Mon, 14 Sep 2026 13:05:17 +0200
Subject: [PATCH 25/34] models : fix mimo2 swa pattern load (#28865)

---
 src/models/mimo2.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/models/mimo2.cpp b/src/models/mimo2.cpp
index 8772319f4..a466984af 100644
--- a/src/models/mimo2.cpp
+++ b/src/models/mimo2.cpp
@@ -9,7 +9,7 @@ void llama_model_mimo2::load_arch_hparams(llama_model_loader & ml) {
     ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW,   hparams.n_swa);
     ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA,         hparams.rope_freq_base_train_swa, false);
 
-    ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer());
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
 
     float value_scale = 0.0f;
     if (ml.get_key(LLM_KV_ATTENTION_VALUE_SCALE, value_scale, false) && value_scale != 1.0f) {

From 97e4ca73582084f2751767f80c86237483ecc381 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Sigbj=C3=B8rn=20Skj=C3=A6ret?=
 <sigbjorn.skjaeret@huggingface.co>
Date: Mon, 14 Sep 2026 13:05:37 +0200
Subject: [PATCH 26/34] models : fix incorrect uses of get_key_or_arr (#28868)

---
 src/models/gemma4-assistant.cpp | 2 +-
 src/models/gemma4.cpp           | 2 +-
 src/models/step35.cpp           | 2 +-
 3 files changed, 3 insertions(+), 3 deletions(-)

diff --git a/src/models/gemma4-assistant.cpp b/src/models/gemma4-assistant.cpp
index 8431ec2a1..74d06151e 100644
--- a/src/models/gemma4-assistant.cpp
+++ b/src/models/gemma4-assistant.cpp
@@ -4,7 +4,7 @@ void llama_model_gemma4_assistant::load_arch_hparams(llama_model_loader & ml) {
     hparams.n_embd_inp_impl = hparams.n_embd_out();
 
     hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
-    ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer_all);
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
 
     uint32_t n_kv_shared_layers = 0;
     ml.get_key(LLM_KV_ATTENTION_SHARED_KV_LAYERS, n_kv_shared_layers, false);
diff --git a/src/models/gemma4.cpp b/src/models/gemma4.cpp
index 39e899aa6..67de74c54 100644
--- a/src/models/gemma4.cpp
+++ b/src/models/gemma4.cpp
@@ -2,7 +2,7 @@
 
 void llama_model_gemma4::load_arch_hparams(llama_model_loader & ml) {
     hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
-    ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer());
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
 
     uint32_t n_kv_shared_layers = 0;
     ml.get_key(LLM_KV_ATTENTION_SHARED_KV_LAYERS, n_kv_shared_layers, false);
diff --git a/src/models/step35.cpp b/src/models/step35.cpp
index 946a36960..ca68855d8 100644
--- a/src/models/step35.cpp
+++ b/src/models/step35.cpp
@@ -23,7 +23,7 @@ void llama_model_step35::load_arch_hparams(llama_model_loader & ml) {
     ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW,  hparams.n_swa);
     ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA,        hparams.rope_freq_base_train_swa, false);
 
-    ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer_all);
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
 
     ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_EXP,   hparams.swiglu_clamp_exp,   hparams.n_layer_all, false);
     ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_SHEXP, hparams.swiglu_clamp_shexp, hparams.n_layer_all, false);

From bbdd9f246e9667f9aeb7ad11cca269466f981184 Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Mon, 14 Sep 2026 15:45:05 +0300
Subject: [PATCH 27/34] tests : add fusion baseline README and broaden fusion
 CI triggers (#28893)

* tests : add README for updating the per-backend fusion baselines

Assisted-by: pi:llama.cpp/Qwen3.8-27B

* ci : trigger fusion on changes to test-llama-archs.cpp and src/models

the dummy models and their architectures drive the fusion baselines, so a
change to either can alter the per-fusion counters and should re-run the
fusion job.

Assisted-by: pi:llama.cpp/Qwen3.8-27B

* tests : merge the fusion build commands in the README

assisted-by: pi:llama.cpp/Qwen3.8-27B

* pi : require explicit permission before posting PR/issue comments

assisted-by: pi:llama.cpp/Qwen3.8-27B
---
 .github/workflows/fusion.yml |  8 ++++++--
 .pi/gg/SYSTEM.md             |  1 +
 tests/fusion/README.md       | 26 ++++++++++++++++++++++++++
 3 files changed, 33 insertions(+), 2 deletions(-)
 create mode 100644 tests/fusion/README.md

diff --git a/.github/workflows/fusion.yml b/.github/workflows/fusion.yml
index ad7d5ab60..7c8596467 100644
--- a/.github/workflows/fusion.yml
+++ b/.github/workflows/fusion.yml
@@ -9,7 +9,9 @@ on:
       '.github/workflows/fusion.yml',
       'ggml/**',
       'tests/fusion/**',
-      'tests/test-fusion.cpp'
+      'tests/test-fusion.cpp',
+      'tests/test-llama-archs.cpp',
+      'src/models/**'
     ]
 
   pull_request:
@@ -18,7 +20,9 @@ on:
       '.github/workflows/fusion.yml',
       'ggml/**',
       'tests/fusion/**',
-      'tests/test-fusion.cpp'
+      'tests/test-fusion.cpp',
+      'tests/test-llama-archs.cpp',
+      'src/models/**'
     ]
 
 concurrency:
diff --git a/.pi/gg/SYSTEM.md b/.pi/gg/SYSTEM.md
index 369b87bcd..bd308ea96 100644
--- a/.pi/gg/SYSTEM.md
+++ b/.pi/gg/SYSTEM.md
@@ -23,6 +23,7 @@ Pull requests (PRs):
 - For the AI usage disclosure section, write "YES. pi:llama.cpp/[MODEL]"
 - If `PI_MODEL_NAME` env var is not set, ask the user to tell you what model was used and write it in place of [MODEL]
 - Always create the pull requests in draft mode
+- Never reply to review comments or post comments on issues/PRs without explicit permission from the user
 
 Commits:
 - On every commit that you make, include a "Assisted-by: pi:llama.cpp/[MODEL]" tag
diff --git a/tests/fusion/README.md b/tests/fusion/README.md
new file mode 100644
index 000000000..3ac02b061
--- /dev/null
+++ b/tests/fusion/README.md
@@ -0,0 +1,26 @@
+# Fusion baselines
+
+Per-device baselines for `test-fusion`, one CSV per backend (e.g. `MTL.csv`). Rows are
+`arch,moe,mode,label,count`. Regenerate a CSV whenever fusion patterns change.
+
+## Update a baseline
+
+```sh
+cmake -B build -DCMAKE_BUILD_TYPE=Release -DGGML_METAL=ON   # enable the target backend
+cmake --build build --config Release --target test-llama-archs --target test-fusion -j
+
+rm -rf build-ci-models && mkdir -p build-ci-models
+./build/bin/test-llama-archs -o build-ci-models
+
+./build/bin/test-fusion --models build-ci-models --device MTL0 --record MTL.csv
+```
+
+## Validate
+
+```sh
+./build/bin/test-fusion --models build-ci-models --device MTL0 --check MTL.csv
+```
+
+Non-zero exit means a row differs from the baseline. Use `--model FILE` to run a single
+architecture. Note `--check` only sees present rows — a fusion that stops matching is not
+reported, so diff the recorded CSV to catch removed patterns.

From eeea731613e4aefd6f37fe9bb70ecc72a59767ba Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Mon, 14 Sep 2026 16:34:19 +0300
Subject: [PATCH 28/34] ggml : bump version to 0.24.0 (ggml/1627)

---
 ggml/CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/ggml/CMakeLists.txt b/ggml/CMakeLists.txt
index ba9bc83b9..7f8a1f706 100644
--- a/ggml/CMakeLists.txt
+++ b/ggml/CMakeLists.txt
@@ -4,7 +4,7 @@ project("ggml" C CXX ASM)
 
 ### GGML Version
 set(GGML_VERSION_MAJOR 0)
-set(GGML_VERSION_MINOR 23)
+set(GGML_VERSION_MINOR 24)
 set(GGML_VERSION_PATCH 0)
 set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")
 

From d9e03f1074dbd2979126d91dce1b5d304ec8394e Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Mon, 14 Sep 2026 16:43:29 +0300
Subject: [PATCH 29/34] sync : ggml

---
 scripts/sync-ggml.last | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/scripts/sync-ggml.last b/scripts/sync-ggml.last
index 7b44a311a..d2a1bd951 100644
--- a/scripts/sync-ggml.last
+++ b/scripts/sync-ggml.last
@@ -1 +1 @@
-e91ded11bdcd78c42f9c8d3978ff6686eb4c1226
+456172ec733a135778adcd32d00e576a58232e45

From b29c606e28a01b1bc8c1351026a0fa6e616bf6c4 Mon Sep 17 00:00:00 2001
From: Georgi Gerganov <ggerganov@gmail.com>
Date: Mon, 14 Sep 2026 17:01:05 +0300
Subject: [PATCH 30/34] llama.cpp : bump version to 0.4.1 (#28900)

---
 CMakeLists.txt | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 4052fa3d6..4feaf083e 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -5,7 +5,7 @@ include(CheckIncludeFileCXX)
 ### llama.cpp version
 set(LLAMA_VERSION_MAJOR 0)
 set(LLAMA_VERSION_MINOR 4)
-set(LLAMA_VERSION_PATCH 0)
+set(LLAMA_VERSION_PATCH 1)
 set(LLAMA_VERSION_BASE "${LLAMA_VERSION_MAJOR}.${LLAMA_VERSION_MINOR}.${LLAMA_VERSION_PATCH}")
 
 # whether this is a development/nightly build

From dfe45163e1c98491bbd7ccab4339422ea3f80100 Mon Sep 17 00:00:00 2001
From: Christian Kastner <ckk@kvr.at>
Date: Mon, 14 Sep 2026 16:02:30 +0200
Subject: [PATCH 31/34] scripts: Add script to verify API/ABI compatibility
 (#28579)

---
 scripts/check-apiabi-compat.sh | 272 +++++++++++++++++++++++++++++++++
 1 file changed, 272 insertions(+)
 create mode 100755 scripts/check-apiabi-compat.sh

diff --git a/scripts/check-apiabi-compat.sh b/scripts/check-apiabi-compat.sh
new file mode 100755
index 000000000..078abb864
--- /dev/null
+++ b/scripts/check-apiabi-compat.sh
@@ -0,0 +1,272 @@
+#!/bin/sh
+# Check for backwards-incompatible API and ABI changes between two builds
+#
+# Backwards-incompatible API changes, such as removing a value from an enum,
+# are checked by abi-compliance-checker. Such changes can break compilation of
+# existing programs.
+#
+# Backwards-incompatible ABI changes, such as the removal of a public function,
+# are checked by libigail-tools. Such changes could break run-time dynamic
+# linking of existing binaries. (We don't use a-c-c for ABI checks because it
+# needs a debug build, whereas abigail does not.)
+#
+# Commands:
+#   --generate <build-dir>: Creates API/ABI dumps in <build-dir>
+#                           <build-dir> is expected to be a CMake build result
+#   --check <dir1> <dir2>:  Compares dumps in <dir1> and <dir2>
+#                           Comparison exit codes
+#                              0: all good
+#                              1: backwards-incompatible changes found
+#
+# Options:
+#   --include-path <dir>:   a-c-c calls gcc on headers; use this option to add
+#                           directories to gcc's search path
+#
+#
+# This script would typically be used before cutting a release:
+#
+#   1. Generate API/ABI dump for the old version
+#
+#      $ check-apiabi-compat.sh --generate <build-dir-old> libfoo [ libbar ...]
+#
+#   2. <update source>
+#
+#   3. Generate API/ABI dump for the new version
+#
+#      $ check-apiabi-compat.sh --generate <build-dir-new> libfoo [ libbar ...]
+#
+#   4. Compare the two dumps
+#
+#      $ check-apiabi-compat.sh --check <old-build-dir> <new-build-dir>
+#
+#      If the check exits 0, all is fine. Otherwise, backwards-incompatible
+#      changes were found, and the librar(ies) need a SOVER bump.
+set -eu
+
+# Preconditions
+if ! command -v abi-compliance-checker >/dev/null 2>&1; then
+    echo "abi-compliance-checker is not installed." >&2
+    exit 1
+elif ! command -v abidw >/dev/null 2>&1; then
+    echo "abigail-tools are not installed." >&2
+    exit 1
+fi
+
+# Some generic functions
+usage() {
+    echo "Usage: $0 [ --include-path <dir> ] --generate <build-dir> libXXX [ libYYY ... ]" >&2
+    echo "       $0 --check <old-build-dir> <new-build-dir>" >&2
+}
+
+get_cmake_project_name() {
+    sed -nr 's/^project\("(.*)".*$/\1/p' CMakeLists.txt
+}
+
+get_cmake_version() {
+    major="$(sed -nr 's/^set\([A-Z]+_VERSION_MAJOR ([0-9]+)\)$/\1/p' CMakeLists.txt)"
+    minor="$(sed -nr 's/^set\([A-Z]+_VERSION_MINOR ([0-9]+)\)$/\1/p' CMakeLists.txt)"
+    patch="$(sed -nr 's/^set\([A-Z]+_VERSION_PATCH ([0-9]+)\)$/\1/p' CMakeLists.txt)"
+    echo "$major.$minor.$patch"
+}
+
+# Option parsing and validation
+DO_GEN=0
+DO_CHECK=0
+BUILD_DIR=
+BUILD_DIR_NEW=
+INCLUDE_PATHS=
+while [ "$#" -gt 0 ]; do
+    case "$1" in
+    --generate=*)
+        DO_GEN=1
+        BUILD_DIR="${1#*=}"
+        shift
+        ;;
+    --generate)
+        DO_GEN=1
+        if [ -z "${2:-}" ]; then
+            usage
+            exit 1
+        fi
+        BUILD_DIR="$2"
+        shift 2
+        ;;
+    --check)
+        DO_CHECK=1
+        if [ -z "${2:-}" ] || [ -z "${3:-}" ]; then
+            usage
+            exit 1
+        elif ! [ -d "$2" ]; then
+            echo "$2 is not a directory." >&2
+            exit 1
+        elif ! [ -d "$3" ]; then
+            echo "$3 is not a directory." >&2
+            exit 1
+        fi
+        BUILD_DIR="$2"
+        BUILD_DIR_NEW="$3"
+        shift 3
+        ;;
+    --include-path=*)
+        INCLUDE_PATHS="$INCLUDE_PATHS ${1#*=}"
+        shift
+        ;;
+    --include-path)
+        if [ -z "${2:-}" ]; then
+            usage
+            exit 1
+        fi
+        INCLUDE_PATHS="$INCLUDE_PATHS $2"
+        shift 2
+        ;;
+    -h | --help)
+        usage
+        exit 1
+        ;;
+    -?*)
+        usage
+        exit 1
+        ;;
+    *)
+        break
+        ;;
+    esac
+done
+if [ $((DO_GEN + DO_CHECK)) -gt 1 ]; then
+    echo "Can only use one --command." >&2
+    exit
+fi
+PROJECT_NAME="$(get_cmake_project_name)"
+PROJECT_VERSION="$(get_cmake_version)"
+LIB_NAMES=""
+while [ "$#" -gt 0 ]; do
+    if [ "${1#lib}" = "$1" ]; then
+        echo "Library to check must start with libXXX." >&2
+        exit 1
+    fi
+    LIB_NAMES="$LIB_NAMES $1"
+    shift
+done
+
+dump_current_api() {
+    echo "Dumping API..."
+
+    DESCRIPTOR="$BUILD_DIR/apiabi/acc-descriptor.xml"
+    mkdir -p "$BUILD_DIR/apiabi"
+    cat >"$DESCRIPTOR" <<EOF
+<version>$PROJECT_VERSION</version>
+<headers>include</headers>
+<add_include_paths>$INCLUDE_PATHS</add_include_paths>
+EOF
+
+    # This addresses a bug between a-c-c and universal-ctags, manifested when
+    # a name is use both for a tag and a function name
+    mkdir -p "$BUILD_DIR/apiabi/.ctags.d"
+    echo "--fields=-t" >"$BUILD_DIR/apiabi/.ctags.d/acc.ctags"
+
+    # Change HOME so that .ctags.d gets picked up by universal-ctags, if used
+    HOME="$BUILD_DIR/apiabi" abi-compliance-checker \
+        -headers-only \
+        -lib "$PROJECT_NAME" \
+        -dump "$DESCRIPTOR" \
+        -log-path "$BUILD_DIR/apiabi/acc.log" \
+        -dump-path "$BUILD_DIR/apiabi/api.dump"
+    # acc generates this file with an ancient timestamp, which confuses gzip
+    touch "$BUILD_DIR/apiabi/api.dump"
+}
+
+dump_current_abi() {
+    echo "Dumping ABIs ..."
+    mkdir -p "$BUILD_DIR/apiabi"
+    # The suppressions are needed to avoid including all the internal C++
+    # symbols, and system types
+    cat >"$BUILD_DIR/apiabi/abidw.suppress" <<EOF
+[suppress_function]
+
+label = suppress internal C++ mangled functions
+symbol_name_regexp = ^_Z
+drop = yes
+
+[suppress_variable]
+label = suppress internal C++ mangled variables
+symbol_name_regexp = ^_Z
+drop = yes
+
+[suppress_type]
+label = Suppress types outside of our own source
+source_location_not_regexp = ^include/
+drop = yes
+EOF
+
+    # In abidw 2.5, handling of undefined stuff was changed a bit
+    abidw_version="$(abidw --version | sed -r 's/^abidw: ([0-9]+\.[0-9]+).*$/\1/')"
+    abidw_major="${abidw_version%.*}"
+    abidw_minor="${abidw_version#*.}"
+    if [ "$abidw_major" -gt 2 ] || [ "$abidw_minor" -gt 4 ]; then
+        abidw_undefined_syms_options="--no-load-undefined-interfaces"
+    else
+        abidw_undefined_syms_options="--drop-undefined-syms"
+    fi
+
+    for lib_name in $LIB_NAMES; do
+        # Depending on where add_library resides, the libraries can end up in
+        # build/src or build/bin
+        lib_path="$BUILD_DIR/src/$lib_name.so"
+        if ! [ -f "$lib_path" ]; then
+            lib_path="$BUILD_DIR/bin/$lib_name.so"
+            if ! [ -f "$lib_path" ]; then
+                echo "Cannot find library $lib_name.so" >&2
+                exit 1
+            fi
+        fi
+        abidw \
+            --headers-dir include \
+            "$abidw_undefined_syms_options" \
+            --suppressions "$BUILD_DIR/apiabi/abidw.suppress" \
+            --out-file "${BUILD_DIR}/apiabi/$lib_name.abi.xml" \
+            "$lib_path"
+    done
+}
+
+# Run the actual commands
+if [ "$DO_GEN" -eq 1 ]; then
+    dump_current_api
+    dump_current_abi
+    exit 0
+elif [ "$DO_CHECK" -eq 1 ]; then
+    # From here on, we don't want to exit on first error
+    set +e
+
+    abi-compliance-checker \
+        -strict \
+        -source \
+        -library "$PROJECT_NAME" \
+        -old "$BUILD_DIR/apiabi/api.dump" \
+        -new "$BUILD_DIR_NEW/apiabi/api.dump" \
+        -src-report-path "$BUILD_DIR_NEW/apiabi/api_compat_report.html"
+    API_RESULT=$?
+
+    ABI_RESULT=0
+    for xml_file in "$BUILD_DIR/apiabi/"lib*.abi.xml; do
+        xml_file_new="$BUILD_DIR_NEW/apiabi/$(basename "$xml_file")"
+
+        if ! [ -f "$xml_file_new" ]; then
+            echo "Cannot compare, missing file: $xml_file_new" >&2
+            exit 1
+        fi
+
+        abidiff "$xml_file" "$xml_file_new"
+        res=$?
+        [ "$((res & 8))" -ne 0 ] && ABI_RESULT=1
+    done
+
+    if [ "$API_RESULT" -gt 0 ]; then
+        echo "ERROR: API changed with possible backwards-compatibility problems." >&2
+    fi
+    if [ "$ABI_RESULT" -gt 0 ]; then
+        echo "ERROR: ABI changed with possible backwards-compatibility problems." >&2
+    fi
+    if [ "$((API_RESULT + ABI_RESULT))" -gt 0 ]; then
+        exit 1
+    fi
+fi

From f3a184b1534e6f4162fe88a030e9bc8dce24af33 Mon Sep 17 00:00:00 2001
From: Daniel Bevenius <daniel.bevenius@gmail.com>
Date: Mon, 14 Sep 2026 16:07:28 +0200
Subject: [PATCH 32/34] cmake : remove precompiled headers (#28892)

This commit removes the precompiled headers that I added in Commit
3bcfeb700  ("cmake : add PCH and unity build to improve build times
(#28091)").

The motivation for this is that this looked good when developing this
but has caused multiple issues that I had taken into consideration and
we have decided to remove it and only keep the unity builds from the
above commit.

Refs: https://github.com/ggml-org/llama.cpp/pull/28882#issuecomment-5662272126
---
 common/CMakeLists.txt       |  2 --
 src/CMakeLists.txt          |  1 -
 tests/CMakeLists.txt        |  2 --
 tools/mtmd/CMakeLists.txt   |  7 -------
 tools/server/CMakeLists.txt | 16 ----------------
 5 files changed, 28 deletions(-)

diff --git a/common/CMakeLists.txt b/common/CMakeLists.txt
index 38dab96ab..2b307c59d 100644
--- a/common/CMakeLists.txt
+++ b/common/CMakeLists.txt
@@ -136,8 +136,6 @@ set_target_properties(${TARGET} PROPERTIES
 target_include_directories(${TARGET} PUBLIC .)
 target_link_libraries     (${TARGET} PUBLIC vendor::nlohmann vendor::sheredom)
 target_compile_features   (${TARGET} PUBLIC cxx_std_17)
-target_precompile_headers (${TARGET} PRIVATE common.h)
-target_precompile_headers (${TARGET} PRIVATE chat.h)
 
 if (LLAMA_SUBPROCESS)
     target_compile_definitions(${TARGET} PUBLIC LLAMA_SUBPROCESS)
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index bc922b6a7..afdaddc79 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -67,7 +67,6 @@ configure_file(llama-version.h.in ${CMAKE_CURRENT_BINARY_DIR}/llama-version.h @O
 target_include_directories(llama PRIVATE . ${CMAKE_CURRENT_BINARY_DIR})
 target_include_directories(llama PUBLIC ../include)
 target_compile_features   (llama PRIVATE cxx_std_17) # don't bump
-target_precompile_headers (llama PRIVATE models/models.h)
 
 target_link_libraries(llama PUBLIC ggml)
 
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 2c05c0c93..cca90ef30 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -275,8 +275,6 @@ llama_build_and_test(
     peg-parser/test-unicode.cpp
     peg-parser/tests.h
 )
-target_precompile_headers(test-peg-parser PRIVATE peg-parser/tests.h)
-
 
 if (NOT ${CMAKE_SYSTEM_PROCESSOR} MATCHES "s390x")
     set(MODEL_NAME "tinyllamas/stories15M-q4_0.gguf")
diff --git a/tools/mtmd/CMakeLists.txt b/tools/mtmd/CMakeLists.txt
index 176eb1505..907468e87 100644
--- a/tools/mtmd/CMakeLists.txt
+++ b/tools/mtmd/CMakeLists.txt
@@ -84,13 +84,6 @@ target_link_libraries     (mtmd PUBLIC ggml llama)
 target_link_libraries     (mtmd PRIVATE Threads::Threads vendor::hash vendor::miniaudio vendor::stb vendor::sheredom)
 target_include_directories(mtmd PUBLIC  .)
 target_compile_features   (mtmd PRIVATE cxx_std_17)
-target_precompile_headers (mtmd PRIVATE models/models.h)
-
-set_source_files_properties(
-    mtmd-helper.cpp
-    mtmd-helper-gen.cpp
-    PROPERTIES SKIP_PRECOMPILE_HEADERS ON
-)
 
 if (MTMD_VIDEO)
     target_compile_definitions(mtmd PRIVATE MTMD_VIDEO)
diff --git a/tools/server/CMakeLists.txt b/tools/server/CMakeLists.txt
index 02607c838..43c245633 100644
--- a/tools/server/CMakeLists.txt
+++ b/tools/server/CMakeLists.txt
@@ -1,13 +1,5 @@
 include_directories(${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_BINARY_DIR})
 
-# MSVC emits a PCH bookkeeping symbol that WINDOWS_EXPORT_ALL_SYMBOLS exports as an ambiguous "__"
-
-set(LLAMA_SERVER_PCH ON)
-
-if (BUILD_SHARED_LIBS AND CMAKE_CXX_COMPILER_ID STREQUAL "MSVC")
-    set(LLAMA_SERVER_PCH OFF)
-endif()
-
 # server-context containing the core server logic, used by llama-server and CLI
 
 set(TARGET server-context)
@@ -41,10 +33,6 @@ target_include_directories(${TARGET} PRIVATE ../mtmd)
 target_include_directories(${TARGET} PRIVATE ${CMAKE_SOURCE_DIR})
 target_link_libraries(${TARGET} PUBLIC llama-common mtmd ${CMAKE_THREAD_LIBS_INIT})
 
-if (LLAMA_SERVER_PCH)
-    target_precompile_headers(${TARGET} PRIVATE ${CMAKE_SOURCE_DIR}/common/common.h)
-endif()
-
 # llama-server-impl: server logic, reusable by app
 
 set(TARGET llama-server-impl)
@@ -62,10 +50,6 @@ target_include_directories(${TARGET} PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
 target_include_directories(${TARGET} PRIVATE ../mtmd ${CMAKE_SOURCE_DIR})
 target_link_libraries(${TARGET} PUBLIC server-context llama-ui cpp-httplib ${CMAKE_THREAD_LIBS_INIT})
 
-if (LLAMA_SERVER_PCH)
-    target_precompile_headers(${TARGET} PRIVATE ${CMAKE_SOURCE_DIR}/common/common.h)
-endif()
-
 add_dependencies(${TARGET} llama-ui-assets)
 
 if(LLAMA_TOOLS_INSTALL)

From b4fa47d226ffe2efc7b58482291667fe89fc93ac Mon Sep 17 00:00:00 2001
From: Apoorv Parle <19315187+apparle@users.noreply.github.com>
Date: Mon, 14 Sep 2026 07:09:46 -0700
Subject: [PATCH 33/34] release : added gfx1103 to ubuntu rocm build (#28423)

---
 .github/workflows/release.yml | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml
index 9b77c2d97..91dcbe48b 100644
--- a/.github/workflows/release.yml
+++ b/.github/workflows/release.yml
@@ -1280,7 +1280,7 @@ jobs:
       matrix:
         include:
           - ROCM_VERSION: "10.0.0"
-            gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
+            gpu_targets: "gfx908;gfx90a;gfx942;gfx950;gfx1010;gfx1011;gfx1012;gfx1030;gfx1031;gfx1032;gfx1033;gfx1034;gfx1035;gfx1036;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1200;gfx1201"
             build: 'x64'
 
     steps:

From 41abbfd599fbdd3470fcae0a1fb6530ad8403cd7 Mon Sep 17 00:00:00 2001
From: Aman Gupta <amangupta052@gmail.com>
Date: Mon, 14 Sep 2026 22:16:30 +0800
Subject: [PATCH 34/34] qwen4exp: enable rms_norm + mul fusion (#28896)

* qwen4exp: enable rms_norm + mul fusion

* use TENSOR_ALLOW_RESHAPE
---
 src/models/qwen4exp.cpp | 25 +++++++++++--------------
 1 file changed, 11 insertions(+), 14 deletions(-)

diff --git a/src/models/qwen4exp.cpp b/src/models/qwen4exp.cpp
index 8ace95f73..e58d35034 100644
--- a/src/models/qwen4exp.cpp
+++ b/src/models/qwen4exp.cpp
@@ -157,7 +157,8 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
     tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), { n_embd, n_vocab }, 0);
 
     // there is no output_norm: the final hyper-connection mixer carries it
-    hc_head_norm = create_tensor(tn(LLM_TENSOR_HC_HEAD_NORM, "weight"), { hc_dim }, 0);
+    // the gammas load as [n_embd, hc] so the grouped norm multiplies them without a graph reshape
+    hc_head_norm = create_tensor(tn(LLM_TENSOR_HC_HEAD_NORM, "weight"), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
     hc_head_down = create_tensor(tn(LLM_TENSOR_HC_HEAD_DOWN, "weight"), { hc_dim, hc_lr }, 0);
     hc_head_up   = create_tensor(tn(LLM_TENSOR_HC_HEAD_UP,   "weight"), { hc_lr, hc_dim }, 0);
 
@@ -203,11 +204,11 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
         const int64_t conv_dim   = key_dim * 2 + value_dim;
 
         // two HC modules per layer: before the token mixer, before the MoE
-        layer.hc_attn_norm   = create_tensor(tn(LLM_TENSOR_HC_ATTN_NORM,   "weight", il), { hc_dim }, 0);
+        layer.hc_attn_norm   = create_tensor(tn(LLM_TENSOR_HC_ATTN_NORM,   "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
         layer.hc_attn_down   = create_tensor(tn(LLM_TENSOR_HC_ATTN_DOWN,   "weight", il), { hc_dim, hc_lr }, 0);
         layer.hc_attn_up     = create_tensor(tn(LLM_TENSOR_HC_ATTN_UP,     "weight", il), { hc_lr, hc_dim }, 0);
         layer.hc_attn_inject = create_tensor(tn(LLM_TENSOR_HC_ATTN_INJECT, "weight", il), { hc_dim, hc }, 0);
-        layer.hc_ffn_norm    = create_tensor(tn(LLM_TENSOR_HC_FFN_NORM,    "weight", il), { hc_dim }, 0);
+        layer.hc_ffn_norm    = create_tensor(tn(LLM_TENSOR_HC_FFN_NORM,    "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
         layer.hc_ffn_down    = create_tensor(tn(LLM_TENSOR_HC_FFN_DOWN,    "weight", il), { hc_dim, hc_lr }, 0);
         layer.hc_ffn_up      = create_tensor(tn(LLM_TENSOR_HC_FFN_UP,      "weight", il), { hc_lr, hc_dim }, 0);
         layer.hc_ffn_inject  = create_tensor(tn(LLM_TENSOR_HC_FFN_INJECT,  "weight", il), { hc_dim, hc }, 0);
@@ -240,9 +241,9 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
         if (hparams.is_ple(il)) {
             layer.ple_key        = create_tensor(tn(LLM_TENSOR_PLE_KEY,        "weight", il), { n_embd, hc_dim }, 0);
             layer.ple_value      = create_tensor(tn(LLM_TENSOR_PLE_VALUE,      "weight", il), { n_embd, n_embd }, 0);
-            layer.ple_norm_key   = create_tensor(tn(LLM_TENSOR_PLE_NORM_KEY,   "weight", il), { hc_dim }, 0);
-            layer.ple_norm_query = create_tensor(tn(LLM_TENSOR_PLE_NORM_QUERY, "weight", il), { hc_dim }, 0);
-            layer.ple_norm_conv  = create_tensor(tn(LLM_TENSOR_PLE_NORM_CONV,  "weight", il), { hc_dim }, 0);
+            layer.ple_norm_key   = create_tensor(tn(LLM_TENSOR_PLE_NORM_KEY,   "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
+            layer.ple_norm_query = create_tensor(tn(LLM_TENSOR_PLE_NORM_QUERY, "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
+            layer.ple_norm_conv  = create_tensor(tn(LLM_TENSOR_PLE_NORM_CONV,  "weight", il), { n_embd, hc }, TENSOR_ALLOW_RESHAPE);
             layer.ple_conv1d     = create_tensor(tn(LLM_TENSOR_PLE_CONV1D,     "weight", il), { hparams.ple_conv_kernel, hc_dim }, 0);
         }
 
@@ -275,11 +276,10 @@ ggml_tensor * llama_model_qwen4exp::graph::build_hc_mix(
     const int64_t hc_dim = hc * n_embd;
     const int64_t nt     = x->ne[2];
 
-    // grouped RMSNorm: reduce over one stream, then scale all streams with the [hc_dim] gamma
+    // grouped RMSNorm: reduce over one stream, then scale all streams with the [n_embd, hc] gamma
     // the converter folded each gamma to (1 + w)
-    ggml_tensor * xn = ggml_rms_norm(ctx0, x, hparams.f_norm_rms_eps);
+    ggml_tensor * xn = ggml_mul(ctx0, ggml_rms_norm(ctx0, x, hparams.f_norm_rms_eps), w_norm);
     xn = ggml_reshape_2d(ctx0, xn, hc_dim, nt);
-    xn = ggml_mul(ctx0, xn, w_norm);
     cb(xn, "hc_norm", il);
 
     ggml_tensor * lo = build_lora_mm(w_down, xn);
@@ -1200,13 +1200,10 @@ ggml_tensor * llama_model_qwen4exp::graph::build_ple(
     ggml_tensor * key   = build_lora_mm(model.layers[il].ple_key,   emb);
     ggml_tensor * value = build_lora_mm(model.layers[il].ple_value, emb);
 
-    // both norms group over one hc stream, with a weight over the whole hc*n_embd layout
+    // both norms group over one hc stream, with a [n_embd, hc] weight
     auto grouped_norm = [&](ggml_tensor * x, ggml_tensor * w) {
         ggml_tensor * t = ggml_reshape_3d(ctx0, x, n_embd, hc, n_tokens);
-        t = ggml_rms_norm(ctx0, t, hparams.f_norm_rms_eps);
-        t = ggml_reshape_2d(ctx0, t, hc_dim, n_tokens);
-        t = ggml_mul(ctx0, t, w);
-        return ggml_reshape_3d(ctx0, t, n_embd, hc, n_tokens);
+        return ggml_mul(ctx0, ggml_rms_norm(ctx0, t, hparams.f_norm_rms_eps), w);
     };
 
     key = grouped_norm(key, model.layers[il].ple_norm_key);