diff --git a/third_party/llama.cpp/patches/0001-Add-Skippy-ABI-and-package-writer-foundation.patch b/third_party/llama.cpp/patches/0001-Add-Skippy-ABI-and-package-writer-foundation.patch index 1265a60825..920444261e 100644 --- a/third_party/llama.cpp/patches/0001-Add-Skippy-ABI-and-package-writer-foundation.patch +++ b/third_party/llama.cpp/patches/0001-Add-Skippy-ABI-and-package-writer-foundation.patch @@ -1,7 +1,7 @@ -From 44fbcee408ecf53a7ee56963368642062715eca4 Mon Sep 17 00:00:00 2001 +From 34be3b6d89d94ac218678e7f6ada201a1bcba4c1 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:52 +1000 -Subject: [PATCH 01/16] Add Skippy ABI and package writer foundation +Subject: [PATCH 01/45] Add Skippy ABI and package writer foundation Squashes the initial Skippy ABI headers, model info entrypoints, single-stage runtime baseline, runtime-slice tensor filtering, activation-frame execution, GGUF slice writing, and package part composition. --- @@ -444,10 +444,10 @@ index 7ed490ce..a7986a87 100644 enum llm_graph_type { LLM_GRAPH_TYPE_DEFAULT, diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp -index 28f8bb79..f4310743 100644 +index 43447f57..7f23dbbc 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp -@@ -17,6 +17,16 @@ static const size_t kiB = 1024; +@@ -18,6 +18,16 @@ static const size_t kiB = 1024; static const size_t MiB = 1024*kiB; static const size_t GiB = 1024*MiB; @@ -464,7 +464,7 @@ index 28f8bb79..f4310743 100644 const char * llama_file_version_name(llama_fver version) { switch (version) { case GGUF_FILE_VERSION_V1: return "GGUF V1 (support until nov 2023)"; -@@ -1107,6 +1117,40 @@ struct ggml_tensor * llama_model_loader::create_tensor( +@@ -1094,6 +1104,40 @@ struct ggml_tensor * llama_model_loader::create_tensor( throw std::runtime_error(format("missing tensor info mapping for %s", tn.str().c_str())); } @@ -506,7 +506,7 @@ index 28f8bb79..f4310743 100644 if (info.op == GGML_OP_NONE || (flags & TENSOR_SKIP)) { const size_t nbytes = ggml_nbytes(t_meta); diff --git a/src/llama-model-loader.h b/src/llama-model-loader.h -index c476026d..3fb84fdd 100644 +index 75a3652d..92ad6430 100644 --- a/src/llama-model-loader.h +++ b/src/llama-model-loader.h @@ -20,6 +20,17 @@ using llama_buf_map = std::unordered_map; @@ -2080,5 +2080,5 @@ index 00000000..26169320 + +} -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0002-Add-early-staged-model-family-and-chat-support.patch b/third_party/llama.cpp/patches/0002-Add-early-staged-model-family-and-chat-support.patch index b8873a3f61..f26e2540fe 100644 --- a/third_party/llama.cpp/patches/0002-Add-early-staged-model-family-and-chat-support.patch +++ b/third_party/llama.cpp/patches/0002-Add-early-staged-model-family-and-chat-support.patch @@ -1,7 +1,7 @@ -From ce280251dfacfcf0f99307690fa0f0114665f572 Mon Sep 17 00:00:00 2001 +From 0f00b168cc09677e844bf8822040aa3d71f304b5 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:52 +1000 -Subject: [PATCH 02/16] Add early staged model-family and chat support +Subject: [PATCH 02/45] Add early staged model-family and chat support Squashes the first Qwen MoE staged-runtime support, native KV page basics, session reset, token EOG exposure, batched token verification, chat-template ABI, and the first wider model-family support pass. --- @@ -32,7 +32,7 @@ Squashes the first Qwen MoE staged-runtime support, native KV page basics, sessi create mode 100644 common/stage-chat.cpp diff --git a/common/CMakeLists.txt b/common/CMakeLists.txt -index 4cf580a0..b06bfeda 100644 +index 99688f53..3928ba17 100644 --- a/common/CMakeLists.txt +++ b/common/CMakeLists.txt @@ -98,6 +98,7 @@ add_library(${TARGET} @@ -42,7 +42,7 @@ index 4cf580a0..b06bfeda 100644 + stage-chat.cpp speculative.cpp speculative.h - unicode.cpp + trie.cpp diff --git a/common/stage-chat.cpp b/common/stage-chat.cpp new file mode 100644 index 00000000..39d6331a @@ -455,7 +455,7 @@ index a7986a87..33d57061 100644 public: llm_graph_input_pos(uint32_t n_pos_per_embd) : n_pos_per_embd(n_pos_per_embd) {} diff --git a/src/llama-kv-cache.cpp b/src/llama-kv-cache.cpp -index e70583e6..e8be39b0 100644 +index e25464c5..d7c7c220 100644 --- a/src/llama-kv-cache.cpp +++ b/src/llama-kv-cache.cpp @@ -4,6 +4,7 @@ @@ -863,10 +863,10 @@ index 531d99db..d9301c19 100644 // graph_build API // diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp -index f4310743..4e43bfbc 100644 +index 7f23dbbc..c94aca78 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp -@@ -1095,13 +1095,6 @@ struct ggml_tensor * llama_model_loader::create_tensor( +@@ -1082,13 +1082,6 @@ struct ggml_tensor * llama_model_loader::create_tensor( }; auto buft_for_tensor = [&](ggml_tensor * t_meta) -> ggml_backend_buffer_type_t { @@ -880,7 +880,7 @@ index f4310743..4e43bfbc 100644 // some models use the token embedding tensor as the output, but since these are used in different layers and with different ops // the tensor is duplicated // to handle this, we check if the tensor is duplicated, and if so, we assume that it is being loaded as the output tensor -@@ -1135,22 +1128,38 @@ struct ggml_tensor * llama_model_loader::create_tensor( +@@ -1122,22 +1115,38 @@ struct ggml_tensor * llama_model_loader::create_tensor( break; } @@ -929,10 +929,10 @@ index f4310743..4e43bfbc 100644 if (info.op == GGML_OP_NONE || (flags & TENSOR_SKIP)) { const size_t nbytes = ggml_nbytes(t_meta); diff --git a/src/llama-model.cpp b/src/llama-model.cpp -index 4c10e412..3c65b60e 100644 +index b100f601..07d3f4f3 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp -@@ -1393,13 +1393,13 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { +@@ -1398,13 +1398,13 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) { } // MoE expert weight scales (per-expert, shape {n_expert}) @@ -2347,5 +2347,5 @@ index 26169320..50e25c8f 100644 const char * path, struct skippy_model_info ** out_info, -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0003-Add-staged-sampling-checkpoints-and-part-loading.patch b/third_party/llama.cpp/patches/0003-Add-staged-sampling-checkpoints-and-part-loading.patch index 07cbef32fd..8a296b15a8 100644 --- a/third_party/llama.cpp/patches/0003-Add-staged-sampling-checkpoints-and-part-loading.patch +++ b/third_party/llama.cpp/patches/0003-Add-staged-sampling-checkpoints-and-part-loading.patch @@ -1,7 +1,7 @@ -From 0b214c33dd63bb1513d74f13e486d9648da1965d Mon Sep 17 00:00:00 2001 +From ab8b97e2a7169c1f6081a27c49f6bc64959f79a1 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:52 +1000 -Subject: [PATCH 03/16] Add staged sampling, checkpoints, and part loading +Subject: [PATCH 03/45] Add staged sampling, checkpoints, and part loading Squashes sampling configuration, batched verification frames, recurrent/native session checkpoints, logit bias, session trim, verification handoff fixes, split-metadata stripping, thinking toggle support, and ordered GGUF part loading. --- @@ -49,10 +49,10 @@ index 39d6331a..60d53fc4 100644 for (size_t i = 0; i < message_count; ++i) { diff --git a/include/llama.h b/include/llama.h -index a311ac20..4460a6a0 100644 +index 9fab6931..e54b64dc 100644 --- a/include/llama.h +++ b/include/llama.h -@@ -506,6 +506,13 @@ extern "C" { +@@ -514,6 +514,13 @@ extern "C" { size_t n_paths, struct llama_model_params params); @@ -245,18 +245,18 @@ index 23a46e0f..4dd985da 100644 size_t output_text_capacity, size_t * out_text_bytes, diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp -index 4e43bfbc..14f64c0f 100644 +index c94aca78..ecafbdf2 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp -@@ -531,6 +531,7 @@ llama_model_loader::llama_model_loader( +@@ -532,6 +532,7 @@ llama_model_loader::llama_model_loader( void * set_tensor_data_ud, const std::string & fname, std::vector & splits, + bool ordered_parts, FILE * file, - bool use_mmap, - bool use_direct_io, -@@ -599,11 +600,50 @@ llama_model_loader::llama_model_loader( + llama_load_mode load_mode, + bool check_tensors, +@@ -588,11 +589,50 @@ llama_model_loader::llama_model_loader( n_bytes += ggml_nbytes(cur); weights_map.emplace(tensor_name, llama_tensor_weight(files.back().get(), 0, metadata, cur)); } @@ -309,7 +309,7 @@ index 4e43bfbc..14f64c0f 100644 uint16_t idx = 0; const std::string kv_split_no = llm_kv(LLM_KV_SPLIT_NO); diff --git a/src/llama-model-loader.h b/src/llama-model-loader.h -index 3fb84fdd..d889c99d 100644 +index 92ad6430..f73fa907 100644 --- a/src/llama-model-loader.h +++ b/src/llama-model-loader.h @@ -136,6 +136,7 @@ struct llama_model_loader { @@ -318,38 +318,38 @@ index 3fb84fdd..d889c99d 100644 std::vector & splits, // optional, only need if the split does not follow naming scheme + bool ordered_parts, FILE * file, - bool use_mmap, - bool use_direct_io, + llama_load_mode load_mode, + bool check_tensors, diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp -index b66759b2..11ee5fa8 100644 +index caf7733a..7a3d105a 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp -@@ -881,7 +881,7 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std:: +@@ -885,7 +885,7 @@ static void llama_model_quantize_impl(const std::string & fname_inp, const std:: const llama_model_kv_override * kv_overrides = params->kv_overrides; std::vector splits = {}; llama_model_loader ml(/*metadata*/ nullptr, /*set_tensor_data*/ nullptr, /*set_tensor_data_ud*/ nullptr, -- fname_inp, splits, /*file*/ nullptr, use_mmap, /*use_direct_io*/ false, /*check_tensors*/ true, /*no_alloc*/ false, kv_overrides, nullptr); -+ fname_inp, splits, /*ordered_parts*/ false, /*file*/ nullptr, use_mmap, /*use_direct_io*/ false, /*check_tensors*/ true, /*no_alloc*/ false, kv_overrides, nullptr); +- fname_inp, splits, /*file*/ nullptr, /*load_mode*/ load_mode, /*check_tensors*/ true, /*no_alloc*/ false, kv_overrides, nullptr); ++ fname_inp, splits, /*ordered_parts*/ false, /*file*/ nullptr, /*load_mode*/ load_mode, /*check_tensors*/ true, /*no_alloc*/ false, kv_overrides, nullptr); ml.init_mappings(false); // no prefetching auto mparams = llama_model_default_params(); diff --git a/src/llama.cpp b/src/llama.cpp -index 0de6048f..900a87f7 100644 +index 11ac9656..48c578ee 100644 --- a/src/llama.cpp +++ b/src/llama.cpp -@@ -277,9 +277,9 @@ static bool llama_prepare_model_devices(const llama_model_params & params, llama +@@ -299,9 +299,9 @@ static bool llama_prepare_model_devices(const llama_model_params & params, llama // Returns 0 on success, -1 on error, and -2 on cancellation via llama_progress_callback static std::pair llama_model_load(struct gguf_context * metadata, llama_model_set_tensor_data_t set_tensor_data, void * set_tensor_data_ud, - const std::string & fname, std::vector & splits, FILE * file, llama_model_params & params) { + const std::string & fname, std::vector & splits, bool ordered_parts, FILE * file, llama_model_params & params) { try { -- llama_model_loader ml(metadata, set_tensor_data, set_tensor_data_ud, fname, splits, file, params.use_mmap, params.use_direct_io, -+ llama_model_loader ml(metadata, set_tensor_data, set_tensor_data_ud, fname, splits, ordered_parts, file, params.use_mmap, params.use_direct_io, +- llama_model_loader ml(metadata, set_tensor_data, set_tensor_data_ud, fname, splits, file, params.load_mode, ++ llama_model_loader ml(metadata, set_tensor_data, set_tensor_data_ud, fname, splits, ordered_parts, file, params.load_mode, params.check_tensors, params.no_alloc, params.kv_overrides, params.tensor_buft_overrides); ml.print_info(); -@@ -344,6 +344,7 @@ static struct llama_model * llama_model_load_from_file_impl( +@@ -366,6 +366,7 @@ static struct llama_model * llama_model_load_from_file_impl( void * set_tensor_data_ud, const std::string & path_model, std::vector & splits, @@ -357,7 +357,7 @@ index 0de6048f..900a87f7 100644 FILE * file, struct llama_model_params params) { { -@@ -386,7 +387,7 @@ static struct llama_model * llama_model_load_from_file_impl( +@@ -408,7 +409,7 @@ static struct llama_model * llama_model_load_from_file_impl( }; } @@ -366,16 +366,16 @@ index 0de6048f..900a87f7 100644 GGML_ASSERT(status <= 0); if (status < 0) { if (status == -1) { -@@ -414,7 +415,7 @@ struct llama_model * llama_model_init_from_user( +@@ -436,7 +437,7 @@ struct llama_model * llama_model_init_from_user( std::vector splits = {}; - params.use_mmap = false; + params.load_mode = LLAMA_LOAD_MODE_NONE; params.use_extra_bufts = false; - return llama_model_load_from_file_impl(metadata, set_tensor_data, set_tensor_data_ud, path_model, splits, /*file*/ nullptr, params); + return llama_model_load_from_file_impl(metadata, set_tensor_data, set_tensor_data_ud, path_model, splits, /*ordered_parts*/ false, /*file*/ nullptr, params); } // deprecated struct llama_model * llama_load_model_from_file( -@@ -427,7 +428,7 @@ struct llama_model * llama_model_load_from_file( +@@ -449,7 +450,7 @@ struct llama_model * llama_model_load_from_file( const char * path_model, struct llama_model_params params) { std::vector splits = {}; @@ -384,7 +384,7 @@ index 0de6048f..900a87f7 100644 } struct llama_model * llama_model_load_from_splits( -@@ -443,7 +444,27 @@ struct llama_model * llama_model_load_from_splits( +@@ -465,7 +466,27 @@ struct llama_model * llama_model_load_from_splits( for (size_t i = 0; i < n_paths; ++i) { splits.push_back(paths[i]); } @@ -413,7 +413,7 @@ index 0de6048f..900a87f7 100644 } struct llama_model * llama_model_load_from_file_ptr(FILE * file, struct llama_model_params params) { -@@ -453,7 +474,7 @@ struct llama_model * llama_model_load_from_file_ptr(FILE * file, struct llama_mo +@@ -475,7 +496,7 @@ struct llama_model * llama_model_load_from_file_ptr(FILE * file, struct llama_mo } std::string path_model; std::vector splits = {}; @@ -422,7 +422,7 @@ index 0de6048f..900a87f7 100644 } void llama_model_save_to_file(const struct llama_model * model, const char * path_model) { -@@ -578,4 +599,3 @@ const char * llama_print_system_info(void) { +@@ -600,4 +621,3 @@ const char * llama_print_system_info(void) { return s.c_str(); } @@ -1330,5 +1330,5 @@ index 50e25c8f..ab5a147c 100644 gguf_set_val_bool(out_ctx, "skippy.package.materialized", true); gguf_set_val_u32(out_ctx, "skippy.package.part_count", static_cast(input_count)); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0004-Add-lanes-external-media-and-chat-grammar-support.patch b/third_party/llama.cpp/patches/0004-Add-lanes-external-media-and-chat-grammar-support.patch index 7f47c577c1..475ea40e7f 100644 --- a/third_party/llama.cpp/patches/0004-Add-lanes-external-media-and-chat-grammar-support.patch +++ b/third_party/llama.cpp/patches/0004-Add-lanes-external-media-and-chat-grammar-support.patch @@ -1,18 +1,18 @@ -From fdaad2b1c0571b49c42a215f9803b161a007398a Mon Sep 17 00:00:00 2001 +From 59cdfb7bf5cf5e112f229ebec9a4e0fa3d954547 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:52 +1000 -Subject: [PATCH 04/16] Add lanes, external media, and chat grammar support +Subject: [PATCH 04/45] Add lanes, external media, and chat grammar support Squashes Qwen3 MoE staged execution, KV cache type and generation signal ABI, backend device selection, external media prefill, shared execution lanes, batch/flash-attention config, tool-aware chat templates, and chat grammar sampling. --- common/stage-chat.cpp | 230 ++++++++ include/skippy-signals.h | 44 ++ include/skippy.h | 165 +++--- - src/llama-kv-cache.cpp | 361 +------------ + src/llama-kv-cache.cpp | 356 +------------ src/llama-kv-cache.h | 19 - src/models/qwen3moe.cpp | 18 +- src/skippy.cpp | 1079 +++++++++++++++++++++++--------------- - 7 files changed, 1019 insertions(+), 897 deletions(-) + 7 files changed, 1017 insertions(+), 894 deletions(-) create mode 100644 include/skippy-signals.h diff --git a/common/stage-chat.cpp b/common/stage-chat.cpp @@ -573,7 +573,7 @@ index 4dd985da..3ffdb879 100644 const char * path, struct skippy_model_info ** out_info, diff --git a/src/llama-kv-cache.cpp b/src/llama-kv-cache.cpp -index e8be39b0..bbcd5b67 100644 +index d7c7c220..f74bb495 100644 --- a/src/llama-kv-cache.cpp +++ b/src/llama-kv-cache.cpp @@ -196,12 +196,10 @@ llama_kv_cache::llama_kv_cache( @@ -593,19 +593,7 @@ index e8be39b0..bbcd5b67 100644 } // [TAG_V_CACHE_VARIABLE] -@@ -323,9 +321,8 @@ llama_kv_cache::llama_kv_cache( - ggml_is_quantized(type_k) && - hparams.n_embd_head_k() % 64 == 0; - -- // always create Hadamard rotation tensors for DeepSeek lightning indexers -- if ((model.arch == LLM_ARCH_DEEPSEEK32 || model.arch == LLM_ARCH_DEEPSEEK4) && -- hparams.n_embd_head_k_full == hparams.indexer_head_size) { -+ // always create Hadamard rotation tensors for DeepSeek V3.2 DSA lightning indexer -+ if (model.arch == LLM_ARCH_DEEPSEEK32 && hparams.n_embd_head_k_full == hparams.indexer_head_size) { - attn_rot_k = true; - } - -@@ -1225,352 +1222,6 @@ ggml_tensor * llama_kv_cache::get_k_storage(int32_t il) const { +@@ -1225,352 +1223,6 @@ ggml_tensor * llama_kv_cache::get_k_storage(int32_t il) const { return layers[ikv].k; } @@ -2392,5 +2380,5 @@ index ab5a147c..e4d08ee4 100644 } -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0005-Add-resident-prefix-cache-and-session-refinements.patch b/third_party/llama.cpp/patches/0005-Add-resident-prefix-cache-and-session-refinements.patch index 60f1001741..b6165645af 100644 --- a/third_party/llama.cpp/patches/0005-Add-resident-prefix-cache-and-session-refinements.patch +++ b/third_party/llama.cpp/patches/0005-Add-resident-prefix-cache-and-session-refinements.patch @@ -1,15 +1,15 @@ -From 6ef5b5c66a88af43d7f4d853a848ce638d0db170 Mon Sep 17 00:00:00 2001 +From 8d4cf568f56593d1fa8223ad12ca2965d680fee6 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:52 +1000 -Subject: [PATCH 05/16] Add resident prefix cache and session refinements +Subject: [PATCH 05/45] Add resident prefix cache and session refinements Squashes resident prefix cache support, exact-prefix state restoration, borrowed prefix sessions, lazy generation signals, single-token allocation avoidance, server-like thread defaults, prefix cache borrowing, lane preservation, and recurrent-state remapping. --- include/skippy.h | 119 +++++- - src/llama-kv-cache.cpp | 357 ++++++++++++++++++ + src/llama-kv-cache.cpp | 369 +++++++++++++++++++ src/llama-kv-cache.h | 21 ++ src/skippy.cpp | 821 ++++++++++++++++++++++++++++++++++++++++- - 4 files changed, 1308 insertions(+), 10 deletions(-) + 4 files changed, 1320 insertions(+), 10 deletions(-) diff --git a/include/skippy.h b/include/skippy.h index 3ffdb879..deb5151e 100644 @@ -196,10 +196,10 @@ index 3ffdb879..deb5151e 100644 struct skippy_model * model, const char * text, diff --git a/src/llama-kv-cache.cpp b/src/llama-kv-cache.cpp -index bbcd5b67..539666f9 100644 +index f74bb495..41a4e0de 100644 --- a/src/llama-kv-cache.cpp +++ b/src/llama-kv-cache.cpp -@@ -1222,6 +1222,363 @@ ggml_tensor * llama_kv_cache::get_k_storage(int32_t il) const { +@@ -1223,6 +1223,375 @@ ggml_tensor * llama_kv_cache::get_k_storage(int32_t il) const { return layers[ikv].k; } @@ -248,6 +248,10 @@ index bbcd5b67..539666f9 100644 + return false; + } + const auto & cells = v_cells[strm]; ++ if (token_count > static_cast(cells.size())) { ++ error = "token_count exceeds KV cache capacity"; ++ return false; ++ } + + std::vector cell_idxs(static_cast(token_count), std::numeric_limits::max()); + uint32_t found_cells = 0; @@ -438,6 +442,10 @@ index bbcd5b67..539666f9 100644 + error = "sequence id is mapped to an invalid KV stream"; + return false; + } ++ if (desc.token_count > static_cast(v_cells[strm].size())) { ++ error = "token_count exceeds KV cache capacity"; ++ return false; ++ } + + std::vector selected; + selected.reserve(layers.size()); @@ -484,6 +492,10 @@ index bbcd5b67..539666f9 100644 + auto udata = std::make_shared(); + const uint32_t n_tokens = static_cast(desc.token_count); + const uint32_t n_pos = hparams.n_pos_per_embd(); ++ if (n_pos != 0 && n_tokens > std::numeric_limits::max() / n_pos) { ++ error = "token position count exceeds addressable memory"; ++ return false; ++ } + udata->token.resize(n_tokens); + udata->pos.resize(static_cast(n_tokens) * n_pos); + udata->n_seq_id.resize(n_tokens, 1); @@ -1569,5 +1581,5 @@ index e4d08ee4..f078f0b9 100644 struct skippy_session * session, uint64_t * out_token_count, -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0006-Expand-staged-execution-across-dense-and-recurrent-f.patch b/third_party/llama.cpp/patches/0006-Expand-staged-execution-across-dense-and-recurrent-f.patch index 67fda87723..709f3e5630 100644 --- a/third_party/llama.cpp/patches/0006-Expand-staged-execution-across-dense-and-recurrent-f.patch +++ b/third_party/llama.cpp/patches/0006-Expand-staged-execution-across-dense-and-recurrent-f.patch @@ -1,7 +1,7 @@ -From 3067d86be715f1933b80b60950f8fe254d441380 Mon Sep 17 00:00:00 2001 +From 7432fb6267d078b2e6c79df8cc301517913cf6c0 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:53 +1000 -Subject: [PATCH 06/16] Expand staged execution across dense and recurrent +Subject: [PATCH 06/45] Expand staged execution across dense and recurrent families Squashes the dense and decoder family support waves through LFM2, Mamba, Jamba, RWKV6, Qwen2MoE, and native sequence id exposure. @@ -1513,5 +1513,5 @@ index f078f0b9..1df8a462 100644 const struct skippy_session * session) { return session != nullptr && session->ctx != nullptr ? llama_n_batch(session->ctx) : 0; -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0007-Expand-staged-execution-across-VL-and-broad-model-fa.patch b/third_party/llama.cpp/patches/0007-Expand-staged-execution-across-VL-and-broad-model-fa.patch index c0fef6cbe7..262b56fc3e 100644 --- a/third_party/llama.cpp/patches/0007-Expand-staged-execution-across-VL-and-broad-model-fa.patch +++ b/third_party/llama.cpp/patches/0007-Expand-staged-execution-across-VL-and-broad-model-fa.patch @@ -1,7 +1,7 @@ -From 021062e62500f50ae968e376d4e45c71cbe60307 Mon Sep 17 00:00:00 2001 +From a4be6bd95fb953270ad66f37355eb8e154387d6e Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:53 +1000 -Subject: [PATCH 07/16] Expand staged execution across VL and broad model +Subject: [PATCH 07/45] Expand staged execution across VL and broad model families Squashes RWKV7 activation sideband support, Phi2, Granite, Hunyuan Dense/MoE, PhiMoE, tied-output handling, Qwen35/Hunyuan-VL, broad llama-family coverage, BitNet and StarCoder fixes, and DeepSeek OCR/Qwen3-VL-MoE support. @@ -244,10 +244,10 @@ index 33d57061..d9e241a1 100644 std::vector t_layer_inp; diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp -index 14f64c0f..683450ac 100644 +index ecafbdf2..91589e1b 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp -@@ -18,6 +18,7 @@ static const size_t MiB = 1024*kiB; +@@ -19,6 +19,7 @@ static const size_t MiB = 1024*kiB; static const size_t GiB = 1024*MiB; static thread_local llama_model_loader_stage_filter g_skippy_filter; @@ -255,7 +255,7 @@ index 14f64c0f..683450ac 100644 void llama_model_loader_set_stage_filter(const llama_model_loader_stage_filter & filter) { g_skippy_filter = filter; -@@ -27,6 +28,10 @@ void llama_model_loader_clear_stage_filter() { +@@ -28,6 +29,10 @@ void llama_model_loader_clear_stage_filter() { g_skippy_filter = {}; } @@ -266,7 +266,7 @@ index 14f64c0f..683450ac 100644 const char * llama_file_version_name(llama_fver version) { switch (version) { case GGUF_FILE_VERSION_V1: return "GGUF V1 (support until nov 2023)"; -@@ -1187,6 +1192,9 @@ struct ggml_tensor * llama_model_loader::create_tensor( +@@ -1174,6 +1179,9 @@ struct ggml_tensor * llama_model_loader::create_tensor( if (!(flags & TENSOR_DUPLICATED)) { n_created++; } @@ -276,7 +276,7 @@ index 14f64c0f..683450ac 100644 } return nullptr; -@@ -1194,12 +1202,15 @@ struct ggml_tensor * llama_model_loader::create_tensor( +@@ -1181,12 +1189,15 @@ struct ggml_tensor * llama_model_loader::create_tensor( } if (!t_meta) { @@ -293,7 +293,7 @@ index 14f64c0f..683450ac 100644 if (info.op == GGML_OP_NONE || (flags & TENSOR_SKIP)) { const size_t nbytes = ggml_nbytes(t_meta); diff --git a/src/llama-model-loader.h b/src/llama-model-loader.h -index d889c99d..d85aa85d 100644 +index f73fa907..5d693719 100644 --- a/src/llama-model-loader.h +++ b/src/llama-model-loader.h @@ -30,6 +30,7 @@ struct llama_model_loader_stage_filter { @@ -305,10 +305,10 @@ index d889c99d..d85aa85d 100644 enum llama_fver { GGUF_FILE_VERSION_V1 = 1, diff --git a/src/llama-model.cpp b/src/llama-model.cpp -index 3c65b60e..474daf17 100644 +index 07d3f4f3..b603f869 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp -@@ -2758,7 +2758,8 @@ void llama_model_base::create_tensor_qkv(llama_layer & layer, int bid, +@@ -2763,7 +2763,8 @@ void llama_model_base::create_tensor_qkv(llama_layer & layer, int bid, int flags) { const int64_t n_embd_qkv = n_embd_q_ + n_embd_k_ + n_embd_v_; layer.wqkv = create_tensor(tn(LLM_TENSOR_ATTN_QKV, "weight", bid), {n_embd_, n_embd_qkv}, TENSOR_NOT_REQUIRED | TENSOR_SKIP_IF_VIRTUAL); @@ -3221,5 +3221,5 @@ index 1df8a462..0290f045 100644 static enum skippy_status skippy_validate_state_range( -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0008-Add-external-decode-media-prefill-and-newer-family-s.patch b/third_party/llama.cpp/patches/0008-Add-external-decode-media-prefill-and-newer-family-s.patch index 85a7b1a35a..946db2893d 100644 --- a/third_party/llama.cpp/patches/0008-Add-external-decode-media-prefill-and-newer-family-s.patch +++ b/third_party/llama.cpp/patches/0008-Add-external-decode-media-prefill-and-newer-family-s.patch @@ -1,7 +1,7 @@ -From c1d935460e27f2522207cd248b0cb95ab10d2b79 Mon Sep 17 00:00:00 2001 +From 3653cfda98f3349e7f48490c90128825c8d47762 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:53 +1000 -Subject: [PATCH 08/16] Add external decode, media prefill, and newer family +Subject: [PATCH 08/45] Add external decode, media prefill, and newer family support Squashes tensor filtering accounting fixes, recurrent KV page export handling, external decode filtering, staged external media prefill, tensor element counts, activation-frame batch padding and position sideband, plus Gemma3n, Llama4, and Mistral4 slices. @@ -322,10 +322,10 @@ index d9e241a1..2e0d9c3b 100644 std::vector t_layer_inp; diff --git a/src/llama-kv-cache.cpp b/src/llama-kv-cache.cpp -index 539666f9..e5c0a061 100644 +index 41a4e0de..0e6902aa 100644 --- a/src/llama-kv-cache.cpp +++ b/src/llama-kv-cache.cpp -@@ -1301,14 +1301,16 @@ bool llama_kv_cache::stage_export_kv_page( +@@ -1302,14 +1302,16 @@ bool llama_kv_cache::stage_export_kv_page( if (il < static_cast(layer_start) || il >= static_cast(layer_end)) { continue; } @@ -347,7 +347,7 @@ index 539666f9..e5c0a061 100644 if (k_row > std::numeric_limits::max() || v_row > std::numeric_limits::max() || -@@ -1321,19 +1323,19 @@ bool llama_kv_cache::stage_export_kv_page( +@@ -1322,19 +1324,19 @@ bool llama_kv_cache::stage_export_kv_page( v_row_bytes = static_cast(v_row); v_element_bytes = static_cast(v_el); k_type = static_cast(layer.k_stream[strm]->type); @@ -372,7 +372,7 @@ index 539666f9..e5c0a061 100644 const uint64_t k_bytes = token_count * k_row; if (payload_bytes > std::numeric_limits::max() - k_bytes || payload_bytes + k_bytes > std::numeric_limits::max() - v_bytes) { -@@ -1387,6 +1389,9 @@ bool llama_kv_cache::stage_export_kv_page( +@@ -1388,6 +1390,9 @@ bool llama_kv_cache::stage_export_kv_page( if (!v_trans) { for (const auto * layer : selected) { @@ -382,7 +382,7 @@ index 539666f9..e5c0a061 100644 auto * v = layer->v_stream[strm]; for (uint32_t cell_idx : cell_idxs) { ggml_backend_tensor_get(v, dst, static_cast(cell_idx) * v_row_bytes, v_row_bytes); -@@ -1465,15 +1470,17 @@ bool llama_kv_cache::stage_import_kv_page( +@@ -1466,15 +1471,17 @@ bool llama_kv_cache::stage_import_kv_page( if (il < static_cast(desc.layer_start) || il >= static_cast(desc.layer_end)) { continue; } @@ -406,7 +406,7 @@ index 539666f9..e5c0a061 100644 desc.k_row_bytes != k_row || desc.v_row_bytes != v_row || desc.v_element_bytes != v_el) { -@@ -1556,6 +1563,9 @@ bool llama_kv_cache::stage_import_kv_page( +@@ -1557,6 +1564,9 @@ bool llama_kv_cache::stage_import_kv_page( auto & cells = v_cells[strm]; if (!v_trans) { for (const auto * layer : selected) { @@ -417,10 +417,10 @@ index 539666f9..e5c0a061 100644 for (uint32_t i = 0; i < n_tokens; ++i) { ggml_backend_tensor_set(v, src, static_cast(sinfo.idxs[0][i]) * desc.v_row_bytes, desc.v_row_bytes); diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp -index 683450ac..c55ed668 100644 +index 91589e1b..a62c3644 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp -@@ -1139,6 +1139,19 @@ struct ggml_tensor * llama_model_loader::create_tensor( +@@ -1126,6 +1126,19 @@ struct ggml_tensor * llama_model_loader::create_tensor( return it->second.get(); }; @@ -440,7 +440,7 @@ index 683450ac..c55ed668 100644 auto buft_for_tensor = [&](ggml_tensor * t_meta) -> ggml_backend_buffer_type_t { // some models use the token embedding tensor as the output, but since these are used in different layers and with different ops // the tensor is duplicated -@@ -1179,18 +1192,34 @@ struct ggml_tensor * llama_model_loader::create_tensor( +@@ -1166,18 +1179,34 @@ struct ggml_tensor * llama_model_loader::create_tensor( tn.tensor == LLM_TENSOR_PER_LAYER_MODEL_PROJ || tn.tensor == LLM_TENSOR_PER_LAYER_PROJ_NORM; } @@ -480,7 +480,7 @@ index 683450ac..c55ed668 100644 g_skippy_last_tensor_filtered = true; } else { diff --git a/src/llama-model-loader.h b/src/llama-model-loader.h -index d85aa85d..6876262d 100644 +index 5d693719..da58661f 100644 --- a/src/llama-model-loader.h +++ b/src/llama-model-loader.h @@ -14,6 +14,7 @@ @@ -1237,5 +1237,5 @@ index 0290f045..27dc78b0 100644 } -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0009-Add-chat-grammar-device-enumeration-and-runtime-even.patch b/third_party/llama.cpp/patches/0009-Add-chat-grammar-device-enumeration-and-runtime-even.patch index 9be7cd87c1..a830e8d881 100644 --- a/third_party/llama.cpp/patches/0009-Add-chat-grammar-device-enumeration-and-runtime-even.patch +++ b/third_party/llama.cpp/patches/0009-Add-chat-grammar-device-enumeration-and-runtime-even.patch @@ -1,7 +1,7 @@ -From bd3343f6b65ad898f32c1179cdd8424275342df6 Mon Sep 17 00:00:00 2001 +From 8619285489ebdbaf28638a8cf598c00e2cbf66ec Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:53 +1000 -Subject: [PATCH 09/16] Add chat grammar, device enumeration, and runtime +Subject: [PATCH 09/45] Add chat grammar, device enumeration, and runtime events Squashes chat grammar exception containment and lazy triggering, backend device enumeration, min-p sampling, KV cell compaction, and model-open runtime events. @@ -20,7 +20,7 @@ Squashes chat grammar exception containment and lazy triggering, backend device create mode 100644 include/skippy/devices.h diff --git a/common/chat-auto-parser-generator.cpp b/common/chat-auto-parser-generator.cpp -index 3f91364c..4b582932 100644 +index af84ff32..a68e2c91 100644 --- a/common/chat-auto-parser-generator.cpp +++ b/common/chat-auto-parser-generator.cpp @@ -23,6 +23,18 @@ static void foreach_function(const json & tools, const std::function & ubatches) { -@@ -1166,6 +1183,93 @@ void llama_kv_cache::apply_ubatch(const slot_info & sinfo, const llama_ubatch & +@@ -1167,6 +1184,93 @@ void llama_kv_cache::apply_ubatch(const slot_info & sinfo, const llama_ubatch & } } @@ -575,7 +575,7 @@ index e5c0a061..43e9492a 100644 bool llama_kv_cache::get_can_shift() const { // Step35 uses per-layer RoPE dims; K-shift assumes a single global n_rot. if (model.arch == LLM_ARCH_STEP35) { -@@ -2878,8 +2982,9 @@ llama_kv_cache_context::llama_kv_cache_context( +@@ -2884,8 +2988,9 @@ llama_kv_cache_context::llama_kv_cache_context( llama_kv_cache * kv, llama_context * lctx, bool do_shift, @@ -1451,18 +1451,18 @@ index 27dc78b0..4dcd0d96 100644 enum skippy_status skippy_model_free( diff --git a/tests/test-chat-auto-parser.cpp b/tests/test-chat-auto-parser.cpp -index 78e42c65..2d854509 100644 +index 4218f8d5..2cb295c1 100644 --- a/tests/test-chat-auto-parser.cpp +++ b/tests/test-chat-auto-parser.cpp -@@ -57,6 +57,7 @@ static void test_seed_oss_tool_with_reasoning(testing & t); - static void test_nemotron_analysis(testing & t); - static void test_nemotron_reasoning_detection(testing & t); - static void test_nemotron_tool_format(testing & t); +@@ -66,6 +66,7 @@ static void test_laguna_s_tool_format(testing & t); + static void test_laguna_xs2_analysis(testing & t); + static void test_laguna_xs2_reasoning_detection(testing & t); + static void test_laguna_xs2_tool_format(testing & t); +static void test_nemotron_lazy_tool_trigger(testing & t); // CohereForAI template analysis tests static void test_cohere_reasoning_detection(testing & t); -@@ -1306,6 +1307,7 @@ static common_chat_template load_nemotron_template(testing & t) { +@@ -1318,6 +1319,7 @@ static common_chat_template load_nemotron_template(testing & t) { static void test_nemotron_analysis(testing & t) { t.test("Nemotron reasoning detection", test_nemotron_reasoning_detection); t.test("Nemotron tool format", test_nemotron_tool_format); @@ -1470,8 +1470,8 @@ index 78e42c65..2d854509 100644 } static void test_nemotron_reasoning_detection(testing & t) { -@@ -1378,6 +1380,30 @@ static void test_nemotron_tool_format(testing & t) { - t.assert_true("should support tools", analysis.jinja_caps.supports_tools); +@@ -1478,6 +1480,30 @@ static void test_laguna_xs2_analysis(testing & t) { + t.test("Laguna-XS.2(v5) tool format", test_laguna_xs2_tool_format); } +static void test_nemotron_lazy_tool_trigger(testing & t) { @@ -1501,11 +1501,11 @@ index 78e42c65..2d854509 100644 static common_chat_template load_cohere_template(testing & t) { return load_template(t, "models/templates/CohereForAI-c4ai-command-r7b-12-2024-tool_use.jinja"); } -@@ -2092,4 +2118,3 @@ static void test_tagged_args_with_embedded_quotes(testing & t) { +@@ -2192,4 +2218,3 @@ static void test_tagged_args_with_embedded_quotes(testing & t) { } } } - -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0010-Add-MTP-execution-support-and-sampling-cleanup.patch b/third_party/llama.cpp/patches/0010-Add-MTP-execution-support-and-sampling-cleanup.patch index c62a10b91e..ba24eab26b 100644 --- a/third_party/llama.cpp/patches/0010-Add-MTP-execution-support-and-sampling-cleanup.patch +++ b/third_party/llama.cpp/patches/0010-Add-MTP-execution-support-and-sampling-cleanup.patch @@ -1,7 +1,7 @@ -From 4e961a921cfe7a5977ac781e8dfbfcdb04a7e725 Mon Sep 17 00:00:00 2001 +From 42f14cd0687b5a880fa2a4f1db51ca035b5ba326 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 26 Jun 2026 13:47:53 +1000 -Subject: [PATCH 10/16] Add MTP execution support and sampling cleanup +Subject: [PATCH 10/45] Add MTP execution support and sampling cleanup Squashes GLM MTP graph support, batched stage execution, native MTP sidecar support, batched MTP verification drafts, sampling and detokenization parity, unused ABI removal, MTP layer-count and sidecar fixes, VerifySpan copy avoidance, greedy sampling fast path, GLM chat fallback, and GLM DSA IndexShare support. --- @@ -13,16 +13,16 @@ Squashes GLM MTP graph support, batched stage execution, native MTP sidecar supp src/llama-model.cpp | 8 + src/models/deepseek2.cpp | 288 ++++++++++- src/models/deepseek32.cpp | 16 +- - src/models/glm-dsa.cpp | 12 +- + src/models/glm-dsa.cpp | 11 +- src/models/models.h | 5 + src/skippy.cpp | 1034 +++++++++++++++++++++++++++++++------ - 11 files changed, 1371 insertions(+), 183 deletions(-) + 11 files changed, 1371 insertions(+), 182 deletions(-) diff --git a/common/chat.cpp b/common/chat.cpp -index 22d2ee4a..4ed58a77 100644 +index 7a6e7238..91658e6d 100644 --- a/common/chat.cpp +++ b/common/chat.cpp -@@ -663,6 +663,113 @@ std::string common_chat_format_example(const struct common_chat_templates * +@@ -665,6 +665,113 @@ std::string common_chat_format_example(const struct common_chat_templates * " {{- '<|im_start|>assistant\n' -}}\n" \ "{%- endif -%}" @@ -136,7 +136,7 @@ index 22d2ee4a..4ed58a77 100644 void common_chat_templates_free(struct common_chat_templates * tmpls) { delete tmpls; } -@@ -727,6 +834,10 @@ common_chat_templates_ptr common_chat_templates_init(const struct llama_model * +@@ -729,6 +836,10 @@ common_chat_templates_ptr common_chat_templates_init(const struct llama_model * } else { default_template_src = chat_template_override; } @@ -337,10 +337,10 @@ index fb804ffd..a75f0b75 100644 ggml_tensor * kq_mask_all = ggml_fill(ctx0, kq_mask, -INFINITY); diff --git a/src/llama-model.cpp b/src/llama-model.cpp -index 474daf17..5b83c0c0 100644 +index b603f869..10633a51 100644 --- a/src/llama-model.cpp +++ b/src/llama-model.cpp -@@ -2190,6 +2190,14 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, +@@ -2196,6 +2196,14 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, } } @@ -747,10 +747,10 @@ index 32262e68..ecec31fb 100644 ggml_tensor * q = ggml_mul_mat(ctx0, model.layers[il].wq_b, qr); diff --git a/src/models/glm-dsa.cpp b/src/models/glm-dsa.cpp -index 32fe6def..10bf235f 100644 +index df190e1f..1d90426d 100644 --- a/src/models/glm-dsa.cpp +++ b/src/models/glm-dsa.cpp -@@ -75,11 +75,6 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { +@@ -110,11 +110,6 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { for (int i = 0; i < n_layer_all; ++i) { int flags = 0; @@ -762,7 +762,7 @@ index 32fe6def..10bf235f 100644 auto & layer = layers[i]; -@@ -132,7 +127,7 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { +@@ -167,7 +162,7 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { layer.ffn_up_shexp = create_tensor(tn(LLM_TENSOR_FFN_UP_SHEXP, "weight", i), {n_embd, n_ff_exp * n_expert_shared}, flags); } @@ -771,7 +771,7 @@ index 32fe6def..10bf235f 100644 if (i >= n_layer) { layer.nextn.eh_proj = create_tensor(tn(LLM_TENSOR_NEXTN_EH_PROJ, "weight", i), { 2 * n_embd, n_embd }, flags); layer.nextn.enorm = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM, "weight", i), { n_embd }, flags); -@@ -147,6 +142,9 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { +@@ -182,6 +177,10 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { } std::unique_ptr llama_model_glm_dsa::build_arch_graph(const llm_graph_params & params) const { @@ -781,9 +781,9 @@ index 32fe6def..10bf235f 100644 + return std::make_unique(*this, params); } -- + diff --git a/src/models/models.h b/src/models/models.h -index a86ae05a..e9a7c0c2 100644 +index 76daa8cc..f3b576c0 100644 --- a/src/models/models.h +++ b/src/models/models.h @@ -1068,6 +1068,10 @@ struct llama_model_deepseek2 : public llama_model_base { @@ -797,10 +797,10 @@ index a86ae05a..e9a7c0c2 100644 std::unique_ptr build_arch_graph(const llm_graph_params & params) const override; }; -@@ -1218,6 +1222,7 @@ struct llama_model_glm_dsa : public llama_model_base { - void load_arch_tensors(llama_model_loader & ml) override; - - using graph = llama_model_deepseek2::graph; +@@ -1220,6 +1224,7 @@ struct llama_model_glm_dsa : public llama_model_base { + struct graph : public llm_graph_context { + graph(const llama_model & model, const llm_graph_params & params); + }; + using graph_mtp = llama_model_deepseek2::graph_mtp; std::unique_ptr build_arch_graph(const llm_graph_params & params) const override; @@ -2208,5 +2208,5 @@ index 4dcd0d96..8af180e5 100644 } -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0011-Pass-reasoning-format-through-stage-chat-templates.patch b/third_party/llama.cpp/patches/0011-Pass-reasoning-format-through-stage-chat-templates.patch index e963607026..c4c30a9fdc 100644 --- a/third_party/llama.cpp/patches/0011-Pass-reasoning-format-through-stage-chat-templates.patch +++ b/third_party/llama.cpp/patches/0011-Pass-reasoning-format-through-stage-chat-templates.patch @@ -1,7 +1,7 @@ -From 5af8ab8b7c670470d36f3f8756d716a3548dc895 Mon Sep 17 00:00:00 2001 +From 3b3d6f4ab0c269931f5f75120f3d8b496917e3b7 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 3 Jul 2026 13:15:19 +1000 -Subject: [PATCH 11/16] Pass reasoning format through stage chat templates +Subject: [PATCH 11/45] Pass reasoning format through stage chat templates --- common/stage-chat.cpp | 4 ++++ @@ -57,5 +57,5 @@ index cfb4a8fa..0e1f0351 100644 enum skippy_feature { SKIPPY_FEATURE_RUNTIME_SLICE = 1 << 0, -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0012-Wire-mmap-and-mlock-runtime-load-options.patch b/third_party/llama.cpp/patches/0012-Wire-mmap-and-mlock-runtime-load-options.patch index a87b7e7259..64904604df 100644 --- a/third_party/llama.cpp/patches/0012-Wire-mmap-and-mlock-runtime-load-options.patch +++ b/third_party/llama.cpp/patches/0012-Wire-mmap-and-mlock-runtime-load-options.patch @@ -1,7 +1,7 @@ -From 1bd70926c77b6c94f59829bd4a38bf01adf13d47 Mon Sep 17 00:00:00 2001 +From 08a6cf87151e9f2460482d2889cb9f3750d47316 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Wed, 1 Jul 2026 23:52:26 -0400 -Subject: [PATCH 12/16] Wire mmap and mlock runtime load options +Subject: [PATCH 12/45] Wire mmap and mlock runtime load options --- include/skippy.h | 3 +++ @@ -64,5 +64,5 @@ index 8af180e5..fcd8a48f 100644 params.use_extra_bufts = false; } -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0013-Add-external-MTP-draft-sidecar-attachment.patch b/third_party/llama.cpp/patches/0013-Add-external-MTP-draft-sidecar-attachment.patch index c5e1e13dee..a824a13519 100644 --- a/third_party/llama.cpp/patches/0013-Add-external-MTP-draft-sidecar-attachment.patch +++ b/third_party/llama.cpp/patches/0013-Add-external-MTP-draft-sidecar-attachment.patch @@ -1,7 +1,7 @@ -From 9543737c2487088776b078eaa94fbd8266870403 Mon Sep 17 00:00:00 2001 +From 4724a2467f59db12a2638a417b70577a5f1df6b4 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 3 Jul 2026 12:18:37 -0400 -Subject: [PATCH 13/16] Add external MTP draft sidecar attachment +Subject: [PATCH 13/45] Add external MTP draft sidecar attachment --- include/skippy.h | 6 ++++ @@ -172,5 +172,5 @@ index fcd8a48f..db17f7d2 100644 llama_free(model->ctx); } -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0014-Add-non-frame-native-MTP-decode-ABI.patch b/third_party/llama.cpp/patches/0014-Add-non-frame-native-MTP-decode-ABI.patch index cf0baa6550..45df7a4c2a 100644 --- a/third_party/llama.cpp/patches/0014-Add-non-frame-native-MTP-decode-ABI.patch +++ b/third_party/llama.cpp/patches/0014-Add-non-frame-native-MTP-decode-ABI.patch @@ -1,7 +1,7 @@ -From defe90c1c6d5152169107ee18cfa76303528f330 Mon Sep 17 00:00:00 2001 +From bc10f74e32c88213d092b45b46b24502a0d397cb Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Fri, 3 Jul 2026 15:03:17 -0400 -Subject: [PATCH 14/16] Add non-frame native MTP decode ABI +Subject: [PATCH 14/45] Add non-frame native MTP decode ABI --- include/skippy.h | 9 +++++++++ @@ -92,5 +92,5 @@ index db17f7d2..5f5596ce 100644 struct skippy_session * const * sessions, const llama_token * token_ids, -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0015-Fix-stage-activation-graph-input-allocation.patch b/third_party/llama.cpp/patches/0015-Fix-stage-activation-graph-input-allocation.patch index 2a712a94d6..308a81b550 100644 --- a/third_party/llama.cpp/patches/0015-Fix-stage-activation-graph-input-allocation.patch +++ b/third_party/llama.cpp/patches/0015-Fix-stage-activation-graph-input-allocation.patch @@ -1,7 +1,7 @@ -From d42eefb175b9e2ce7445391fa6fd059dc4521a13 Mon Sep 17 00:00:00 2001 +From 76ff9a9cdd614c8bf4a082d06b9d144bc6096dd4 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Wed, 15 Jul 2026 07:54:20 +1000 -Subject: [PATCH 15/16] Fix stage activation graph input allocation +Subject: [PATCH 15/45] Fix stage activation graph input allocation --- src/llama-graph.cpp | 39 +++++++++++++++++++++++++++------------ @@ -102,5 +102,5 @@ index a75f0b75..6613747a 100644 if (n_embd_inp != n_embd) { cur = ggml_view_2d(ctx0, cur, n_embd, n_tokens, cur->nb[1], 0); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0016-Recognize-thinking-field-in-chat-auto-parser.patch b/third_party/llama.cpp/patches/0016-Recognize-thinking-field-in-chat-auto-parser.patch index ac46410881..7a87d97086 100644 --- a/third_party/llama.cpp/patches/0016-Recognize-thinking-field-in-chat-auto-parser.patch +++ b/third_party/llama.cpp/patches/0016-Recognize-thinking-field-in-chat-auto-parser.patch @@ -1,7 +1,7 @@ -From 4a64533c4fc75234c0d7aaf9a453a90ee43f25fa Mon Sep 17 00:00:00 2001 +From 587f96c77401abfb93f56147b306c2f5641f2a54 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Tue, 14 Jul 2026 19:32:50 +1000 -Subject: [PATCH 16/16] Recognize thinking field in chat auto-parser +Subject: [PATCH 16/45] Recognize thinking field in chat auto-parser --- common/chat-diff-analyzer.cpp | 10 +++++++--- @@ -9,10 +9,10 @@ Subject: [PATCH 16/16] Recognize thinking field in chat auto-parser 2 files changed, 28 insertions(+), 3 deletions(-) diff --git a/common/chat-diff-analyzer.cpp b/common/chat-diff-analyzer.cpp -index 127278df..bc9abb70 100644 +index 7db1dcb0..e1bd5775 100644 --- a/common/chat-diff-analyzer.cpp +++ b/common/chat-diff-analyzer.cpp -@@ -456,7 +456,8 @@ void analyze_reasoning::compare_reasoning_presence() { +@@ -476,7 +476,8 @@ void analyze_reasoning::compare_reasoning_presence() { json assistant_with_reasoning = json{ { "role", "assistant" }, { "content", ASSISTANT_MSG }, @@ -22,7 +22,7 @@ index 127278df..bc9abb70 100644 }; template_params params; -@@ -589,13 +590,15 @@ void analyze_reasoning::compare_reasoning_scope() { +@@ -609,13 +610,15 @@ void analyze_reasoning::compare_reasoning_scope() { json assistant_reasoning_content = json{ { "role", "assistant" }, { "content", ASSISTANT_MSG }, @@ -39,7 +39,7 @@ index 127278df..bc9abb70 100644 { "tool_calls", json::array({ build_tool_call(FUN_FIRST, json{ { ARG_FIRST, "VVVV" }, { ARG_SECOND, "XXXX" } }) }) } }; -@@ -664,7 +667,8 @@ analyze_content::analyze_content(const common_chat_template & tmpl, const analyz +@@ -684,7 +687,8 @@ analyze_content::analyze_content(const common_chat_template & tmpl, const analyz json assistant_with_reasoning = json{ { "role", "assistant" }, { "content", "" }, @@ -50,10 +50,10 @@ index 127278df..bc9abb70 100644 template_params params_content_only; diff --git a/tests/test-chat-auto-parser.cpp b/tests/test-chat-auto-parser.cpp -index 2d854509..312038a3 100644 +index 2cb295c1..3273e66c 100644 --- a/tests/test-chat-auto-parser.cpp +++ b/tests/test-chat-auto-parser.cpp -@@ -61,6 +61,7 @@ static void test_nemotron_lazy_tool_trigger(testing & t); +@@ -70,6 +70,7 @@ static void test_nemotron_lazy_tool_trigger(testing & t); // CohereForAI template analysis tests static void test_cohere_reasoning_detection(testing & t); @@ -61,7 +61,7 @@ index 2d854509..312038a3 100644 static void test_cohere_analysis(testing & t); // SmolLM3 template analysis tests -@@ -1410,6 +1411,26 @@ static common_chat_template load_cohere_template(testing & t) { +@@ -1510,6 +1511,26 @@ static common_chat_template load_cohere_template(testing & t) { static void test_cohere_analysis(testing & t) { t.test("Cohere reasoning detection", test_cohere_reasoning_detection); @@ -89,5 +89,5 @@ index 2d854509..312038a3 100644 static void test_cohere_reasoning_detection(testing & t) { -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0018-Expose-stateful-N-gram-cache-ABI.patch b/third_party/llama.cpp/patches/0017-Expose-stateful-N-gram-cache-ABI.patch similarity index 98% rename from third_party/llama.cpp/patches/0018-Expose-stateful-N-gram-cache-ABI.patch rename to third_party/llama.cpp/patches/0017-Expose-stateful-N-gram-cache-ABI.patch index 39b6d25e83..780dbc8fb4 100644 --- a/third_party/llama.cpp/patches/0018-Expose-stateful-N-gram-cache-ABI.patch +++ b/third_party/llama.cpp/patches/0017-Expose-stateful-N-gram-cache-ABI.patch @@ -1,7 +1,7 @@ -From 6f8336c5b93f9c7888162e36497a8c00d4bcf569 Mon Sep 17 00:00:00 2001 +From c83720f81a6cebdfcba89d761afabb72bb6796f6 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 14:07:01 +1000 -Subject: [PATCH 1/2] Expose stateful N-gram cache ABI +Subject: [PATCH 17/45] Expose stateful N-gram cache ABI --- common/ngram-cache.cpp | 136 ++++++++++++++++++++++++++++++++++++++++ @@ -254,5 +254,5 @@ index 5f5596ce..79e3e73f 100644 void skippy_error_free(struct skippy_error * error) { -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0019-Remove-legacy-session-checkpoint-ABI.patch b/third_party/llama.cpp/patches/0018-Remove-legacy-session-checkpoint-ABI.patch similarity index 98% rename from third_party/llama.cpp/patches/0019-Remove-legacy-session-checkpoint-ABI.patch rename to third_party/llama.cpp/patches/0018-Remove-legacy-session-checkpoint-ABI.patch index f8294ae559..bd91f1403b 100644 --- a/third_party/llama.cpp/patches/0019-Remove-legacy-session-checkpoint-ABI.patch +++ b/third_party/llama.cpp/patches/0018-Remove-legacy-session-checkpoint-ABI.patch @@ -1,7 +1,7 @@ -From b56e6f768f917feef60e2a7578d62091ae03d9bd Mon Sep 17 00:00:00 2001 +From ac55c0b40d0c582878f2f0068315e5889c97f845 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 14:07:56 +1000 -Subject: [PATCH 2/2] Remove legacy session checkpoint ABI +Subject: [PATCH 18/45] Remove legacy session checkpoint ABI --- include/skippy.h | 10 ---- @@ -237,5 +237,5 @@ index 79e3e73f..32fb95ad 100644 struct skippy_model * model, const char * text, -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0020-Re-prime-native-MTP-after-state-restoration.patch b/third_party/llama.cpp/patches/0019-Re-prime-native-MTP-after-state-restoration.patch similarity index 96% rename from third_party/llama.cpp/patches/0020-Re-prime-native-MTP-after-state-restoration.patch rename to third_party/llama.cpp/patches/0019-Re-prime-native-MTP-after-state-restoration.patch index b1bde9c8e2..3d858bab5c 100644 --- a/third_party/llama.cpp/patches/0020-Re-prime-native-MTP-after-state-restoration.patch +++ b/third_party/llama.cpp/patches/0019-Re-prime-native-MTP-after-state-restoration.patch @@ -1,7 +1,7 @@ -From 595a5b64c64f85f6dba0a0a61de1868a8b0d78d6 Mon Sep 17 00:00:00 2001 +From 7462ace1495a306a658712c27034b7a70f936c6d Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Tue, 21 Jul 2026 01:38:34 +1000 -Subject: [PATCH] Re-prime native MTP after state restoration +Subject: [PATCH 19/45] Re-prime native MTP after state restoration --- src/skippy.cpp | 31 ++++++++++++++++++++++++++++++- @@ -106,5 +106,5 @@ index 32fb95ad..7f561a8d 100644 *out_session = session; return skippy_success(out_error); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0021-Fix-N-gram-confidence-threshold-indexing.patch b/third_party/llama.cpp/patches/0020-Fix-N-gram-confidence-threshold-indexing.patch similarity index 95% rename from third_party/llama.cpp/patches/0021-Fix-N-gram-confidence-threshold-indexing.patch rename to third_party/llama.cpp/patches/0020-Fix-N-gram-confidence-threshold-indexing.patch index 0e780d1f6a..5583220831 100644 --- a/third_party/llama.cpp/patches/0021-Fix-N-gram-confidence-threshold-indexing.patch +++ b/third_party/llama.cpp/patches/0020-Fix-N-gram-confidence-threshold-indexing.patch @@ -1,7 +1,7 @@ -From ce5f8e97f7be0a9e3ce2f6a5f97ebc035a749cf1 Mon Sep 17 00:00:00 2001 +From 99a2f74d756dc4b93b650b92609a2ccae1719fd2 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Tue, 21 Jul 2026 02:44:00 +1000 -Subject: [PATCH] Fix N-gram confidence threshold indexing +Subject: [PATCH 20/45] Fix N-gram confidence threshold indexing When the configured N-gram range starts above one, ngrams_primary is a zero-based slice of that range. Indexing the confidence tables directly by @@ -58,5 +58,5 @@ index a0447210..ce5f8e97 100644 if (drafted_token == LLAMA_TOKEN_NULL) { drafted_token = try_draft(nc_static, ngram_static); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0022-ggml-add-GLM-DSA-sparse-execution-primitives.patch b/third_party/llama.cpp/patches/0021-ggml-add-GLM-DSA-sparse-execution-primitives.patch similarity index 99% rename from third_party/llama.cpp/patches/0022-ggml-add-GLM-DSA-sparse-execution-primitives.patch rename to third_party/llama.cpp/patches/0021-ggml-add-GLM-DSA-sparse-execution-primitives.patch index 13c30015b1..a67b1afff4 100644 --- a/third_party/llama.cpp/patches/0022-ggml-add-GLM-DSA-sparse-execution-primitives.patch +++ b/third_party/llama.cpp/patches/0021-ggml-add-GLM-DSA-sparse-execution-primitives.patch @@ -1,7 +1,7 @@ -From dfe419c41613f4cd640998b391899a145abdfc10 Mon Sep 17 00:00:00 2001 +From 8bdf03a6a77a8acb38fd4bb35349a9d05f0f0038 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Tue, 14 Jul 2026 09:15:57 +1000 -Subject: [PATCH 19/24] ggml: add GLM-DSA sparse execution primitives +Subject: [PATCH 21/45] ggml: add GLM-DSA sparse execution primitives --- ggml/include/ggml-rpc.h | 2 +- @@ -5169,7 +5169,7 @@ index b36fa811..03567502 100644 void ggml_metal_buffer_set_tensor (ggml_metal_buffer_t buf, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size); void ggml_metal_buffer_get_tensor (ggml_metal_buffer_t buf, const struct ggml_tensor * tensor, void * data, size_t offset, size_t size); diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m -index 80e47f2c..8becf966 100644 +index 4edd77c6..4531d2d5 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -2,7 +2,6 @@ @@ -5196,7 +5196,7 @@ index 80e47f2c..8becf966 100644 void ggml_metal_encoder_end_encoding(ggml_metal_encoder_t encoder) { [encoder->obj endEncoding]; } -@@ -1272,6 +1280,137 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te +@@ -1273,6 +1281,137 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te return true; case GGML_OP_GATED_DELTA_NET: return has_simdgroup_reduction && op->src[2]->ne[0] % 32 == 0; @@ -5334,7 +5334,7 @@ index 80e47f2c..8becf966 100644 case GGML_OP_SOLVE_TRI: case GGML_OP_MUL_MAT: case GGML_OP_MUL_MAT_ID: -@@ -1411,6 +1550,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te +@@ -1412,6 +1551,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te // pointers to global device ggml_metal_device_t dev; @@ -5342,7 +5342,7 @@ index 80e47f2c..8becf966 100644 }; static void ggml_metal_log_allocated_size(id device, size_t size_aligned) { -@@ -1515,7 +1655,6 @@ ggml_metal_buffer_t ggml_metal_buffer_init(ggml_metal_device_t dev, size_t size, +@@ -1516,7 +1656,6 @@ ggml_metal_buffer_t ggml_metal_buffer_init(ggml_metal_device_t dev, size_t size, ggml_metal_buffer_t res = calloc(1, sizeof(struct ggml_metal_buffer)); res->dev = dev; @@ -5350,7 +5350,7 @@ index 80e47f2c..8becf966 100644 const size_t size_page = sysconf(_SC_PAGESIZE); size_t size_aligned = size; -@@ -1585,7 +1724,6 @@ ggml_metal_buffer_t ggml_metal_buffer_map(ggml_metal_device_t dev, void * ptr, s +@@ -1586,7 +1725,6 @@ ggml_metal_buffer_t ggml_metal_buffer_map(ggml_metal_device_t dev, void * ptr, s ggml_metal_buffer_t res = calloc(1, sizeof(struct ggml_metal_buffer)); res->dev = dev; @@ -5358,7 +5358,7 @@ index 80e47f2c..8becf966 100644 res->all_data = ptr; res->all_size = size; -@@ -1913,3 +2051,11 @@ struct ggml_metal_buffer_id ggml_metal_buffer_get_id(ggml_metal_buffer_t buf, co +@@ -1914,3 +2052,11 @@ struct ggml_metal_buffer_id ggml_metal_buffer_get_id(ggml_metal_buffer_t buf, co return res; } @@ -31466,7 +31466,7 @@ index a7d1fe7d..cdc85b08 100644 struct ggml_hash_set ggml_hash_set_new(size_t size) { diff --git a/src/llama-hparams.h b/src/llama-hparams.h -index 8be5f28f..92626b4d 100644 +index 747754fc..6cca1d33 100644 --- a/src/llama-hparams.h +++ b/src/llama-hparams.h @@ -226,6 +226,10 @@ struct llama_hparams { @@ -31478,7 +31478,8 @@ index 8be5f28f..92626b4d 100644 + bool indexer_types_present = false; + std::array indexer_types; // -1 unknown, 0 shared, 1 full - // DeepSeek-V4 - uint32_t dsv4_o_group_count = 0; + // Indexer is "full" (1) or "shared" (0) + // Shared indexers reuse top-k from previous full layer -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/patches/0023-skippy-expose-GLM-DSA-staged-runtime-controls.patch b/third_party/llama.cpp/patches/0022-skippy-expose-GLM-DSA-staged-runtime-controls.patch similarity index 96% rename from third_party/llama.cpp/patches/0023-skippy-expose-GLM-DSA-staged-runtime-controls.patch rename to third_party/llama.cpp/patches/0022-skippy-expose-GLM-DSA-staged-runtime-controls.patch index 1ed508cf72..88487057d3 100644 --- a/third_party/llama.cpp/patches/0023-skippy-expose-GLM-DSA-staged-runtime-controls.patch +++ b/third_party/llama.cpp/patches/0022-skippy-expose-GLM-DSA-staged-runtime-controls.patch @@ -1,17 +1,17 @@ -From e87de13ba5a2bb71fb5b1b874ca61d7338ec3f62 Mon Sep 17 00:00:00 2001 +From 02b458479060c5ae055c4a9973a66884ed55c6e9 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Tue, 14 Jul 2026 09:16:19 +1000 -Subject: [PATCH 20/24] skippy: expose GLM-DSA staged runtime controls +Subject: [PATCH 22/45] skippy: expose GLM-DSA staged runtime controls --- include/skippy.h | 20 + src/llama-graph.cpp | 45 + src/llama-graph.h | 43 + - src/skippy.cpp | 2006 ++++++++++++++++++++++++++++++++++++++++--- - 4 files changed, 1986 insertions(+), 128 deletions(-) + src/skippy.cpp | 2015 ++++++++++++++++++++++++++++++++++++++++--- + 4 files changed, 1989 insertions(+), 134 deletions(-) diff --git a/include/skippy.h b/include/skippy.h -index c9ecc349..cd0fe21a 100644 +index 50c08515..8f166e9c 100644 --- a/include/skippy.h +++ b/include/skippy.h @@ -43,6 +43,15 @@ enum skippy_activation_layout { @@ -235,7 +235,7 @@ index 2e0d9c3b..ab6e111e 100644 std::vector t_layer_inp; diff --git a/src/skippy.cpp b/src/skippy.cpp -index 46da7007..7e11d92b 100644 +index 7f561a8d..a8cc03cb 100644 --- a/src/skippy.cpp +++ b/src/skippy.cpp @@ -9,6 +9,7 @@ @@ -334,7 +334,7 @@ index 46da7007..7e11d92b 100644 bool executable = true; uint32_t lane_count = 1; std::vector lane_in_use; -@@ -936,14 +1005,27 @@ struct skippy_filter_scope { +@@ -932,14 +1001,27 @@ struct skippy_filter_scope { }; struct skippy_graph_filter_scope { @@ -369,7 +369,7 @@ index 46da7007..7e11d92b 100644 skippy_graph_set_filter(filter); enabled = true; } -@@ -1070,6 +1152,55 @@ struct skippy_gemma3n_altup_scope { +@@ -1066,6 +1148,55 @@ struct skippy_gemma3n_altup_scope { bool enabled = false; }; @@ -425,7 +425,7 @@ index 46da7007..7e11d92b 100644 static bool skippy_is_filtered(const skippy_session * session) { return session != nullptr && session->stage_model != nullptr && -@@ -1095,6 +1226,124 @@ static bool skippy_is_gemma3n_activation_model(const skippy_session * session) { +@@ -1091,6 +1222,124 @@ static bool skippy_is_gemma3n_activation_model(const skippy_session * session) { return session->stage_model->model->arch == LLM_ARCH_GEMMA3N; } @@ -550,7 +550,7 @@ index 46da7007..7e11d92b 100644 static size_t skippy_activation_hidden_bytes(const skippy_session * session, size_t token_count) { if (session == nullptr || session->stage_model == nullptr || session->stage_model->model == nullptr) { return 0; -@@ -1105,6 +1354,113 @@ static size_t skippy_activation_hidden_bytes(const skippy_session * session, siz +@@ -1101,6 +1350,113 @@ static size_t skippy_activation_hidden_bytes(const skippy_session * session, siz sizeof(float); } @@ -664,7 +664,7 @@ index 46da7007..7e11d92b 100644 static size_t skippy_gemma3n_altup_bytes(const skippy_session * session, size_t token_count) { if (session == nullptr || session->stage_model == nullptr || session->stage_model->model == nullptr) { return 0; -@@ -1122,6 +1478,11 @@ static uint64_t skippy_output_activation_flags( +@@ -1118,6 +1474,11 @@ static uint64_t skippy_output_activation_flags( if (skippy_emits_activation_frame(session) && skippy_is_gemma3n_activation_model(session)) { return SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP; } @@ -676,7 +676,7 @@ index 46da7007..7e11d92b 100644 if (!skippy_emits_activation_frame(session) || !skippy_is_rwkv7_activation_model(session)) { return 0; } -@@ -1145,6 +1506,9 @@ static size_t skippy_activation_payload_bytes( +@@ -1141,6 +1502,9 @@ static size_t skippy_activation_payload_bytes( return skippy_gemma3n_altup_bytes(session, token_count); } size_t payload_bytes = hidden_bytes; @@ -686,7 +686,7 @@ index 46da7007..7e11d92b 100644 if ((flags & SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST) != 0) { payload_bytes += hidden_bytes; } -@@ -1155,11 +1519,12 @@ static bool skippy_has_activation_payload(const skippy_activation_desc * desc, c +@@ -1151,11 +1515,12 @@ static bool skippy_has_activation_payload(const skippy_activation_desc * desc, c return desc != nullptr && desc->payload_bytes > 0 && payload != nullptr; } @@ -700,23 +700,22 @@ index 46da7007..7e11d92b 100644 struct skippy_error ** out_error) { if (session == nullptr || session->stage_model == nullptr) { skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "session is required"); -@@ -1185,7 +1550,7 @@ static enum skippy_status skippy_validate_frame_input( +@@ -1181,140 +1546,1186 @@ static enum skippy_status skippy_validate_frame_input( skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame must be version 1 F32 token-major"); return SKIPPY_STATUS_INVALID_ARGUMENT; } - if (input_desc->sequence_count != 1 || input_desc->token_count != expected_token_count) { -+ if (input_desc->sequence_count != expected_sequence_count || input_desc->token_count != expected_token_count) { - skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame token or sequence count does not match request"); - return SKIPPY_STATUS_INVALID_ARGUMENT; - } -@@ -1194,131 +1559,1177 @@ static enum skippy_status skippy_validate_frame_input( - return SKIPPY_STATUS_INVALID_ARGUMENT; - } - -- const uint64_t supported_flags = SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST | SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP; -- if ((input_desc->flags & ~supported_flags) != 0) { -- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame has unsupported sideband flags"); +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame token or sequence count does not match request"); - return SKIPPY_STATUS_INVALID_ARGUMENT; ++ if (input_desc->sequence_count != expected_sequence_count || input_desc->token_count != expected_token_count) { ++ skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame token or sequence count does not match request"); ++ return SKIPPY_STATUS_INVALID_ARGUMENT; ++ } ++ if (input_desc->layer_end != config.layer_start) { ++ skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame layer_end must match this stage layer_start"); ++ return SKIPPY_STATUS_INVALID_ARGUMENT; ++ } ++ + const uint64_t supported_flags = + SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST | + SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP | @@ -1500,8 +1499,8 @@ index 46da7007..7e11d92b 100644 + hash ^= static_cast(data[i]); + hash *= 1099511628211ull; } -- if ((input_desc->flags & SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP) != 0 && !skippy_is_gemma3n_activation_model(session)) { -- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "Gemma3n AltUp activation payload is only valid for Gemma3n stages"); +- if (input_desc->layer_end != config.layer_start) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame layer_end must match this stage layer_start"); - return SKIPPY_STATUS_INVALID_ARGUMENT; + return hash; +} @@ -1514,9 +1513,6 @@ index 46da7007..7e11d92b 100644 + sum += value; + max_abs = std::max(max_abs, std::fabs(value)); } -- if (skippy_is_gemma3n_activation_model(session) && (input_desc->flags & SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP) == 0) { -- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "non-first Gemma3n runtime slices require AltUp activation payload"); -- return SKIPPY_STATUS_INVALID_ARGUMENT; + const double mean = count == 0 ? 0.0 : sum / static_cast(count); + line += format( + " count=%zu sum=%.17g mean=%.17g max_abs=%.17g", @@ -1525,15 +1521,18 @@ index 46da7007..7e11d92b 100644 + mean, + max_abs); +} -+ + +- const uint64_t supported_flags = SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST | SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP; +- if ((input_desc->flags & ~supported_flags) != 0) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame has unsupported sideband flags"); +- return SKIPPY_STATUS_INVALID_ARGUMENT; +static void skippy_glm_dsa_trace_tensor_stats(std::string & line, const ggml_tensor * tensor, uint32_t max_bytes) { + if (!skippy_glm_dsa_tensor_has_buffer(tensor)) { + line += " stats=unallocated"; + return; } -- if ((input_desc->flags & SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP) != 0 && -- (input_desc->flags & SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST) != 0) { -- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "Gemma3n AltUp and RWKV7 v_first activation flags cannot be combined"); +- if ((input_desc->flags & SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP) != 0 && !skippy_is_gemma3n_activation_model(session)) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "Gemma3n AltUp activation payload is only valid for Gemma3n stages"); - return SKIPPY_STATUS_INVALID_ARGUMENT; + + const size_t nbytes = ggml_nbytes(tensor); @@ -1541,8 +1540,8 @@ index 46da7007..7e11d92b 100644 + line += format(" stats=skipped:nbytes>%u", max_bytes); + return; } -- if ((input_desc->flags & SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST) != 0 && !skippy_is_rwkv7_activation_model(session)) { -- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "RWKV7 v_first sideband is only valid for RWKV7 stages"); +- if (skippy_is_gemma3n_activation_model(session) && (input_desc->flags & SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP) == 0) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "non-first Gemma3n runtime slices require AltUp activation payload"); - return SKIPPY_STATUS_INVALID_ARGUMENT; + + std::vector raw(nbytes); @@ -1587,8 +1586,9 @@ index 46da7007..7e11d92b 100644 + default: + break; } -- if (skippy_is_rwkv7_activation_model(session) && (input_desc->flags & SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST) == 0) { -- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "non-first RWKV7 runtime slices require v_first activation sideband"); +- if ((input_desc->flags & SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP) != 0 && +- (input_desc->flags & SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST) != 0) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "Gemma3n AltUp and RWKV7 v_first activation flags cannot be combined"); - return SKIPPY_STATUS_INVALID_ARGUMENT; +} + @@ -1600,17 +1600,19 @@ index 46da7007..7e11d92b 100644 + !skippy_glm_dsa_tensor_has_buffer(tensor)) { + return; } - -- const size_t expected_bytes = skippy_activation_payload_bytes(session, expected_token_count, input_desc->flags); -- if (input_desc->payload_bytes != expected_bytes) { -- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame payload size does not match model hidden size"); +- if ((input_desc->flags & SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST) != 0 && !skippy_is_rwkv7_activation_model(session)) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "RWKV7 v_first sideband is only valid for RWKV7 stages"); - return SKIPPY_STATUS_INVALID_ARGUMENT; ++ + std::string filename = tensor->name; + for (char & c : filename) { + if (!std::isalnum(static_cast(c)) && c != '-' && c != '_') { + c = '_'; + } } +- if (skippy_is_rwkv7_activation_model(session) && (input_desc->flags & SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST) == 0) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "non-first RWKV7 runtime slices require v_first activation sideband"); +- return SKIPPY_STATUS_INVALID_ARGUMENT; + filename += format( + "_%lld_%lld_%lld_%lld_%s.bin", + static_cast(tensor->ne[0]), @@ -1618,29 +1620,20 @@ index 46da7007..7e11d92b 100644 + static_cast(tensor->ne[2]), + static_cast(tensor->ne[3]), + ggml_type_name(tensor->type)); - -- return SKIPPY_STATUS_OK; ++ + const size_t nbytes = ggml_nbytes(tensor); + std::vector raw(nbytes); + ggml_backend_tensor_get(tensor, raw.data(), 0, nbytes); + std::ofstream output(std::string(dump_dir) + "/" + filename, std::ios::binary | std::ios::trunc); + if (output) { + output.write(reinterpret_cast(raw.data()), static_cast(raw.size())); -+ } - } + } ++} --static enum skippy_status skippy_prepare_output_activation_frame( -- skippy_session * session, -- size_t token_count, -- void * output_payload, -- size_t output_payload_capacity, -- size_t * out_output_payload_bytes, -- skippy_activation_desc * output_desc, -- const skippy_activation_desc * input_desc, -- struct skippy_error ** out_error) { -- const uint64_t output_flags = skippy_output_activation_flags(session, input_desc); -- const size_t payload_bytes = skippy_emits_activation_frame(session) ? -- skippy_activation_payload_bytes(session, token_count, output_flags) : 0; +- const size_t expected_bytes = skippy_activation_payload_bytes(session, expected_token_count, input_desc->flags); +- if (input_desc->payload_bytes != expected_bytes) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "activation frame payload size does not match model hidden size"); +- return SKIPPY_STATUS_INVALID_ARGUMENT; +static void skippy_glm_dsa_trace_tensor( + skippy_glm_dsa_op_timing & timing, + skippy_glm_dsa_op_kind kind, @@ -1649,10 +1642,10 @@ index 46da7007..7e11d92b 100644 + timing.trace_nodes >= timing.trace_node_limit || + !skippy_glm_dsa_trace_filter_matches(timing.trace_filter, tensor->name)) { + return; -+ } + } -- if (out_output_payload_bytes != nullptr) { -- *out_output_payload_bytes = payload_bytes; +- return SKIPPY_STATUS_OK; +-} + timing.trace_nodes += 1; + skippy_glm_dsa_dump_trace_tensor(timing, tensor); + std::string line = format( @@ -1683,6 +1676,35 @@ index 46da7007..7e11d92b 100644 +static bool skippy_glm_dsa_op_timing_cb(ggml_tensor * tensor, bool ask, void * user_data) { + skippy_model * model = static_cast(user_data); + if (model == nullptr || !model->glm_dsa_timing.enabled || tensor == nullptr) { ++ return false; ++ } + +-static enum skippy_status skippy_prepare_output_activation_frame( +- skippy_session * session, +- size_t token_count, +- void * output_payload, +- size_t output_payload_capacity, +- size_t * out_output_payload_bytes, +- skippy_activation_desc * output_desc, +- const skippy_activation_desc * input_desc, +- struct skippy_error ** out_error) { +- const uint64_t output_flags = skippy_output_activation_flags(session, input_desc); +- const size_t payload_bytes = skippy_emits_activation_frame(session) ? +- skippy_activation_payload_bytes(session, token_count, output_flags) : 0; ++ skippy_glm_dsa_op_timing & timing = model->glm_dsa_timing; ++ const skippy_glm_dsa_op_kind kind = skippy_glm_dsa_op_kind_for_tensor(tensor->name); ++ const bool trace_match = timing.trace_tensors && ++ skippy_glm_dsa_trace_filter_matches(timing.trace_filter, tensor->name); ++ if (kind == SKIPPY_GLM_DSA_OP_UNKNOWN && !trace_match) { ++ return false; ++ } + +- if (out_output_payload_bytes != nullptr) { +- *out_output_payload_bytes = payload_bytes; ++ // Tensor tracing can be used as a narrow synchronization probe. Keep the ++ // normal op-timing callback broad, but only split the graph at tensors the ++ // trace filter actually requests when timing is disabled. ++ if (timing.trace_tensors && !timing.print_timing && !trace_match) { + return false; } @@ -1698,12 +1720,10 @@ index 46da7007..7e11d92b 100644 - } else if (output_payload_capacity > 0 && output_payload == nullptr) { - skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "output_payload is required when output capacity is non-zero"); - return SKIPPY_STATUS_INVALID_ARGUMENT; -+ skippy_glm_dsa_op_timing & timing = model->glm_dsa_timing; -+ const skippy_glm_dsa_op_kind kind = skippy_glm_dsa_op_kind_for_tensor(tensor->name); -+ const bool trace_match = timing.trace_tensors && -+ skippy_glm_dsa_trace_filter_matches(timing.trace_filter, tensor->name); -+ if (kind == SKIPPY_GLM_DSA_OP_UNKNOWN && !trace_match) { -+ return false; ++ if (ask) { ++ timing.pending_kind = kind; ++ timing.pending_start_us = ggml_time_us(); ++ return true; } - if (output_desc != nullptr) { @@ -1718,29 +1738,6 @@ index 46da7007..7e11d92b 100644 - output_desc->sequence_count = token_count > 0 ? 1 : 0; - output_desc->payload_bytes = payload_bytes; - output_desc->flags = output_flags; -+ // Tensor tracing can be used as a narrow synchronization probe. Keep the -+ // normal op-timing callback broad, but only split the graph at tensors the -+ // trace filter actually requests when timing is disabled. -+ if (timing.trace_tensors && !timing.print_timing && !trace_match) { -+ return false; - } - -- return SKIPPY_STATUS_OK; --} -+ if (ask) { -+ timing.pending_kind = kind; -+ timing.pending_start_us = ggml_time_us(); -+ return true; -+ } - --static enum skippy_status skippy_decode_batch( -- skippy_session * session, -- llama_batch batch, -- size_t token_count, -- struct skippy_error ** out_error) { -- if (session == nullptr || session->ctx == nullptr || token_count == 0) { -- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "session and at least one token are required"); -- return SKIPPY_STATUS_INVALID_ARGUMENT; + const int64_t elapsed_us = std::max(0, ggml_time_us() - timing.pending_start_us); + if (kind != SKIPPY_GLM_DSA_OP_UNKNOWN) { + timing.stats[kind].nodes += 1; @@ -1755,11 +1752,7 @@ index 46da7007..7e11d92b 100644 + return true; +} -- skippy_graph_filter_scope graph_filter_scope(&session->stage_model->config); -- const int32_t rc = llama_decode(session->ctx, batch); -- if (rc != 0) { -- skippy_set_error(out_error, SKIPPY_STATUS_RUNTIME_ERROR, "llama_decode failed"); -- return SKIPPY_STATUS_RUNTIME_ERROR; +- return SKIPPY_STATUS_OK; +static skippy_glm_dsa_op_stat skippy_glm_dsa_sparse_mask_stat_from_stats(const skippy_glm_dsa_op_stat stats[]) { + skippy_glm_dsa_op_stat stat = {}; + stat.nodes += stats[SKIPPY_GLM_DSA_OP_SPARSE_MASK_FILL].nodes; @@ -1769,12 +1762,26 @@ index 46da7007..7e11d92b 100644 + stat.elapsed_us += stats[SKIPPY_GLM_DSA_OP_SPARSE_MASK_TOPK].elapsed_us; + stat.elapsed_us += stats[SKIPPY_GLM_DSA_OP_SPARSE_MASK_ADD].elapsed_us; + return stat; -+} -+ + } + +-static enum skippy_status skippy_decode_batch( +- skippy_session * session, +- llama_batch batch, +- size_t token_count, +- struct skippy_error ** out_error) { +- if (session == nullptr || session->ctx == nullptr || token_count == 0) { +- skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "session and at least one token are required"); +- return SKIPPY_STATUS_INVALID_ARGUMENT; +- } +static skippy_glm_dsa_op_stat skippy_glm_dsa_sparse_mask_stat(const skippy_glm_dsa_op_timing & timing) { + return skippy_glm_dsa_sparse_mask_stat_from_stats(timing.stats); +} -+ + +- skippy_graph_filter_scope graph_filter_scope(&session->stage_model->config); +- const int32_t rc = llama_decode(session->ctx, batch); +- if (rc != 0) { +- skippy_set_error(out_error, SKIPPY_STATUS_RUNTIME_ERROR, "llama_decode failed"); +- return SKIPPY_STATUS_RUNTIME_ERROR; +static skippy_glm_dsa_op_stat skippy_glm_dsa_indexer_topk_stat_from_stats(const skippy_glm_dsa_op_stat stats[]) { + skippy_glm_dsa_op_stat stat = {}; + stat.nodes += stats[SKIPPY_GLM_DSA_OP_INDEXER].nodes; @@ -1982,7 +1989,7 @@ index 46da7007..7e11d92b 100644 } static void skippy_mtp_clear_session_state(skippy_session * session) { -@@ -1460,6 +2871,7 @@ static enum skippy_status skippy_decode_tokens( +@@ -1477,6 +2888,7 @@ static enum skippy_status skippy_decode_tokens( const llama_token * token_ids, size_t token_count, bool request_logits, @@ -1990,7 +1997,7 @@ index 46da7007..7e11d92b 100644 struct skippy_error ** out_error) { (void) request_logits; -@@ -1493,7 +2905,7 @@ static enum skippy_status skippy_decode_tokens( +@@ -1510,7 +2922,7 @@ static enum skippy_status skippy_decode_tokens( /*seq_id =*/ &seq_ids, /*logits =*/ &logits, }; @@ -1999,7 +2006,7 @@ index 46da7007..7e11d92b 100644 if (status == SKIPPY_STATUS_OK) { status = skippy_mtp_sync_target_tokens(session, token_ids, token_count, pos, out_error); } -@@ -1514,7 +2929,7 @@ static enum skippy_status skippy_decode_tokens( +@@ -1531,7 +2943,7 @@ static enum skippy_status skippy_decode_tokens( batch.logits[i] = request_logits && i == n_tokens - 1 ? 1 : 0; } @@ -2008,7 +2015,7 @@ index 46da7007..7e11d92b 100644 llama_batch_free(batch); if (status == SKIPPY_STATUS_OK) { status = skippy_mtp_sync_target_tokens(session, token_ids, token_count, token_start, out_error); -@@ -1554,7 +2969,12 @@ static enum skippy_status skippy_verify_token_batch( +@@ -1571,7 +2983,12 @@ static enum skippy_status skippy_verify_token_batch( batch.logits[i] = 1; } @@ -2022,7 +2029,7 @@ index 46da7007..7e11d92b 100644 llama_batch_free(batch); return status; } -@@ -2201,6 +3621,8 @@ static enum skippy_status skippy_copy_output_activation_frame( +@@ -2221,6 +3638,8 @@ static enum skippy_status skippy_copy_output_activation_frame( void * output_payload, const skippy_activation_desc * input_desc, const void * input_payload, @@ -2031,7 +2038,7 @@ index 46da7007..7e11d92b 100644 struct skippy_error ** out_error) { if (!skippy_emits_activation_frame(session)) { return skippy_success(out_error); -@@ -2208,6 +3630,7 @@ static enum skippy_status skippy_copy_output_activation_frame( +@@ -2228,6 +3647,7 @@ static enum skippy_status skippy_copy_output_activation_frame( const uint64_t output_flags = skippy_output_activation_flags(session, input_desc); const size_t hidden_bytes = skippy_activation_hidden_bytes(session, token_count); @@ -2039,7 +2046,7 @@ index 46da7007..7e11d92b 100644 if ((output_flags & SKIPPY_ACTIVATION_FLAG_GEMMA3N_ALTUP) != 0) { llm_graph_result * res = session->ctx->get_gf_res_prev(); ggml_tensor * altup = res != nullptr ? res->get_skippy_gemma3n_altup() : nullptr; -@@ -2221,6 +3644,12 @@ static enum skippy_status skippy_copy_output_activation_frame( +@@ -2241,6 +3661,12 @@ static enum skippy_status skippy_copy_output_activation_frame( return SKIPPY_STATUS_RUNTIME_ERROR; } ggml_backend_tensor_get(altup, output_payload, 0, altup_bytes); @@ -2052,7 +2059,7 @@ index 46da7007..7e11d92b 100644 return skippy_success(out_error); } -@@ -2232,6 +3661,66 @@ static enum skippy_status skippy_copy_output_activation_frame( +@@ -2252,6 +3678,66 @@ static enum skippy_status skippy_copy_output_activation_frame( std::memcpy(output_payload, embeddings, hidden_bytes); @@ -2119,7 +2126,7 @@ index 46da7007..7e11d92b 100644 if ((output_flags & SKIPPY_ACTIVATION_FLAG_RWKV7_V_FIRST) != 0) { uint8_t * sideband_output = static_cast(output_payload) + hidden_bytes; const skippy_runtime_config & config = session->stage_model->config; -@@ -2257,6 +3746,14 @@ static enum skippy_status skippy_copy_output_activation_frame( +@@ -2277,6 +3763,14 @@ static enum skippy_status skippy_copy_output_activation_frame( } std::memcpy(sideband_output, static_cast(input_payload) + hidden_bytes, hidden_bytes); } @@ -2134,7 +2141,7 @@ index 46da7007..7e11d92b 100644 } return skippy_success(out_error); } -@@ -2270,6 +3767,7 @@ static enum skippy_status skippy_decode_activation_frame( +@@ -2290,6 +3784,7 @@ static enum skippy_status skippy_decode_activation_frame( size_t position_count, size_t token_count, bool request_logits, @@ -2142,7 +2149,7 @@ index 46da7007..7e11d92b 100644 struct skippy_error ** out_error) { if (session == nullptr || session->ctx == nullptr || input_desc == nullptr || input_payload == nullptr || token_count == 0) { skippy_set_error(out_error, SKIPPY_STATUS_INVALID_ARGUMENT, "session and activation frame input are required"); -@@ -2350,7 +3848,19 @@ static enum skippy_status skippy_decode_activation_frame( +@@ -2370,7 +3865,19 @@ static enum skippy_status skippy_decode_activation_frame( skippy_rwkv7_v_first_scope rwkv7_v_first_scope(input_desc, input_payload, hidden_bytes, n_embd); skippy_gemma3n_altup_scope gemma3n_altup_scope(input_desc, input_payload, n_embd, n_altup); const llama_pos token_start = pos_storage.empty() ? session->n_past : pos_storage[0]; @@ -2163,7 +2170,7 @@ index 46da7007..7e11d92b 100644 if (status == SKIPPY_STATUS_OK && token_ids != nullptr) { status = skippy_mtp_sync_target_tokens(session, token_ids, token_count, token_start, out_error); } -@@ -2416,7 +3926,19 @@ static enum skippy_status skippy_verify_activation_frame( +@@ -2436,7 +3943,19 @@ static enum skippy_status skippy_verify_activation_frame( skippy_rwkv7_v_first_scope rwkv7_v_first_scope(input_desc, input_payload, hidden_bytes, n_embd); skippy_gemma3n_altup_scope gemma3n_altup_scope(input_desc, input_payload, n_embd, n_altup); const llama_pos token_start = session->n_past; @@ -2184,7 +2191,7 @@ index 46da7007..7e11d92b 100644 if (!alias_input_payload) { llama_batch_free(batch); } -@@ -2719,8 +4241,28 @@ static enum skippy_status skippy_finish_model_open( +@@ -2737,8 +4256,28 @@ static enum skippy_status skippy_finish_model_open( params.type_v = config != nullptr && config->cache_type_v > 0 ? static_cast(config->cache_type_v) : GGML_TYPE_F16; params.flash_attn_type = config != nullptr ? static_cast(config->flash_attn_type) : LLAMA_FLASH_ATTN_TYPE_AUTO; params.embeddings = config != nullptr && config->filter_tensors_on_load && !config->include_output; @@ -2214,7 +2221,7 @@ index 46da7007..7e11d92b 100644 params.n_rs_seq = std::max(params.n_rs_seq, 2); params.kv_unified = true; } -@@ -3414,7 +4956,13 @@ enum skippy_status skippy_prefill_chunk( +@@ -3423,7 +4962,13 @@ enum skippy_status skippy_prefill_chunk( *out_output_activation_bytes = 0; } @@ -2229,7 +2236,7 @@ index 46da7007..7e11d92b 100644 } enum skippy_status skippy_decode_step_sampled( -@@ -3437,7 +4985,13 @@ enum skippy_status skippy_decode_step_sampled( +@@ -3446,7 +4991,13 @@ enum skippy_status skippy_decode_step_sampled( *out_output_activation_bytes = 0; } @@ -2244,7 +2251,7 @@ index 46da7007..7e11d92b 100644 if (status != SKIPPY_STATUS_OK) { return status; } -@@ -3540,7 +5094,9 @@ enum skippy_status skippy_decode_batch_sampled( +@@ -3549,7 +5100,9 @@ enum skippy_status skippy_decode_batch_sampled( batch.logits[i] = 1; } @@ -2255,7 +2262,7 @@ index 46da7007..7e11d92b 100644 const int32_t rc = llama_decode(first->ctx, batch); if (rc != 0) { llama_batch_free(batch); -@@ -3651,9 +5207,25 @@ static enum skippy_status skippy_prefill_chunk_frame_impl( +@@ -3660,9 +5213,25 @@ static enum skippy_status skippy_prefill_chunk_frame_impl( } if (skippy_is_filtered(session) && session->stage_model->config.layer_start > 0) { @@ -2283,7 +2290,7 @@ index 46da7007..7e11d92b 100644 } if (status != SKIPPY_STATUS_OK) { return status; -@@ -3664,7 +5236,8 @@ static enum skippy_status skippy_prefill_chunk_frame_impl( +@@ -3673,7 +5242,8 @@ static enum skippy_status skippy_prefill_chunk_frame_impl( skippy_sample_token(session, sampling) : -1; } @@ -2293,7 +2300,7 @@ index 46da7007..7e11d92b 100644 } enum skippy_status skippy_prefill_chunk_frame( -@@ -3827,9 +5400,25 @@ enum skippy_status skippy_decode_step_frame_sampled( +@@ -3836,9 +5406,25 @@ enum skippy_status skippy_decode_step_frame_sampled( } if (skippy_is_filtered(session) && session->stage_model->config.layer_start > 0) { @@ -2321,7 +2328,7 @@ index 46da7007..7e11d92b 100644 } if (status != SKIPPY_STATUS_OK) { return status; -@@ -3839,7 +5428,8 @@ enum skippy_status skippy_decode_step_frame_sampled( +@@ -3848,7 +5434,8 @@ enum skippy_status skippy_decode_step_frame_sampled( *out_predicted_token = session->stage_model->config.include_output ? skippy_sample_token(session, sampling) : -1; } @@ -2331,7 +2338,7 @@ index 46da7007..7e11d92b 100644 } enum skippy_status skippy_decode_step_frame_sampled_mtp( -@@ -3932,6 +5522,9 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( +@@ -3941,6 +5528,9 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( } const size_t hidden_bytes_per_request = skippy_activation_hidden_bytes(first, 1); @@ -2341,7 +2348,7 @@ index 46da7007..7e11d92b 100644 for (int32_t i = 0; i < n_tokens; ++i) { skippy_session * session = sessions[i]; if (session == nullptr || session->ctx != first->ctx || session->stage_model != first->stage_model) { -@@ -3944,14 +5537,30 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( +@@ -3953,14 +5543,30 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( if (status != SKIPPY_STATUS_OK) { return status; } @@ -2376,7 +2383,7 @@ index 46da7007..7e11d92b 100644 status = skippy_prepare_output_activation_frame( session, 1, -@@ -3967,9 +5576,13 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( +@@ -3976,9 +5582,13 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( } std::vector embd_storage; @@ -2390,7 +2397,7 @@ index 46da7007..7e11d92b 100644 for (int32_t i = 0; i < n_tokens; ++i) { const skippy_activation_desc * input_desc = input_descs[i]; const void * input_payload = input_payloads[i]; -@@ -3982,6 +5595,14 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( +@@ -3991,6 +5601,14 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( std::memcpy(dst, input_payload, static_cast(n_embd)*sizeof(float)); std::memset(dst + n_embd, 0, static_cast(n_embd_inp - n_embd)*sizeof(float)); } @@ -2405,7 +2412,7 @@ index 46da7007..7e11d92b 100644 } } else { token_storage.assign(token_ids, token_ids + request_count); -@@ -4019,7 +5640,16 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( +@@ -4028,7 +5646,16 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( }; skippy_activation_tokens_scope activation_tokens_scope(token_ids, request_count); @@ -2423,7 +2430,7 @@ index 46da7007..7e11d92b 100644 const int32_t rc = llama_decode(first->ctx, batch); if (rc != 0) { skippy_set_error(out_error, SKIPPY_STATUS_RUNTIME_ERROR, "llama_decode failed"); -@@ -4035,6 +5665,75 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( +@@ -4044,6 +5671,75 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( } } @@ -2499,7 +2506,7 @@ index 46da7007..7e11d92b 100644 for (int32_t i = 0; i < n_tokens; ++i) { skippy_session * session = sessions[i]; session->n_past += 1; -@@ -4054,6 +5753,23 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( +@@ -4063,6 +5759,23 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( output_payloads[i], embeddings + static_cast(i)*n_embd, hidden_bytes_per_request); @@ -2523,7 +2530,7 @@ index 46da7007..7e11d92b 100644 } } -@@ -4129,17 +5845,34 @@ enum skippy_status skippy_verify_tokens_frame_sampled( +@@ -4138,17 +5851,34 @@ enum skippy_status skippy_verify_tokens_frame_sampled( if (skippy_is_filtered(session) && session->stage_model->config.layer_start > 0) { status = session->stage_model->config.include_output ? skippy_verify_activation_frame(session, input_desc, input_payload, token_ids, token_count, out_error) : @@ -2561,7 +2568,7 @@ index 46da7007..7e11d92b 100644 if (status != SKIPPY_STATUS_OK) { return status; } -@@ -4198,7 +5931,8 @@ enum skippy_status skippy_session_copy_output_activation_frame( +@@ -4207,7 +5937,8 @@ enum skippy_status skippy_session_copy_output_activation_frame( if (status != SKIPPY_STATUS_OK) { return status; } @@ -2571,7 +2578,7 @@ index 46da7007..7e11d92b 100644 } static enum skippy_status skippy_validate_state_range( -@@ -4566,6 +6300,13 @@ static llama_kv_cache * skippy_get_kv_cache( +@@ -4575,6 +6306,13 @@ static llama_kv_cache * skippy_get_kv_cache( if (auto * kv = dynamic_cast(memory)) { return kv; } @@ -2585,7 +2592,7 @@ index 46da7007..7e11d92b 100644 skippy_set_error(out_error, SKIPPY_STATUS_UNSUPPORTED, "runtime memory type is not supported for native KV pages"); return nullptr; -@@ -4635,6 +6376,10 @@ enum skippy_status skippy_import_kv_page( +@@ -4644,6 +6382,10 @@ enum skippy_status skippy_import_kv_page( skippy_set_error(out_error, SKIPPY_STATUS_RUNTIME_ERROR, error.c_str()); return SKIPPY_STATUS_RUNTIME_ERROR; } @@ -2596,7 +2603,7 @@ index 46da7007..7e11d92b 100644 session->n_past = std::max( session->n_past, static_cast(std::min( -@@ -4679,6 +6424,11 @@ enum skippy_status skippy_trim_session( +@@ -4689,6 +6431,11 @@ enum skippy_status skippy_trim_session( skippy_set_error(out_error, SKIPPY_STATUS_RUNTIME_ERROR, "failed to trim hybrid ISWA memory suffix"); return SKIPPY_STATUS_RUNTIME_ERROR; } @@ -2609,4 +2616,5 @@ index 46da7007..7e11d92b 100644 if (!kv->seq_rm(session->seq_id, p0, -1)) { skippy_set_error(out_error, SKIPPY_STATUS_RUNTIME_ERROR, "failed to trim native KV suffix"); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/patches/0024-tests-cover-native-GLM-DSA-execution-paths.patch b/third_party/llama.cpp/patches/0023-tests-cover-native-GLM-DSA-execution-paths.patch similarity index 99% rename from third_party/llama.cpp/patches/0024-tests-cover-native-GLM-DSA-execution-paths.patch rename to third_party/llama.cpp/patches/0023-tests-cover-native-GLM-DSA-execution-paths.patch index 7fc0e78feb..537f0285a0 100644 --- a/third_party/llama.cpp/patches/0024-tests-cover-native-GLM-DSA-execution-paths.patch +++ b/third_party/llama.cpp/patches/0023-tests-cover-native-GLM-DSA-execution-paths.patch @@ -1,7 +1,7 @@ -From 356609cab13dccdfb2069bcd739a3fff39e55e66 Mon Sep 17 00:00:00 2001 +From dbce275745b31af70ea3af5d023d2012da818c6d Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Tue, 14 Jul 2026 09:16:32 +1000 -Subject: [PATCH 21/24] tests: cover native GLM-DSA execution paths +Subject: [PATCH 23/45] tests: cover native GLM-DSA execution paths --- scripts/glm52-phase-d-policy-report.sh | 287 ++ @@ -1262,7 +1262,7 @@ index 00000000..8a1e0f38 + +echo "wrote ${summary}" diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp -index e3553c5e..18e06172 100644 +index e7cd6d0c..0a7c2997 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -51,6 +51,87 @@ @@ -1714,7 +1714,7 @@ index e3553c5e..18e06172 100644 // GGML_OP_GATED_LINEAR_ATTN struct test_gla : public test_case { const ggml_type type; -@@ -6060,120 +6373,4029 @@ struct test_topk_moe : public test_case { +@@ -6062,120 +6375,4029 @@ struct test_topk_moe : public test_case { } }; @@ -5830,7 +5830,7 @@ index e3553c5e..18e06172 100644 out = ggml_add(ctx, out, bias2); ggml_set_name(out, "out"); -@@ -6747,20 +10969,618 @@ struct test_timestep_embedding : public test_case { +@@ -6749,20 +10971,618 @@ struct test_timestep_embedding : public test_case { return VARS_TO_STR4(type, ne_a, dim, max_period); } @@ -6458,7 +6458,7 @@ index e3553c5e..18e06172 100644 }; // GGML_OP_LEAKY_RELU -@@ -6912,6 +11732,424 @@ struct test_flash_attn_ext : public test_case { +@@ -6914,6 +11734,424 @@ struct test_flash_attn_ext : public test_case { } }; @@ -6883,7 +6883,7 @@ index e3553c5e..18e06172 100644 // GGML_OP_CROSS_ENTROPY_LOSS struct test_cross_entropy_loss : public test_case { const ggml_type type; -@@ -8097,6 +13335,13 @@ static std::vector> make_test_cases_eval() { +@@ -8099,6 +13337,13 @@ static std::vector> make_test_cases_eval() { } test_cases.emplace_back(new test_get_rows(GGML_TYPE_F32, 1, 8, 2, 1, 1, false)); @@ -6897,7 +6897,7 @@ index e3553c5e..18e06172 100644 for (ggml_type type : all_types) { for (int b : {1, 7}) { for (bool v : {false, true}) { -@@ -8805,6 +14050,18 @@ static std::vector> make_test_cases_eval() { +@@ -8807,6 +14052,18 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 8192, 512, 5120, {128, 1}, {1, 1})); #endif @@ -6916,7 +6916,7 @@ index e3553c5e..18e06172 100644 for (ggml_type type_a : all_types) { for (int i = 1; i < 10; ++i) { test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 16, i, 256, { 1, 1}, {1, 1})); -@@ -9601,6 +14858,218 @@ static std::vector> make_test_cases_eval() { +@@ -9604,6 +14861,218 @@ static std::vector> make_test_cases_eval() { } } } @@ -7135,7 +7135,7 @@ index e3553c5e..18e06172 100644 test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 128, 1, 1)); test_cases.emplace_back(new test_gated_delta_net(GGML_TYPE_F32, 32, 16, 1, 1)); -@@ -9668,6 +15137,10 @@ static std::vector> make_test_cases_eval() { +@@ -9671,6 +15140,10 @@ static std::vector> make_test_cases_eval() { } } @@ -7146,7 +7146,7 @@ index e3553c5e..18e06172 100644 return test_cases; } #ifdef _MSC_VER -@@ -9824,6 +15297,21 @@ static std::vector> make_test_cases_perf() { +@@ -9827,6 +15300,21 @@ static std::vector> make_test_cases_perf() { } } @@ -7168,7 +7168,7 @@ index e3553c5e..18e06172 100644 // qwen3-30b-a3b for (int bs : {1, 4, 8, 32, 64, 128, 256, 512}) { for (ggml_type type_a : {GGML_TYPE_F32, GGML_TYPE_F16, GGML_TYPE_Q4_0, GGML_TYPE_Q8_0, GGML_TYPE_Q4_K, GGML_TYPE_Q6_K, GGML_TYPE_IQ2_XS}) { -@@ -9843,6 +15331,266 @@ static std::vector> make_test_cases_perf() { +@@ -9846,6 +15334,266 @@ static std::vector> make_test_cases_perf() { } } @@ -7436,7 +7436,7 @@ index e3553c5e..18e06172 100644 // gpt-oss-20b for (int bs : {1, 4, 8, 512}) { diff --git a/tests/test-llama-archs.cpp b/tests/test-llama-archs.cpp -index 57d33a62..fe0accc8 100644 +index 86c3051c..a2e79125 100644 --- a/tests/test-llama-archs.cpp +++ b/tests/test-llama-archs.cpp @@ -14,6 +14,7 @@ @@ -8433,7 +8433,7 @@ index 57d33a62..fe0accc8 100644 static std::vector get_logits( llama_model * model, llama_context * lctx, const std::vector & tokens, bool encode = false) { const uint32_t n_vocab = llama_vocab_n_tokens(llama_model_get_vocab(model)); -@@ -429,6 +1360,20 @@ static bool arch_supported(const llm_arch arch) { +@@ -430,6 +1361,20 @@ static bool arch_supported(const llm_arch arch) { return true; } @@ -8454,7 +8454,7 @@ index 57d33a62..fe0accc8 100644 static int save_models(const llm_arch target_arch, const size_t seed, const ggml_log_level log_level, const std::string & dir) { struct user_data_t { struct { -@@ -542,6 +1487,33 @@ static int test_backends(const llm_arch target_arch, const size_t seed, const gg +@@ -543,6 +1488,33 @@ static int test_backends(const llm_arch target_arch, const size_t seed, const gg const std::string template_row_res = "%15s %10s|%20s|\n"; bool all_ok = true; @@ -8488,7 +8488,7 @@ index 57d33a62..fe0accc8 100644 common_log_flush(common_log_main()); printf(template_header.c_str(), "Model arch.", "Device", "Config", "NMSE vs. CPU", "Roundtrip"); printf("|"); -@@ -591,6 +1563,9 @@ static int test_backends(const llm_arch target_arch, const size_t seed, const gg +@@ -592,6 +1564,9 @@ static int test_backends(const llm_arch target_arch, const size_t seed, const gg std::string status_roundtrip = "\033[1;33mSKIP\033[0m"; char nmse_str[12] = {0}; bool skip = !arch_supported(arch) || (dc.split_mode == LLAMA_SPLIT_MODE_TENSOR && dc.devs.empty()); @@ -8499,4 +8499,5 @@ index 57d33a62..fe0accc8 100644 skip = true; // FIXME #endif // GGML_USE_WEBGPU -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/patches/0025-Support-GLM-DSA-fused-KV_B-tensors.patch b/third_party/llama.cpp/patches/0024-Support-GLM-DSA-fused-KV_B-tensors.patch similarity index 93% rename from third_party/llama.cpp/patches/0025-Support-GLM-DSA-fused-KV_B-tensors.patch rename to third_party/llama.cpp/patches/0024-Support-GLM-DSA-fused-KV_B-tensors.patch index 71600508da..6cbe4bfd16 100644 --- a/third_party/llama.cpp/patches/0025-Support-GLM-DSA-fused-KV_B-tensors.patch +++ b/third_party/llama.cpp/patches/0024-Support-GLM-DSA-fused-KV_B-tensors.patch @@ -1,7 +1,7 @@ -From 320112e2be2f38c92559b4eaeb482816c6037874 Mon Sep 17 00:00:00 2001 +From 358fdff04cab3ca3d85cd51dbc5592a1378df9fc Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Wed, 24 Jun 2026 15:35:23 +1000 -Subject: [PATCH 22/24] Support GLM-DSA fused KV_B tensors +Subject: [PATCH 24/45] Support GLM-DSA fused KV_B tensors --- src/models/deepseek2.cpp | 23 +++++++++++++++++------ @@ -64,10 +64,10 @@ index 180fb082..feb05570 100644 cb(q_nope, "mtp_q_nope_perm", il); diff --git a/src/models/glm-dsa.cpp b/src/models/glm-dsa.cpp -index 10bf235f..a975da0d 100644 +index 1d90426d..b682d670 100644 --- a/src/models/glm-dsa.cpp +++ b/src/models/glm-dsa.cpp -@@ -87,9 +87,13 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { +@@ -122,9 +122,13 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { layer.wkv_a_mqa = create_tensor(tn(LLM_TENSOR_ATTN_KV_A_MQA, "weight", i), {n_embd, kv_lora_rank + n_embd_head_qk_rope}, flags); @@ -85,5 +85,5 @@ index 10bf235f..a975da0d 100644 layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_head * n_embd_head_v_mla, n_embd}, flags); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0025-Run-GLM-DSA-through-IndexShare-graph.patch b/third_party/llama.cpp/patches/0025-Run-GLM-DSA-through-IndexShare-graph.patch new file mode 100644 index 0000000000..cb59f5881c --- /dev/null +++ b/third_party/llama.cpp/patches/0025-Run-GLM-DSA-through-IndexShare-graph.patch @@ -0,0 +1,93 @@ +From 89bf8ceca0a8dfdccf9abed3f859a1c80af21948 Mon Sep 17 00:00:00 2001 +From: Mesh-LLM CI +Date: Thu, 25 Jun 2026 15:51:20 +1000 +Subject: [PATCH 25/45] Run GLM-DSA through IndexShare graph + +--- + src/models/glm-dsa.cpp | 43 ++++++++++++++++++++++++++++++++++++++---- + 1 file changed, 39 insertions(+), 4 deletions(-) + +diff --git a/src/models/glm-dsa.cpp b/src/models/glm-dsa.cpp +index b682d670..82c6a093 100644 +--- a/src/models/glm-dsa.cpp ++++ b/src/models/glm-dsa.cpp +@@ -224,20 +224,37 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par + ggml_tensor * cur; + ggml_tensor * inpL; + ++ const int effective_n_layers = n_layer; ++ const skippy_graph_filter & stage_filter = skippy_graph_get_filter(); ++ const bool stage_filtered = stage_filter.enabled; ++ const int il_start = stage_filtered ? stage_filter.layer_start : 0; ++ const int il_end = stage_filtered ? ++ std::min(stage_filter.layer_end, effective_n_layers) : ++ effective_n_layers; ++ ++ if (stage_filtered && il_start < il_end && !hparams.is_indexer_full(il_start)) { ++ throw std::runtime_error("GLM_DSA staged execution must start on a full indexer layer"); ++ } ++ for (int il = il_start; il < il_end; ++il) { ++ if (!model.layers[il].wk_b || !model.layers[il].wv_b) { ++ throw std::runtime_error("GLM_DSA sparse attention requires split K_B and V_B tensors"); ++ } ++ } ++ + // {n_embd, n_tokens} +- inpL = build_inp_embd(model.tok_embd); ++ inpL = build_inp_embd(stage_filtered && il_start > 0 ? nullptr : model.tok_embd); + + // inp_pos - contains the positions + ggml_tensor * inp_pos = build_inp_pos(); + + llm_graph_input_attn_k_dsa * inp_attn_dsa = build_attn_inp_k_dsa(); + +- ggml_tensor * inp_out_ids = build_inp_out_ids(); ++ ggml_tensor * inp_out_ids = (!stage_filtered || stage_filter.include_output) ? build_inp_out_ids() : nullptr; + + // Difference vs Deepseek 3.2: shared indexer layers reuse the top_k from the previous full indexer layers + // See https://huggingface.co/zai-org/GLM-5.2/blob/main/config.json#L30 + ggml_tensor * prev_top_k = nullptr; +- for (int il = 0; il < n_layer; ++il) { ++ for (int il = il_start; il < il_end; ++il) { + ggml_tensor * inpSA = inpL; + + // norm +@@ -472,7 +489,7 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par + Qcur, Kcur, Vcur, nullptr, nullptr, model.layers[il].wv_b, top_k, kq_scale, il); + } + } +- if (il == n_layer - 1 && inp_out_ids) { ++ if (il == il_end - 1 && inp_out_ids && cparams.embeddings_nextn_masked) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); + inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids); + } +@@ -533,8 +550,26 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par + } + cur = inpL; + ++ if (stage_filtered && !stage_filter.include_output) { ++ if (!cparams.embeddings_nextn_masked && inp_out_ids) { ++ cur = ggml_get_rows(ctx0, cur, inp_out_ids); ++ } ++ ++ cb(cur, "stage_boundary", il_end - 1); ++ res->t_embd = cur; ++ ggml_build_forward_expand(gf, cur); ++ return; ++ } ++ + cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1); + ++ cb(cur, "h_nextn", -1); ++ res->t_h_nextn = cur; ++ ++ if (!cparams.embeddings_nextn_masked && inp_out_ids) { ++ cur = ggml_get_rows(ctx0, cur, inp_out_ids); ++ } ++ + cb(cur, "result_norm", -1); + res->t_embd = cur; + +-- +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/patches/0028-Bump-Skippy-ABI-for-GLM-DSA-runtime-config.patch b/third_party/llama.cpp/patches/0026-Bump-Skippy-ABI-for-GLM-DSA-runtime-config.patch similarity index 74% rename from third_party/llama.cpp/patches/0028-Bump-Skippy-ABI-for-GLM-DSA-runtime-config.patch rename to third_party/llama.cpp/patches/0026-Bump-Skippy-ABI-for-GLM-DSA-runtime-config.patch index b5eb4614f6..7ee7eb102c 100644 --- a/third_party/llama.cpp/patches/0028-Bump-Skippy-ABI-for-GLM-DSA-runtime-config.patch +++ b/third_party/llama.cpp/patches/0026-Bump-Skippy-ABI-for-GLM-DSA-runtime-config.patch @@ -1,14 +1,14 @@ -From 0fb232f0f632084ab4fce7ed82ad12e4afe6a56f Mon Sep 17 00:00:00 2001 +From f9e182d9e48b35b1e9e445b6267bac73fc472291 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Tue, 21 Jul 2026 15:31:56 +1000 -Subject: [PATCH] Bump Skippy ABI for GLM DSA runtime config +Subject: [PATCH 26/45] Bump Skippy ABI for GLM DSA runtime config --- include/skippy/common.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/skippy/common.h b/include/skippy/common.h -index 2d73bf63..3346b6a5 100644 +index 17b7203c..aeea033d 100644 --- a/include/skippy/common.h +++ b/include/skippy/common.h @@ -26,7 +26,7 @@ extern "C" { @@ -21,5 +21,5 @@ index 2d73bf63..3346b6a5 100644 enum skippy_feature { SKIPPY_FEATURE_RUNTIME_SLICE = 1 << 0, -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0026-Run-GLM-DSA-through-IndexShare-graph.patch b/third_party/llama.cpp/patches/0026-Run-GLM-DSA-through-IndexShare-graph.patch deleted file mode 100644 index 79eb6fbb08..0000000000 --- a/third_party/llama.cpp/patches/0026-Run-GLM-DSA-through-IndexShare-graph.patch +++ /dev/null @@ -1,370 +0,0 @@ -From e3ec7639933e124f86f056e61b5a6b375f108e70 Mon Sep 17 00:00:00 2001 -From: Mesh-LLM CI -Date: Thu, 25 Jun 2026 15:51:20 +1000 -Subject: [PATCH 23/24] Run GLM-DSA through IndexShare graph - ---- - src/models/glm-dsa.cpp | 322 +++++++++++++++++++++++++++++++++++++++++ - src/models/models.h | 5 +- - 2 files changed, 326 insertions(+), 1 deletion(-) - -diff --git a/src/models/glm-dsa.cpp b/src/models/glm-dsa.cpp -index a975da0d..adde84fd 100644 ---- a/src/models/glm-dsa.cpp -+++ b/src/models/glm-dsa.cpp -@@ -1,5 +1,7 @@ - #include "models.h" - -+#include "llama-kv-cache-dsa.h" -+ - void llama_model_glm_dsa::load_arch_hparams(llama_model_loader & ml) { - ml.get_key(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp); - ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps); -@@ -145,6 +147,326 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { - } - } - -+llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_params & params) : -+ llm_graph_context(params) { -+ const bool is_mla = hparams.is_mla(); -+ GGML_ASSERT(is_mla); -+ -+ const int64_t n_embd_head_k = hparams.n_embd_head_k_mla(); -+ const int64_t n_embd_head_v = hparams.n_embd_head_v_mla(); -+ GGML_UNUSED(n_embd_head_v); -+ -+ const int64_t n_embd_head_qk_rope = hparams.n_rot(); -+ const int64_t n_embd_head_qk_nope = n_embd_head_k - n_embd_head_qk_rope; -+ -+ const int64_t n_indexer_head = hparams.indexer_n_head; -+ const int64_t n_embd_indexer_head = hparams.indexer_head_size; -+ const int64_t n_embd_indexer_head_rope = hparams.n_rot(); -+ const int64_t n_embd_indexer_head_nope = n_embd_indexer_head - n_embd_indexer_head_rope; -+ const uint32_t n_indexer_top_k = hparams.indexer_top_k; -+ -+ const uint32_t kv_lora_rank = hparams.n_lora_kv; -+ -+ GGML_ASSERT(ext_factor >= 0.0f); -+ const float attn_factor_org = attn_factor * (1.0f + 0.1f * logf(1.0f / freq_scale)); -+ const float mscale = attn_factor_org * (1.0f + 0.1f * hparams.rope_yarn_log_mul * logf(1.0f / freq_scale)); -+ const float kq_scale = 1.0f * mscale * mscale / sqrtf(float(n_embd_head_k)); -+ -+ ggml_tensor * cur; -+ ggml_tensor * inpL; -+ -+ const int effective_n_layers = n_layer; -+ const skippy_graph_filter & stage_filter = skippy_graph_get_filter(); -+ const bool stage_filtered = stage_filter.enabled; -+ const int il_start = stage_filtered ? stage_filter.layer_start : 0; -+ const int il_end = stage_filtered ? -+ std::min(stage_filter.layer_end, effective_n_layers) : -+ effective_n_layers; -+ -+ for (int il = il_start; il < il_end; ++il) { -+ if (!model.layers[il].wk_b || !model.layers[il].wv_b) { -+ throw std::runtime_error("GLM_DSA sparse attention requires split K_B and V_B tensors"); -+ } -+ } -+ -+ inpL = build_inp_embd(stage_filtered && il_start > 0 ? nullptr : model.tok_embd); -+ -+ ggml_tensor * inp_pos = build_inp_pos(); -+ llm_graph_input_attn_k_dsa * inp_attn_dsa = build_attn_inp_k_dsa(); -+ ggml_tensor * inp_out_ids = (!stage_filtered || stage_filter.include_output) ? build_inp_out_ids() : nullptr; -+ -+ for (int il = il_start; il < il_end; ++il) { -+ ggml_tensor * inpSA = inpL; -+ -+ cur = build_norm(inpL, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il); -+ cb(cur, "attn_norm", il); -+ -+ { -+ ggml_tensor * qr = ggml_mul_mat(ctx0, model.layers[il].wq_a, cur); -+ cb(qr, "qr", il); -+ -+ qr = build_norm(qr, model.layers[il].attn_q_a_norm, nullptr, LLM_NORM_RMS, il); -+ cb(qr, "qr", il); -+ -+ ggml_tensor * top_k = nullptr; -+ -+ { -+ ggml_tensor * indexer_q = ggml_mul_mat(ctx0, model.layers[il].indexer_attn_q_b, qr); -+ cb(indexer_q, "indexer_q", il); -+ -+ ggml_tensor * indexer_q_pe = -+ ggml_view_3d(ctx0, indexer_q, n_embd_indexer_head_rope, n_indexer_head, n_tokens, -+ ggml_row_size(indexer_q->type, n_embd_indexer_head), -+ ggml_row_size(indexer_q->type, n_embd_indexer_head) * n_indexer_head, 0); -+ cb(indexer_q_pe, "indexer_q_pe", il); -+ -+ ggml_tensor * indexer_q_nope = -+ ggml_view_3d(ctx0, indexer_q, n_embd_indexer_head_nope, n_indexer_head, n_tokens, -+ ggml_row_size(indexer_q->type, n_embd_indexer_head), -+ ggml_row_size(indexer_q->type, n_embd_indexer_head) * n_indexer_head, -+ ggml_row_size(indexer_q->type, n_embd_indexer_head_nope)); -+ cb(indexer_q_nope, "indexer_q_nope", il); -+ -+ indexer_q_pe = ggml_rope_ext(ctx0, indexer_q_pe, inp_pos, nullptr, n_rot, -+ LLAMA_ROPE_TYPE_NEOX, n_ctx_orig, freq_base, freq_scale, -+ ext_factor, attn_factor, beta_fast, beta_slow); -+ cb(indexer_q_pe, "indexer_q_pe", il); -+ -+ indexer_q = ggml_concat(ctx0, indexer_q_pe, indexer_q_nope, 0); -+ cb(indexer_q, "indexer_q", il); -+ -+ ggml_tensor * indexer_k = ggml_mul_mat(ctx0, model.layers[il].indexer_attn_k, cur); -+ cb(indexer_k, "indexer_k", il); -+ -+ indexer_k = build_norm(indexer_k, model.layers[il].indexer_k_norm, model.layers[il].indexer_k_norm_b, LLM_NORM, il); -+ cb(indexer_k, "indexer_k", il); -+ -+ ggml_tensor * indexer_k_pe = -+ ggml_view_3d(ctx0, indexer_k, n_embd_indexer_head_rope, 1, n_tokens, -+ ggml_row_size(indexer_k->type, n_embd_indexer_head), -+ ggml_row_size(indexer_k->type, n_embd_indexer_head) * 1, 0); -+ cb(indexer_k_pe, "indexer_k_pe", il); -+ -+ ggml_tensor * indexer_k_nope = -+ ggml_view_3d(ctx0, indexer_k, n_embd_indexer_head_nope, 1, n_tokens, -+ ggml_row_size(indexer_k->type, n_embd_indexer_head), -+ ggml_row_size(indexer_k->type, n_embd_indexer_head) * 1, -+ ggml_row_size(indexer_k->type, n_embd_indexer_head_nope)); -+ cb(indexer_k_nope, "indexer_k_nope", il); -+ -+ indexer_k_pe = ggml_rope_ext(ctx0, indexer_k_pe, inp_pos, nullptr, n_rot, -+ LLAMA_ROPE_TYPE_NEOX, n_ctx_orig, freq_base, freq_scale, -+ ext_factor, attn_factor, beta_fast, beta_slow); -+ cb(indexer_k_pe, "indexer_k_pe", il); -+ -+ indexer_k = ggml_concat(ctx0, indexer_k_pe, indexer_k_nope, 0); -+ cb(indexer_k, "indexer_k", il); -+ -+ indexer_q = ggml_mul_mat(ctx0, inp_attn_dsa->self_k_rot_lid, indexer_q); -+ cb(indexer_q, "indexer_q", il); -+ indexer_k = ggml_mul_mat(ctx0, inp_attn_dsa->self_k_rot_lid, indexer_k); -+ cb(indexer_k, "indexer_k", il); -+ -+ const auto * mctx_lid = inp_attn_dsa->mctx->get_lid(); -+ const auto & k_idxs_lid = inp_attn_dsa->get_k_idxs_lid(); -+ ggml_build_forward_expand(gf, mctx_lid->cpy_k(ctx0, indexer_k, k_idxs_lid, il)); -+ -+ ggml_tensor * indexer_weights = ggml_mul_mat(ctx0, model.layers[il].indexer_proj, cur); -+ cb(indexer_weights, "indexer_weights", il); -+ -+ indexer_k = mctx_lid->get_k(ctx0, il); -+ -+ const auto n_stream = indexer_k->ne[3]; -+ indexer_q = ggml_view_4d(ctx0, indexer_q, indexer_q->ne[0], indexer_q->ne[1], indexer_q->ne[2]/n_stream, n_stream, indexer_q->nb[1], indexer_q->nb[2], indexer_q->nb[3]/n_stream, 0); -+ indexer_weights = ggml_view_4d(ctx0, indexer_weights, indexer_weights->ne[0], indexer_weights->ne[1]/n_stream, indexer_weights->ne[2], n_stream, indexer_weights->nb[1], indexer_weights->nb[2]/n_stream, indexer_weights->nb[3]/n_stream, 0); -+ -+ indexer_q = ggml_permute(ctx0, indexer_q, 0, 2, 1, 3); -+ cb(indexer_q, "indexer_q", il); -+ indexer_k = ggml_permute(ctx0, indexer_k, 0, 2, 1, 3); -+ cb(indexer_k, "indexer_k", il); -+ -+ ggml_tensor * indexer_kq = ggml_mul_mat(ctx0, indexer_k, indexer_q); -+ cb(indexer_kq, "indexer_kq", il); -+ -+ indexer_kq = ggml_cont(ctx0, ggml_permute(ctx0, indexer_kq, 2, 1, 0, 3)); -+ cb(indexer_kq, "indexer_kq", il); -+ -+ ggml_tensor * indexer_score = ggml_relu(ctx0, indexer_kq); -+ cb(indexer_score, "indexer_score", il); -+ -+ indexer_weights = ggml_scale(ctx0, indexer_weights, 1.0f / sqrtf(float(n_embd_indexer_head * n_indexer_head))); -+ cb(indexer_weights, "indexer_weights", il); -+ -+ indexer_score = ggml_mul(ctx0, indexer_score, indexer_weights); -+ cb(indexer_score, "indexer_score", il); -+ -+ indexer_score = ggml_sum_rows(ctx0, indexer_score); -+ cb(indexer_score, "indexer_score", il); -+ -+ indexer_score = ggml_cont(ctx0, ggml_permute(ctx0, indexer_score, 2, 1, 0, 3)); -+ cb(indexer_score, "indexer_score", il); -+ -+ ggml_tensor * indexer_kq_mask = inp_attn_dsa->get_kq_mask_lid(); -+ indexer_score = ggml_add(ctx0, indexer_score, indexer_kq_mask); -+ cb(indexer_score, "indexer_score", il); -+ -+ uint32_t n_top_k = indexer_score->ne[0] < n_indexer_top_k ? indexer_score->ne[0] : n_indexer_top_k; -+ top_k = ggml_cont(ctx0, ggml_top_k(ctx0, indexer_score, n_top_k)); -+ cb(top_k, "top_k", il); -+ } -+ -+ ggml_tensor * q = ggml_mul_mat(ctx0, model.layers[il].wq_b, qr); -+ cb(q, "q", il); -+ -+ ggml_tensor * q_nope = -+ ggml_view_3d(ctx0, q, n_embd_head_qk_nope, n_head, n_tokens, ggml_row_size(q->type, n_embd_head_k), -+ ggml_row_size(q->type, n_embd_head_k) * n_head, 0); -+ cb(q_nope, "q_nope", il); -+ -+ ggml_tensor * q_pe = ggml_view_3d( -+ ctx0, q, n_embd_head_qk_rope, n_head, n_tokens, ggml_row_size(q->type, n_embd_head_k), -+ ggml_row_size(q->type, n_embd_head_k) * n_head, ggml_row_size(q->type, n_embd_head_qk_nope)); -+ cb(q_pe, "q_pe", il); -+ -+ ggml_tensor * kv_cmpr_pe = ggml_mul_mat(ctx0, model.layers[il].wkv_a_mqa, cur); -+ cb(kv_cmpr_pe, "kv_cmpr_pe", il); -+ -+ ggml_tensor * kv_cmpr = -+ ggml_view_2d(ctx0, kv_cmpr_pe, kv_lora_rank, n_tokens, -+ ggml_row_size(kv_cmpr_pe->type, kv_lora_rank + n_embd_head_qk_rope), 0); -+ cb(kv_cmpr, "kv_cmpr", il); -+ -+ ggml_tensor * k_pe = ggml_view_3d(ctx0, kv_cmpr_pe, n_embd_head_qk_rope, 1, n_tokens, -+ ggml_row_size(kv_cmpr_pe->type, kv_lora_rank + n_embd_head_qk_rope), -+ ggml_row_size(kv_cmpr_pe->type, kv_lora_rank + n_embd_head_qk_rope), -+ ggml_row_size(kv_cmpr_pe->type, kv_lora_rank)); -+ cb(k_pe, "k_pe", il); -+ -+ q_pe = ggml_rope_ext(ctx0, q_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, -+ ext_factor, attn_factor, beta_fast, beta_slow); -+ cb(q_pe, "q_pe", il); -+ -+ k_pe = ggml_rope_ext(ctx0, k_pe, inp_pos, nullptr, n_rot, rope_type, n_ctx_orig, freq_base, freq_scale, -+ ext_factor, attn_factor, beta_fast, beta_slow); -+ cb(k_pe, "k_pe", il); -+ -+ kv_cmpr = build_norm(kv_cmpr, model.layers[il].attn_kv_a_norm, nullptr, LLM_NORM_RMS, il); -+ cb(kv_cmpr, "kv_cmpr", il); -+ -+ q_nope = ggml_permute(ctx0, q_nope, 0, 2, 1, 3); -+ cb(q_nope, "q_nope_perm", il); -+ -+ ggml_tensor * q_nope_absorbed = ggml_mul_mat(ctx0, model.layers[il].wk_b, q_nope); -+ cb(q_nope_absorbed, "q_nope_absorbed", il); -+ -+ q_nope_absorbed = ggml_permute(ctx0, q_nope_absorbed, 0, 2, 1, 3); -+ cb(q_nope_absorbed, "q_nope_absorbed_perm", il); -+ -+ ggml_tensor * Qcur = ggml_concat(ctx0, q_nope_absorbed, q_pe, 0); -+ cb(Qcur, "Qcur", il); -+ -+ kv_cmpr = ggml_reshape_3d(ctx0, kv_cmpr, kv_lora_rank, 1, n_tokens); -+ cb(kv_cmpr, "kv_cmpr_reshape", il); -+ -+ ggml_tensor * Kcur = ggml_concat(ctx0, kv_cmpr, k_pe, 0); -+ cb(Kcur, "Kcur", il); -+ -+ ggml_tensor * Vcur = kv_cmpr; -+ cb(Vcur, "Vcur", il); -+ -+ cur = build_attn(inp_attn_dsa, -+ model.layers[il].wo, NULL, model.layers[il].wo_s, -+ Qcur, Kcur, Vcur, nullptr, nullptr, model.layers[il].wv_b, top_k, kq_scale, il); -+ } -+ -+ if (il == il_end - 1 && inp_out_ids && cparams.embeddings_nextn_masked) { -+ 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, NULL, LLM_NORM_RMS, il); -+ cb(cur, "ffn_norm", il); -+ -+ if ((uint32_t) il < hparams.n_layer_dense_lead) { -+ cur = build_ffn(cur, -+ model.layers[il].ffn_up, NULL, model.layers[il].ffn_up_s, -+ model.layers[il].ffn_gate, NULL, model.layers[il].ffn_gate_s, -+ model.layers[il].ffn_down, NULL, model.layers[il].ffn_down_s, -+ NULL, LLM_FFN_SILU, LLM_FFN_PAR, il); -+ cb(cur, "ffn_out", il); -+ } else { -+ ggml_tensor * moe_out = 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, -+ model.layers[il].ffn_exp_probs_b, -+ n_expert, n_expert_used, -+ LLM_FFN_SILU, hparams.expert_weights_norm, -+ hparams.expert_weights_scale, -+ (llama_expert_gating_func_type) hparams.expert_gating_func, -+ il, -+ nullptr, -+ model.layers[il].ffn_gate_up_exps, -+ model.layers[il].ffn_up_exps_s, -+ model.layers[il].ffn_gate_exps_s, -+ model.layers[il].ffn_down_exps_s); -+ cb(moe_out, "ffn_moe_out", il); -+ -+ ggml_tensor * ffn_shexp = -+ build_ffn(cur, -+ model.layers[il].ffn_up_shexp, NULL, model.layers[il].ffn_up_shexp_s, -+ model.layers[il].ffn_gate_shexp, NULL, model.layers[il].ffn_gate_shexp_s, -+ model.layers[il].ffn_down_shexp, NULL, model.layers[il].ffn_down_shexp_s, -+ NULL, LLM_FFN_SILU, LLM_FFN_PAR, il); -+ cb(ffn_shexp, "ffn_shexp", il); -+ -+ cur = ggml_add(ctx0, moe_out, ffn_shexp); -+ cb(cur, "ffn_out", il); -+ } -+ cur = ggml_add(ctx0, cur, ffn_inp); -+ -+ cur = build_cvec(cur, il); -+ cb(cur, "l_out", il); -+ -+ inpL = cur; -+ } -+ cur = inpL; -+ -+ if (stage_filtered && !stage_filter.include_output) { -+ if (!cparams.embeddings_nextn_masked && inp_out_ids) { -+ cur = ggml_get_rows(ctx0, cur, inp_out_ids); -+ } -+ -+ cb(cur, "stage_boundary", il_end - 1); -+ res->t_embd = cur; -+ ggml_build_forward_expand(gf, cur); -+ return; -+ } -+ -+ cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1); -+ -+ cb(cur, "h_nextn", -1); -+ res->t_h_nextn = cur; -+ -+ if (!cparams.embeddings_nextn_masked && inp_out_ids) { -+ cur = ggml_get_rows(ctx0, cur, inp_out_ids); -+ } -+ -+ cb(cur, "result_norm", -1); -+ res->t_embd = cur; -+ -+ cur = ggml_mul_mat(ctx0, model.output, cur); -+ -+ cb(cur, "result_output", -1); -+ res->t_logits = cur; -+ -+ ggml_build_forward_expand(gf, cur); -+} -+ - std::unique_ptr llama_model_glm_dsa::build_arch_graph(const llm_graph_params & params) const { - if (params.gtype == LLM_GRAPH_TYPE_DECODER_MTP) { - return std::make_unique(*this, params); -diff --git a/src/models/models.h b/src/models/models.h -index e9a7c0c2..d8120d9d 100644 ---- a/src/models/models.h -+++ b/src/models/models.h -@@ -1221,9 +1221,12 @@ struct llama_model_glm_dsa : public llama_model_base { - void load_arch_hparams(llama_model_loader & ml) override; - void load_arch_tensors(llama_model_loader & ml) override; - -- using graph = llama_model_deepseek2::graph; - using graph_mtp = llama_model_deepseek2::graph_mtp; - -+ struct graph : public llm_graph_context { -+ graph(const llama_model & model, const llm_graph_params & params); -+ }; -+ - std::unique_ptr build_arch_graph(const llm_graph_params & params) const override; - }; - --- -2.54.0 (Apple Git-157) - diff --git a/third_party/llama.cpp/patches/0029-Fix-GLM-DSA-Metal-get_rows-placement.patch b/third_party/llama.cpp/patches/0027-Fix-GLM-DSA-Metal-get_rows-placement.patch similarity index 96% rename from third_party/llama.cpp/patches/0029-Fix-GLM-DSA-Metal-get_rows-placement.patch rename to third_party/llama.cpp/patches/0027-Fix-GLM-DSA-Metal-get_rows-placement.patch index 4f8b563e44..50416e2eb2 100644 --- a/third_party/llama.cpp/patches/0029-Fix-GLM-DSA-Metal-get_rows-placement.patch +++ b/third_party/llama.cpp/patches/0027-Fix-GLM-DSA-Metal-get_rows-placement.patch @@ -1,7 +1,7 @@ -From c982f6c4823c99f194690c87f5498bdb03c90a65 Mon Sep 17 00:00:00 2001 +From 9b1b59b34069b19c59f280e1fb4c4480e3b173d0 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Wed, 22 Jul 2026 15:13:16 +1000 -Subject: [PATCH] Fix GLM DSA Metal get_rows placement +Subject: [PATCH 27/45] Fix GLM DSA Metal get_rows placement --- ggml/src/ggml-metal/ggml-metal.metal | 22 ---------------------- @@ -41,5 +41,5 @@ index b7b7e065..7a6df158 100644 if (tid < stride) { reduce[reduce_offset + tid] = max( -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0027-Use-DSA-KV-cache-for-GLM-DSA.patch b/third_party/llama.cpp/patches/0027-Use-DSA-KV-cache-for-GLM-DSA.patch deleted file mode 100644 index 3004688e4b..0000000000 --- a/third_party/llama.cpp/patches/0027-Use-DSA-KV-cache-for-GLM-DSA.patch +++ /dev/null @@ -1,41 +0,0 @@ -From 847934296d90daed13b31e8b76b0a5f44b801b05 Mon Sep 17 00:00:00 2001 -From: Mesh-LLM CI -Date: Thu, 25 Jun 2026 15:58:34 +1000 -Subject: [PATCH 24/24] Use DSA KV cache for GLM-DSA - ---- - src/llama-kv-cache.cpp | 5 +++-- - src/llama-model.cpp | 1 + - 2 files changed, 4 insertions(+), 2 deletions(-) - -diff --git a/src/llama-kv-cache.cpp b/src/llama-kv-cache.cpp -index 43e9492a..9df4984b 100644 ---- a/src/llama-kv-cache.cpp -+++ b/src/llama-kv-cache.cpp -@@ -321,8 +321,9 @@ llama_kv_cache::llama_kv_cache( - ggml_is_quantized(type_k) && - hparams.n_embd_head_k() % 64 == 0; - -- // always create Hadamard rotation tensors for DeepSeek V3.2 DSA lightning indexer -- if (model.arch == LLM_ARCH_DEEPSEEK32 && hparams.n_embd_head_k_full == hparams.indexer_head_size) { -+ // always create Hadamard rotation tensors for the DSA lightning indexer -+ if ((model.arch == LLM_ARCH_DEEPSEEK32 || model.arch == LLM_ARCH_GLM_DSA) && -+ hparams.n_embd_head_k_full == hparams.indexer_head_size) { - attn_rot_k = true; - } - -diff --git a/src/llama-model.cpp b/src/llama-model.cpp -index 5b83c0c0..a325a4ed 100644 ---- a/src/llama-model.cpp -+++ b/src/llama-model.cpp -@@ -2061,6 +2061,7 @@ llama_memory_i * llama_model::create_memory(const llama_memory_params & params, - res = nullptr; - } break; - case LLM_ARCH_DEEPSEEK32: -+ case LLM_ARCH_GLM_DSA: - { - res = new llama_kv_cache_dsa( - *this, --- -2.54.0 (Apple Git-157) - diff --git a/third_party/llama.cpp/patches/0030-ggml-default-GLM-MoE-two-phase-Metal-path.patch b/third_party/llama.cpp/patches/0028-ggml-default-GLM-MoE-two-phase-Metal-path.patch similarity index 88% rename from third_party/llama.cpp/patches/0030-ggml-default-GLM-MoE-two-phase-Metal-path.patch rename to third_party/llama.cpp/patches/0028-ggml-default-GLM-MoE-two-phase-Metal-path.patch index e5d3caa626..2d5d3b6e2a 100644 --- a/third_party/llama.cpp/patches/0030-ggml-default-GLM-MoE-two-phase-Metal-path.patch +++ b/third_party/llama.cpp/patches/0028-ggml-default-GLM-MoE-two-phase-Metal-path.patch @@ -1,7 +1,7 @@ -From 11fb71bbe0c6a1b525097a455d75810f0a4e43e4 Mon Sep 17 00:00:00 2001 +From 1f9c781f915ef3ea0f71f48660272b1c31e19d9e Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Sun, 19 Jul 2026 11:27:53 +1000 -Subject: [PATCH 25/36] ggml: default GLM MoE two-phase Metal path +Subject: [PATCH 28/45] ggml: default GLM MoE two-phase Metal path --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 9 ++++++++- @@ -29,10 +29,10 @@ index 36378f8e..684efdc6 100644 static bool ggml_metal_glm_dsa_moe_dual_lane_enabled() { diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp -index 8618ea52..56fb7b23 100644 +index 0a7c2997..cc7703c8 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp -@@ -9832,6 +9832,17 @@ struct test_glm_ffn_two_phase_consumer : public test_glm_ffn_consumer { +@@ -9835,6 +9835,17 @@ struct test_glm_ffn_two_phase_consumer : public test_glm_ffn_consumer { return std::string("GLM_FFN_TWO_PHASE_") + variant; } @@ -50,7 +50,7 @@ index 8618ea52..56fb7b23 100644 bool eval_perf(ggml_backend_t backend, const char * op_names_filter, printer * output_printer) override { scoped_test_env motif_reference_env("GGML_METAL_ENABLE_GLM_MOE_DECODE_MOTIF_REFERENCE", "0"); scoped_test_env two_phase_env( -@@ -15138,6 +15149,7 @@ static std::vector> make_test_cases_eval() { +@@ -15142,6 +15153,7 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_glm_ffn_chain_native_down(true, "CANDIDATE")); test_cases.emplace_back(new test_glm_ffn_chain_two_phase(true, "CANDIDATE")); @@ -59,5 +59,5 @@ index 8618ea52..56fb7b23 100644 return test_cases; -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0031-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch b/third_party/llama.cpp/patches/0029-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch similarity index 95% rename from third_party/llama.cpp/patches/0031-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch rename to third_party/llama.cpp/patches/0029-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch index cd3effc48c..34f60b209c 100644 --- a/third_party/llama.cpp/patches/0031-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch +++ b/third_party/llama.cpp/patches/0029-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch @@ -1,7 +1,7 @@ -From 8797805107cd1e49bc018a498ce9d008493000f8 Mon Sep 17 00:00:00 2001 +From d2550d768c3d538b50266b13861dc1ad39e266fa Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Sun, 19 Jul 2026 16:31:29 +1000 -Subject: [PATCH 26/36] ggml: add GLM MoE Metal selector diagnostics +Subject: [PATCH 29/45] ggml: add GLM MoE Metal selector diagnostics --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 75 +++++++++++++++++++++----- @@ -139,10 +139,10 @@ index 684efdc6..ab41e1ec 100644 fusion_scratch : ggml_metal_get_buffer_id(weighted_sum); diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp -index 56fb7b23..53489947 100644 +index cc7703c8..50ec5225 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp -@@ -9575,6 +9575,7 @@ struct test_glm_ffn_consumer : public test_case { +@@ -9578,6 +9578,7 @@ struct test_glm_ffn_consumer : public test_case { ggml_tensor * selected_experts = nullptr; ggml_tensor * weights = nullptr; ggml_tensor * routed_out = nullptr; @@ -150,7 +150,7 @@ index 56fb7b23..53489947 100644 ggml_tensor * out = nullptr; test_glm_ffn_consumer( -@@ -9762,9 +9763,9 @@ struct test_glm_ffn_consumer : public test_case { +@@ -9765,9 +9766,9 @@ struct test_glm_ffn_consumer : public test_case { } } @@ -163,7 +163,7 @@ index 56fb7b23..53489947 100644 } ggml_tensor * build_graph(ggml_context * ctx) override { -@@ -9851,6 +9852,36 @@ struct test_glm_ffn_two_phase_consumer : public test_glm_ffn_consumer { +@@ -9854,6 +9855,36 @@ struct test_glm_ffn_two_phase_consumer : public test_glm_ffn_consumer { } }; @@ -200,7 +200,7 @@ index 56fb7b23..53489947 100644 struct test_glm_ffn_dual_lane_consumer : public test_glm_ffn_consumer { const bool dual_lane; const char * variant; -@@ -15150,6 +15181,7 @@ static std::vector> make_test_cases_eval() { +@@ -15154,6 +15185,7 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_glm_ffn_chain_native_down(true, "CANDIDATE")); test_cases.emplace_back(new test_glm_ffn_chain_two_phase(true, "CANDIDATE")); test_cases.emplace_back(new test_glm_ffn_two_phase_consumer(true, "CANDIDATE")); @@ -208,7 +208,7 @@ index 56fb7b23..53489947 100644 test_cases.emplace_back(new test_glm_ffn_dual_lane_consumer(true, "CANDIDATE")); return test_cases; -@@ -15520,6 +15552,7 @@ static std::vector> make_test_cases_perf() { +@@ -15524,6 +15556,7 @@ static std::vector> make_test_cases_perf() { GGML_TYPE_Q4_K, GGML_TYPE_Q4_K, GGML_TYPE_Q4_K, 6144, 2048, true, true)); test_cases.emplace_back(new test_glm_ffn_consumer(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q4_K, GGML_TYPE_Q4_K, 6144, 2048, true, false, true)); @@ -217,5 +217,5 @@ index 56fb7b23..53489947 100644 GGML_TYPE_Q4_K, GGML_TYPE_Q4_K, GGML_TYPE_Q4_K, 6144, 2048, true, false)); test_cases.emplace_back(new test_glm_ffn_consumer(GGML_TYPE_Q1_0, GGML_TYPE_Q1_0, GGML_TYPE_Q1_0, -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0032-ggml-skip-zero-weight-GLM-MoE-gate-up-slots.patch b/third_party/llama.cpp/patches/0030-ggml-skip-zero-weight-GLM-MoE-gate-up-slots.patch similarity index 91% rename from third_party/llama.cpp/patches/0032-ggml-skip-zero-weight-GLM-MoE-gate-up-slots.patch rename to third_party/llama.cpp/patches/0030-ggml-skip-zero-weight-GLM-MoE-gate-up-slots.patch index 2ed38ec876..2bb8c01fc4 100644 --- a/third_party/llama.cpp/patches/0032-ggml-skip-zero-weight-GLM-MoE-gate-up-slots.patch +++ b/third_party/llama.cpp/patches/0030-ggml-skip-zero-weight-GLM-MoE-gate-up-slots.patch @@ -1,14 +1,14 @@ -From 4523e19eecd43c47ca650939bb0c2dfc87c43920 Mon Sep 17 00:00:00 2001 +From 8c3753398c77d59baa5cad70aa89875354d9573e Mon Sep 17 00:00:00 2001 From: James Dumay Date: Sun, 19 Jul 2026 00:00:00 +1000 -Subject: [PATCH 27/36] ggml: skip zero-weight GLM MoE gate-up slots +Subject: [PATCH 30/45] ggml: skip zero-weight GLM MoE gate-up slots --- ggml/src/ggml-metal/ggml-metal.metal | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal -index b7b7e065..1b62d0dc 100644 +index 7a6df158..b5e37c11 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -10773,9 +10773,15 @@ void kernel_mul_mv_id_q2_K_gate_up_swiglu_pair_sg_dual_impl( @@ -41,5 +41,5 @@ index b7b7e065..1b62d0dc 100644 const float up_value = simd_sum(sum_up[row]); const float gate_value = simd_sum(sum_gate[row]); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0033-tests-cover-full-GLM-MoE-selected-chain.patch b/third_party/llama.cpp/patches/0031-tests-cover-full-GLM-MoE-selected-chain.patch similarity index 80% rename from third_party/llama.cpp/patches/0033-tests-cover-full-GLM-MoE-selected-chain.patch rename to third_party/llama.cpp/patches/0031-tests-cover-full-GLM-MoE-selected-chain.patch index 15808800a4..007865421a 100644 --- a/third_party/llama.cpp/patches/0033-tests-cover-full-GLM-MoE-selected-chain.patch +++ b/third_party/llama.cpp/patches/0031-tests-cover-full-GLM-MoE-selected-chain.patch @@ -1,17 +1,17 @@ -From 37ee00411144af6570074a352ac8584ab8527ad6 Mon Sep 17 00:00:00 2001 +From a00ce995595582f131c501bb930578b3e049f879 Mon Sep 17 00:00:00 2001 From: James Dumay Date: Sun, 19 Jul 2026 00:00:00 +1000 -Subject: [PATCH 28/36] tests: cover full GLM MoE selected chain +Subject: [PATCH 31/45] tests: cover full GLM MoE selected chain --- tests/test-backend-ops.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp -index 53489947..2aaded1c 100644 +index 50ec5225..32f1215d 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp -@@ -7824,7 +7824,9 @@ struct test_glm_moe_selected_chain : public test_case { +@@ -7826,7 +7826,9 @@ struct test_glm_moe_selected_chain : public test_case { } } @@ -22,7 +22,7 @@ index 53489947..2aaded1c 100644 bool eval_perf(ggml_backend_t backend, const char * op_names_filter, printer * output_printer) override { scoped_test_env motif_reference_env("GGML_METAL_ENABLE_GLM_MOE_DECODE_MOTIF_REFERENCE", "1"); -@@ -14945,6 +14947,7 @@ static std::vector> make_test_cases_eval() { +@@ -14949,6 +14951,7 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(true, "FIXED_IDS")); test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(false, "ROUTE_B")); test_cases.emplace_back(new test_glm_moe_selected_chain(1, false, "CORRECTNESS")); @@ -31,5 +31,5 @@ index 53489947..2aaded1c 100644 test_cases.emplace_back(new test_glm_attention_projection_chain(type, 2)); } -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0034-tests-add-GLM-Q2Q3-selected-weight-roofline.patch b/third_party/llama.cpp/patches/0032-tests-add-GLM-Q2Q3-selected-weight-roofline.patch similarity index 89% rename from third_party/llama.cpp/patches/0034-tests-add-GLM-Q2Q3-selected-weight-roofline.patch rename to third_party/llama.cpp/patches/0032-tests-add-GLM-Q2Q3-selected-weight-roofline.patch index 71db34b48a..4ca91999ac 100644 --- a/third_party/llama.cpp/patches/0034-tests-add-GLM-Q2Q3-selected-weight-roofline.patch +++ b/third_party/llama.cpp/patches/0032-tests-add-GLM-Q2Q3-selected-weight-roofline.patch @@ -1,17 +1,17 @@ -From e9f260788b857c0f86f41229b3bb2da52f3d2187 Mon Sep 17 00:00:00 2001 +From 223525c271739859819dd21e8164e3247196b7ad Mon Sep 17 00:00:00 2001 From: James Dumay Date: Sun, 19 Jul 2026 00:00:00 +1000 -Subject: [PATCH 29/36] tests: add GLM Q2Q3 selected weight roofline +Subject: [PATCH 32/45] tests: add GLM Q2Q3 selected weight roofline --- tests/test-backend-ops.cpp | 30 ++++++++++++++++++++++++++++++ 1 file changed, 30 insertions(+) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp -index 2aaded1c..c80ddaf9 100644 +index 32f1215d..9f566612 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp -@@ -7683,6 +7683,30 @@ struct test_glm_moe_q2_weight_roofline : public test_glm_moe_routed_q2_final_out +@@ -7685,6 +7685,30 @@ struct test_glm_moe_q2_weight_roofline : public test_glm_moe_routed_q2_final_out } }; @@ -42,7 +42,7 @@ index 2aaded1c..c80ddaf9 100644 struct test_glm_moe_selected_chain : public test_case { static constexpr int64_t n_expert = 32; static constexpr int64_t n_expert_used = 8; -@@ -14946,6 +14970,9 @@ static std::vector> make_test_cases_eval() { +@@ -14950,6 +14974,9 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(false, "ROUTE_A")); test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(true, "FIXED_IDS")); test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(false, "ROUTE_B")); @@ -52,7 +52,7 @@ index 2aaded1c..c80ddaf9 100644 test_cases.emplace_back(new test_glm_moe_selected_chain(1, false, "CORRECTNESS")); test_cases.emplace_back(new test_glm_moe_selected_chain(3, false, "USEFUL_A")); for (ggml_type type : {GGML_TYPE_Q8_0, GGML_TYPE_Q6_K, GGML_TYPE_Q5_K, GGML_TYPE_Q4_K}) { -@@ -15431,6 +15458,9 @@ static std::vector> make_test_cases_perf() { +@@ -15435,6 +15462,9 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(false, "ROUTE_A")); test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(true, "FIXED_IDS")); test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(false, "ROUTE_B")); @@ -63,5 +63,5 @@ index 2aaded1c..c80ddaf9 100644 test_cases.emplace_back(new test_glm_moe_selected_chain(3, true, "BYTE84")); test_cases.emplace_back(new test_glm_moe_selected_chain(3, true, "BYTE128", 128)); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0035-ggml-add-active-count-Q3-GLM-MoE-down-kernels.patch b/third_party/llama.cpp/patches/0033-ggml-add-active-count-Q3-GLM-MoE-down-kernels.patch similarity index 97% rename from third_party/llama.cpp/patches/0035-ggml-add-active-count-Q3-GLM-MoE-down-kernels.patch rename to third_party/llama.cpp/patches/0033-ggml-add-active-count-Q3-GLM-MoE-down-kernels.patch index c908a18f9c..9137658be6 100644 --- a/third_party/llama.cpp/patches/0035-ggml-add-active-count-Q3-GLM-MoE-down-kernels.patch +++ b/third_party/llama.cpp/patches/0033-ggml-add-active-count-Q3-GLM-MoE-down-kernels.patch @@ -1,7 +1,7 @@ -From 66c8184a2a33a5fb60bfcadc47d4f8088669a5fe Mon Sep 17 00:00:00 2001 +From 83bd6794157379b2ac6fb6121d3940219a2ce5f3 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 01:02:36 +1000 -Subject: [PATCH 30/36] ggml: add active-count Q3 GLM MoE down kernels +Subject: [PATCH 33/45] ggml: add active-count Q3 GLM MoE down kernels --- ggml/src/ggml-metal/ggml-metal-device.cpp | 49 +++++++++++ @@ -224,7 +224,7 @@ index ab41e1ec..b4fc441a 100644 ggml_metal_encoder_set_pipeline(enc, pipeline_slot_parallel); ggml_metal_encoder_set_bytes(enc, &args, sizeof(args), 0); diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal -index 1b62d0dc..77119fc9 100644 +index b5e37c11..000fe0b2 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -11754,7 +11754,7 @@ kernel void kernel_mul_mv_id_q3_K_weighted_reduce( @@ -310,10 +310,10 @@ index 1b62d0dc..77119fc9 100644 } diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp -index c80ddaf9..3a2448a8 100644 +index 9f566612..f0b94ded 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp -@@ -8919,6 +8919,15 @@ struct test_glm_moe_static_motif : public test_case { +@@ -8922,6 +8922,15 @@ struct test_glm_moe_static_motif : public test_case { } return weighted_down ? "GLM_MOE_STATIC_MOTIF_Q2_WEIGHTED" : "GLM_MOE_STATIC_MOTIF_Q2"; } @@ -329,7 +329,7 @@ index c80ddaf9..3a2448a8 100644 if (gate_type == GGML_TYPE_Q4_K && up_type == GGML_TYPE_Q4_K && down_type == GGML_TYPE_Q4_K) { if (n_embd == 6144 && n_ff == 2048) { return "GLM_MOE_STATIC_MOTIF_Q4_GLM"; -@@ -9085,6 +9094,7 @@ struct test_glm_moe_down_reduce : public test_case { +@@ -9088,6 +9097,7 @@ struct test_glm_moe_down_reduce : public test_case { const int64_t n_ff; const bool already_weighted; const bool f16_act; @@ -337,7 +337,7 @@ index c80ddaf9..3a2448a8 100644 ggml_tensor * selected_experts = nullptr; ggml_tensor * weights = nullptr; -@@ -9097,23 +9107,38 @@ struct test_glm_moe_down_reduce : public test_case { +@@ -9100,23 +9110,38 @@ struct test_glm_moe_down_reduce : public test_case { int64_t n_embd = 6144, int64_t n_ff = 2048, bool already_weighted = false, @@ -379,7 +379,7 @@ index c80ddaf9..3a2448a8 100644 if (f16_act) { return "GLM_MOE_DOWN_REDUCE_F16_ACT"; } -@@ -9197,7 +9222,9 @@ struct test_glm_moe_down_reduce : public test_case { +@@ -9200,7 +9225,9 @@ struct test_glm_moe_down_reduce : public test_case { for (int64_t token = 0; token < n_tokens; ++token) { float sum = 0.0f; for (int64_t slot = 0; slot < n_expert_used; ++slot) { @@ -390,7 +390,7 @@ index c80ddaf9..3a2448a8 100644 data[token*n_expert_used + slot] = value; sum += value; } -@@ -15033,6 +15060,8 @@ static std::vector> make_test_cases_eval() { +@@ -15037,6 +15064,8 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_glm_moe_gate_up_swiglu_rowtile_q2(false, true)); test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 256, 256, false)); test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 256, 256, true)); @@ -399,7 +399,7 @@ index c80ddaf9..3a2448a8 100644 test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, true, true, false, 2)); test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, true, true, false, 4)); test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, true, true, false, 6)); -@@ -15053,6 +15082,9 @@ static std::vector> make_test_cases_eval() { +@@ -15057,6 +15086,9 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 256, 256)); test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 256, 256, true)); test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048)); @@ -409,7 +409,7 @@ index c80ddaf9..3a2448a8 100644 test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048, true)); test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q2_K, 6144, 2048)); test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q2_K, 6144, 2048, true)); -@@ -15517,6 +15549,7 @@ static std::vector> make_test_cases_perf() { +@@ -15521,6 +15553,7 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, false)); test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, true)); test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, false, true)); @@ -417,7 +417,7 @@ index c80ddaf9..3a2448a8 100644 test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, true, true, false, 2)); test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, true, true, false, 4)); test_cases.emplace_back(new test_glm_moe_static_motif(GGML_TYPE_Q2_K, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, 6144, 2048, true, true, false, 6)); -@@ -15539,6 +15572,9 @@ static std::vector> make_test_cases_perf() { +@@ -15543,6 +15576,9 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 2048, 1024)); test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 2048, 1024, true)); test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048)); @@ -428,5 +428,5 @@ index c80ddaf9..3a2448a8 100644 test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_TQ2_0, 6144, 2048)); test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q2_K, 6144, 2048)); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0036-ggml-use-GLM-max-active-policy-for-Q3-down-kernels.patch b/third_party/llama.cpp/patches/0034-ggml-use-GLM-max-active-policy-for-Q3-down-kernels.patch similarity index 95% rename from third_party/llama.cpp/patches/0036-ggml-use-GLM-max-active-policy-for-Q3-down-kernels.patch rename to third_party/llama.cpp/patches/0034-ggml-use-GLM-max-active-policy-for-Q3-down-kernels.patch index 24832c6f74..91829fed15 100644 --- a/third_party/llama.cpp/patches/0036-ggml-use-GLM-max-active-policy-for-Q3-down-kernels.patch +++ b/third_party/llama.cpp/patches/0034-ggml-use-GLM-max-active-policy-for-Q3-down-kernels.patch @@ -1,7 +1,7 @@ -From 3fc13c2ed2cecf50726e7d10927f5070539c9aa1 Mon Sep 17 00:00:00 2001 +From d8a655fca4f47b48fc617b3302beb3403b743593 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 05:19:37 +1000 -Subject: [PATCH 31/36] ggml: use GLM max-active policy for Q3 down kernels +Subject: [PATCH 34/45] ggml: use GLM max-active policy for Q3 down kernels --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 30 ++++++++++++++++++++++++++ @@ -85,10 +85,10 @@ index b4fc441a..2e42a168 100644 pipeline_slot_parallel = ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_f16(lib, weighted_sum); diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp -index 3a2448a8..44f36f73 100644 +index f0b94ded..12e2059f 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp -@@ -7331,6 +7331,7 @@ struct test_topk_moe_glm_component : public test_case { +@@ -7333,6 +7333,7 @@ struct test_topk_moe_glm_component : public test_case { getenv("SKIPPY_GLM_DSA_EXPERIMENTAL_Q3_DOWN_SLOT_PARALLEL_REDUCE_R8") != nullptr || getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_R8_NB8") != nullptr || getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_R8_NB8_W0") != nullptr || @@ -96,7 +96,7 @@ index 3a2448a8..44f36f73 100644 getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_R12_NB8_W0") != nullptr || getenv("SKIPPY_GLM_DSA_EXPERIMENTAL_Q3_DOWN_SLOT_PARALLEL_REDUCE_R16") != nullptr || getenv("SKIPPY_GLM_DSA_EXPERIMENTAL_Q3_DOWN_SLOT_SPLIT2_REDUCE") != nullptr || -@@ -9021,6 +9022,7 @@ struct test_glm_moe_static_motif : public test_case { +@@ -9024,6 +9025,7 @@ struct test_glm_moe_static_motif : public test_case { getenv("SKIPPY_GLM_DSA_EXPERIMENTAL_Q3_DOWN_SLOT_PARALLEL_REDUCE_R8") != nullptr || getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_R8_NB8") != nullptr || getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_R8_NB8_W0") != nullptr || @@ -105,5 +105,5 @@ index 3a2448a8..44f36f73 100644 getenv("SKIPPY_GLM_DSA_EXPERIMENTAL_Q3_DOWN_SLOT_PARALLEL_REDUCE_R16") != nullptr || getenv("SKIPPY_GLM_DSA_EXPERIMENTAL_Q3_DOWN_SLOT_SPLIT2_REDUCE") != nullptr || -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0037-ggml-honor-explicit-Q2-gate-up-Metal-variant-flags.patch b/third_party/llama.cpp/patches/0035-ggml-honor-explicit-Q2-gate-up-Metal-variant-flags.patch similarity index 92% rename from third_party/llama.cpp/patches/0037-ggml-honor-explicit-Q2-gate-up-Metal-variant-flags.patch rename to third_party/llama.cpp/patches/0035-ggml-honor-explicit-Q2-gate-up-Metal-variant-flags.patch index c2d92b5d5b..46dd3dcddf 100644 --- a/third_party/llama.cpp/patches/0037-ggml-honor-explicit-Q2-gate-up-Metal-variant-flags.patch +++ b/third_party/llama.cpp/patches/0035-ggml-honor-explicit-Q2-gate-up-Metal-variant-flags.patch @@ -1,7 +1,7 @@ -From 8fec61c2696f179dbe598552bef7da0237ebedbb Mon Sep 17 00:00:00 2001 +From f8152becddf60982cc3a69f9db8baecd399bdacb Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 10:44:00 +1000 -Subject: [PATCH 32/36] ggml: honor explicit Q2 gate/up Metal variant flags +Subject: [PATCH 35/45] ggml: honor explicit Q2 gate/up Metal variant flags --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 12 +++++++++++- @@ -38,5 +38,5 @@ index 2e42a168..9469192d 100644 pair_sg_slot8_requested; const bool pair_sg_slot2 = pair_sg && dst->type == GGML_TYPE_F32 && -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0038-ggml-make-GLM-MoE-roofline-honor-active-experts.patch b/third_party/llama.cpp/patches/0036-ggml-make-GLM-MoE-roofline-honor-active-experts.patch similarity index 95% rename from third_party/llama.cpp/patches/0038-ggml-make-GLM-MoE-roofline-honor-active-experts.patch rename to third_party/llama.cpp/patches/0036-ggml-make-GLM-MoE-roofline-honor-active-experts.patch index 85ef7a71ae..82b1789c4c 100644 --- a/third_party/llama.cpp/patches/0038-ggml-make-GLM-MoE-roofline-honor-active-experts.patch +++ b/third_party/llama.cpp/patches/0036-ggml-make-GLM-MoE-roofline-honor-active-experts.patch @@ -1,7 +1,7 @@ -From 6208005ffe6bdc5b2d187f2d6c31138773bfdbc9 Mon Sep 17 00:00:00 2001 +From 36ee61b78bfa044c8701f3db79d8bc51bc1ca3d3 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 13:10:00 +1000 -Subject: [PATCH 33/36] ggml: make GLM MoE roofline honor active experts +Subject: [PATCH 36/45] ggml: make GLM MoE roofline honor active experts --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 18 ++++++++++++------ @@ -61,5 +61,5 @@ index 9469192d..2121075e 100644 return motif.n_fuse; -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0039-ggml-avoid-GLM-Q3-fused-tail-under-active-policy.patch b/third_party/llama.cpp/patches/0037-ggml-avoid-GLM-Q3-fused-tail-under-active-policy.patch similarity index 92% rename from third_party/llama.cpp/patches/0039-ggml-avoid-GLM-Q3-fused-tail-under-active-policy.patch rename to third_party/llama.cpp/patches/0037-ggml-avoid-GLM-Q3-fused-tail-under-active-policy.patch index 6bb6fadf3e..89022fc2fa 100644 --- a/third_party/llama.cpp/patches/0039-ggml-avoid-GLM-Q3-fused-tail-under-active-policy.patch +++ b/third_party/llama.cpp/patches/0037-ggml-avoid-GLM-Q3-fused-tail-under-active-policy.patch @@ -1,7 +1,7 @@ -From dd76e903b54de937336fbea5f1758753a379109f Mon Sep 17 00:00:00 2001 +From 3bcadd4535d6e48229e2edf4a10e849b23475083 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 14:05:00 +1000 -Subject: [PATCH 34/36] ggml: avoid GLM Q3 fused tail under active policy +Subject: [PATCH 37/45] ggml: avoid GLM Q3 fused tail under active policy --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 3 ++ @@ -23,10 +23,10 @@ index 2121075e..91976954 100644 motif.gate == nullptr || motif.up == nullptr || motif.glu == nullptr || motif.down == nullptr || motif.out == nullptr || diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp -index 44f36f73..1bbd3713 100644 +index 12e2059f..43ce1a0e 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp -@@ -7660,6 +7660,44 @@ struct test_glm_moe_routed_q2_final_output : public test_topk_moe_glm_component +@@ -7662,6 +7662,44 @@ struct test_glm_moe_routed_q2_final_output : public test_topk_moe_glm_component } }; @@ -71,7 +71,7 @@ index 44f36f73..1bbd3713 100644 struct test_glm_moe_q2_weight_roofline : public test_glm_moe_routed_q2_final_output { const bool bypass_route; const char * variant; -@@ -14996,6 +15034,10 @@ static std::vector> make_test_cases_eval() { +@@ -15000,6 +15038,10 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_glm_moe_routed_q2_final_output(true, false)); test_cases.emplace_back(new test_glm_moe_routed_q2_final_output(false, true)); test_cases.emplace_back(new test_glm_moe_routed_q2_final_output(true, true)); @@ -82,7 +82,7 @@ index 44f36f73..1bbd3713 100644 test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(false, "ROUTE_A")); test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(true, "FIXED_IDS")); test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(false, "ROUTE_B")); -@@ -15489,6 +15531,10 @@ static std::vector> make_test_cases_perf() { +@@ -15493,6 +15535,10 @@ static std::vector> make_test_cases_perf() { test_cases.emplace_back(new test_glm_moe_routed_q2_final_output(true, false)); test_cases.emplace_back(new test_glm_moe_routed_q2_final_output(false, true)); test_cases.emplace_back(new test_glm_moe_routed_q2_final_output(true, true)); @@ -94,5 +94,5 @@ index 44f36f73..1bbd3713 100644 test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(true, "FIXED_IDS")); test_cases.emplace_back(new test_glm_moe_q2_weight_roofline(false, "ROUTE_B")); -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0040-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch b/third_party/llama.cpp/patches/0038-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch similarity index 91% rename from third_party/llama.cpp/patches/0040-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch rename to third_party/llama.cpp/patches/0038-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch index 088eff5174..5c293f158e 100644 --- a/third_party/llama.cpp/patches/0040-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch +++ b/third_party/llama.cpp/patches/0038-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch @@ -1,7 +1,7 @@ -From aacb3421872735fcd2f1a22e4560a60b44f53046 Mon Sep 17 00:00:00 2001 +From 9fa1c53b208fec4bf4ac863c38135451a1fca6f3 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 03:51:17 +1000 -Subject: [PATCH 35/36] ggml: shrink GLM Q2 gate-up active dispatch +Subject: [PATCH 38/45] ggml: shrink GLM Q2 gate-up active dispatch --- ggml/src/ggml-metal/ggml-metal-ops.cpp | 6 ++++-- @@ -27,7 +27,7 @@ index 91976954..646c983a 100644 const int grid_y = (1 + nr1 - 1)/nr1; const int grid_z = ne123; diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal -index 77119fc9..338e8f3e 100644 +index 000fe0b2..5dfeed66 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -10331,9 +10331,11 @@ void kernel_mul_mv_id_q2_K_gate_up_swiglu_pair_sg_impl( @@ -45,5 +45,5 @@ index 77119fc9..338e8f3e 100644 const bool is_gate = (sgitg & 1) != 0; const int first_row = r0 * nr0; -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0041-ggml-skip-inactive-GLM-Q3-down-slots.patch b/third_party/llama.cpp/patches/0039-ggml-skip-inactive-GLM-Q3-down-slots.patch similarity index 82% rename from third_party/llama.cpp/patches/0041-ggml-skip-inactive-GLM-Q3-down-slots.patch rename to third_party/llama.cpp/patches/0039-ggml-skip-inactive-GLM-Q3-down-slots.patch index 249be4fc9b..7ba0a89cd2 100644 --- a/third_party/llama.cpp/patches/0041-ggml-skip-inactive-GLM-Q3-down-slots.patch +++ b/third_party/llama.cpp/patches/0039-ggml-skip-inactive-GLM-Q3-down-slots.patch @@ -1,14 +1,14 @@ -From 00fa96c1d75aa4ee9885eb64d88cbebe861d9fe5 Mon Sep 17 00:00:00 2001 +From eb68fbbabf64915d6283683ae25fefd7475d46c2 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Mon, 20 Jul 2026 05:30:42 +1000 -Subject: [PATCH 36/36] ggml: skip inactive GLM Q3 down slots +Subject: [PATCH 39/45] ggml: skip inactive GLM Q3 down slots --- ggml/src/ggml-metal/ggml-metal.metal | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal -index 338e8f3e..3bdec47f 100644 +index 5dfeed66..b098231d 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -11802,7 +11802,7 @@ void kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_impl( @@ -21,5 +21,5 @@ index 338e8f3e..3bdec47f 100644 if (expert_id >= 0 && expert_id < args.ne02) { const uint64_t offset0 = uint64_t(expert_id)*args.nb02 + uint64_t(first_row)*args.nb01; -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0042-Fix-GLM-DSA-native-MTP-execution.patch b/third_party/llama.cpp/patches/0040-Fix-GLM-DSA-native-MTP-execution.patch similarity index 95% rename from third_party/llama.cpp/patches/0042-Fix-GLM-DSA-native-MTP-execution.patch rename to third_party/llama.cpp/patches/0040-Fix-GLM-DSA-native-MTP-execution.patch index ce33e44923..457cb41f9d 100644 --- a/third_party/llama.cpp/patches/0042-Fix-GLM-DSA-native-MTP-execution.patch +++ b/third_party/llama.cpp/patches/0040-Fix-GLM-DSA-native-MTP-execution.patch @@ -1,14 +1,14 @@ -From 0489d807735c3eb284cd072934bc94f0b6919f95 Mon Sep 17 00:00:00 2001 +From 6b9709b217341e1229a3a56b1a6132b7391c1a49 Mon Sep 17 00:00:00 2001 From: Mesh-LLM CI Date: Wed, 22 Jul 2026 08:41:09 +1000 -Subject: [PATCH] Fix GLM DSA native MTP execution +Subject: [PATCH 40/45] Fix GLM DSA native MTP execution --- src/llama-context.cpp | 16 ++ src/llama-graph.cpp | 20 ++- src/models/glm-dsa.cpp | 321 +++++++++++++++++++++++++++++++++++++++++ - src/models/models.h | 4 +- - 4 files changed, 355 insertions(+), 6 deletions(-) + src/models/models.h | 5 +- + 4 files changed, 356 insertions(+), 6 deletions(-) diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 9ab7be4b..4342e96e 100644 @@ -73,11 +73,11 @@ index 244a400a..6cef391d 100644 bool llm_graph_input_attn_k_dsa::can_reuse(const llm_graph_params & params) { diff --git a/src/models/glm-dsa.cpp b/src/models/glm-dsa.cpp -index adde84fd..16d605c4 100644 +index 82c6a093..32996613 100644 --- a/src/models/glm-dsa.cpp +++ b/src/models/glm-dsa.cpp -@@ -147,6 +147,327 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { - } +@@ -188,6 +188,327 @@ std::unique_ptr llama_model_glm_dsa::build_arch_graph(const l + return std::make_unique(*this, params); } +llama_model_glm_dsa::graph_mtp::graph_mtp(const llama_model & model, const llm_graph_params & params) : @@ -181,7 +181,7 @@ index adde84fd..16d605c4 100644 + cb(indexer_q_nope, "mtp_indexer_q_nope", il); + + indexer_q_pe = ggml_rope_ext(ctx0, indexer_q_pe, inp_pos, nullptr, n_rot, -+ LLAMA_ROPE_TYPE_NEOX, n_ctx_orig, freq_base, freq_scale, ++ LLAMA_ROPE_TYPE_NORM, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + cb(indexer_q_pe, "mtp_indexer_q_pe", il); + @@ -208,7 +208,7 @@ index adde84fd..16d605c4 100644 + cb(indexer_k_nope, "mtp_indexer_k_nope", il); + + indexer_k_pe = ggml_rope_ext(ctx0, indexer_k_pe, inp_pos, nullptr, n_rot, -+ LLAMA_ROPE_TYPE_NEOX, n_ctx_orig, freq_base, freq_scale, ++ LLAMA_ROPE_TYPE_NORM, n_ctx_orig, freq_base, freq_scale, + ext_factor, attn_factor, beta_fast, beta_slow); + cb(indexer_k_pe, "mtp_indexer_k_pe", il); + @@ -405,20 +405,24 @@ index adde84fd..16d605c4 100644 llm_graph_context(params) { const bool is_mla = hparams.is_mla(); diff --git a/src/models/models.h b/src/models/models.h -index d8120d9d..fa1e36c1 100644 +index f3b576c0..ca471dcc 100644 --- a/src/models/models.h +++ b/src/models/models.h -@@ -1221,7 +1221,9 @@ struct llama_model_glm_dsa : public llama_model_base { +@@ -1221,10 +1221,13 @@ struct llama_model_glm_dsa : public llama_model_base { void load_arch_hparams(llama_model_loader & ml) override; void load_arch_tensors(llama_model_loader & ml) override; -- using graph_mtp = llama_model_deepseek2::graph_mtp; + struct graph_mtp : public llm_graph_context { + graph_mtp(const llama_model & model, const llm_graph_params & params); + }; - ++ struct graph : public llm_graph_context { graph(const llama_model & model, const llm_graph_params & params); + }; +- using graph_mtp = llama_model_deepseek2::graph_mtp; + + std::unique_ptr build_arch_graph(const llm_graph_params & params) const override; + }; -- -2.54.0 (Apple Git-157) +2.50.1 (Apple Git-155) diff --git a/third_party/llama.cpp/patches/0041-Adapt-Skippy-model-loading-to-load-modes.patch b/third_party/llama.cpp/patches/0041-Adapt-Skippy-model-loading-to-load-modes.patch new file mode 100644 index 0000000000..d2a9becf94 --- /dev/null +++ b/third_party/llama.cpp/patches/0041-Adapt-Skippy-model-loading-to-load-modes.patch @@ -0,0 +1,31 @@ +From 43fe1861b1b58d04444da0886fb5fddd1a14f05c Mon Sep 17 00:00:00 2001 +From: Mesh-LLM CI +Date: Sun, 26 Jul 2026 19:30:12 +1000 +Subject: [PATCH 41/45] Adapt Skippy model loading to load modes + +--- + src/skippy.cpp | 7 +++---- + 1 file changed, 3 insertions(+), 4 deletions(-) + +diff --git a/src/skippy.cpp b/src/skippy.cpp +index a8cc03cb..9bd81828 100644 +--- a/src/skippy.cpp ++++ b/src/skippy.cpp +@@ -950,11 +950,10 @@ static void skippy_apply_model_load_memory_config( + if (config == nullptr) { + return; + } +- if (config->has_mmap_override) { +- params.use_mmap = config->use_mmap; +- } + if (config->use_mlock) { +- params.use_mlock = true; ++ params.load_mode = LLAMA_LOAD_MODE_MLOCK; ++ } else if (config->has_mmap_override) { ++ params.load_mode = config->use_mmap ? LLAMA_LOAD_MODE_MMAP : LLAMA_LOAD_MODE_NONE; + } + } + +-- +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/patches/0042-Harden-staged-session-and-sideband-bookkeeping.patch b/third_party/llama.cpp/patches/0042-Harden-staged-session-and-sideband-bookkeeping.patch new file mode 100644 index 0000000000..91f747296b --- /dev/null +++ b/third_party/llama.cpp/patches/0042-Harden-staged-session-and-sideband-bookkeeping.patch @@ -0,0 +1,129 @@ +From 66275c899d65c5fb9c8f44efaf2f4e423bac045c Mon Sep 17 00:00:00 2001 +From: Mesh-LLM CI +Date: Sun, 26 Jul 2026 21:13:49 +1000 +Subject: [PATCH 42/45] Harden staged session and sideband bookkeeping + +--- + ggml/include/ggml-rpc.h | 2 +- + src/skippy.cpp | 41 ++++++++++++++++++----------------------- + 2 files changed, 19 insertions(+), 24 deletions(-) + +diff --git a/ggml/include/ggml-rpc.h b/ggml/include/ggml-rpc.h +index 82d27410..83936373 100644 +--- a/ggml/include/ggml-rpc.h ++++ b/ggml/include/ggml-rpc.h +@@ -8,7 +8,7 @@ extern "C" { + + #define RPC_PROTO_MAJOR_VERSION 4 + #define RPC_PROTO_MINOR_VERSION 0 +-#define RPC_PROTO_PATCH_VERSION 3 ++#define RPC_PROTO_PATCH_VERSION 4 + + #ifdef __cplusplus + static_assert(GGML_OP_COUNT == 107, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); +diff --git a/src/skippy.cpp b/src/skippy.cpp +index 9bd81828..47f0a549 100644 +--- a/src/skippy.cpp ++++ b/src/skippy.cpp +@@ -1410,27 +1410,13 @@ static uint32_t skippy_glm_dsa_top_k_count_for_sideband_bytes(size_t token_count + } + + static bool skippy_glm_dsa_sideband_top_k_count_is_valid( +- size_t token_count, + uint32_t actual_top_k, + uint32_t expected_top_k) { +- if (actual_top_k == 0 || expected_top_k == 0) { +- return false; +- } +- +- // Prefill sidebands can be chunk-local: llama.cpp emits the actual +- // per-chunk top-k width, while the receiving staged session may have a +- // larger accumulated KV prefix. +- if (token_count > 1) { +- return actual_top_k <= expected_top_k; +- } +- +- // Decode producers export the actual top-k tensor width from their graph. +- // The receiving session can have a larger accumulated KV prefix than the +- // producer's current sideband width, especially when an OpenAI generation +- // loop advances through cached or fused decode steps. The sideband is still +- // valid as long as it is non-empty, token-major, and does not claim more +- // rows than the receiver can make visible. +- return actual_top_k <= expected_top_k; ++ // Sidebands can be chunk-local: the producer exports the actual tensor ++ // width, while the receiving staged session may have a larger accumulated ++ // KV prefix. The sideband is valid when it is non-empty and does not claim ++ // more rows than the receiver can make visible. ++ return actual_top_k != 0 && actual_top_k <= expected_top_k; + } + + static size_t skippy_glm_dsa_top_k_bytes(const skippy_session * session, size_t token_count) { +@@ -1611,7 +1597,7 @@ static enum skippy_status skippy_validate_frame_input_sequences( + return SKIPPY_STATUS_INVALID_ARGUMENT; + } + const uint32_t expected_top_k = skippy_glm_dsa_sideband_top_k_count(session, expected_token_count); +- if (!skippy_glm_dsa_sideband_top_k_count_is_valid(expected_token_count, actual_top_k, expected_top_k)) { ++ if (!skippy_glm_dsa_sideband_top_k_count_is_valid(actual_top_k, expected_top_k)) { + char message[256]; + std::snprintf( + message, +@@ -2989,6 +2975,9 @@ static enum skippy_status skippy_verify_token_batch( + SKIPPY_GLM_DSA_PHASE_HINT_VERIFY, + out_error); + llama_batch_free(batch); ++ if (status == SKIPPY_STATUS_OK) { ++ skippy_record_tokens(session, token_ids, token_count); ++ } + return status; + } + +@@ -3880,6 +3869,9 @@ static enum skippy_status skippy_decode_activation_frame( + if (status == SKIPPY_STATUS_OK && token_ids != nullptr) { + status = skippy_mtp_sync_target_tokens(session, token_ids, token_count, token_start, out_error); + } ++ if (status == SKIPPY_STATUS_OK && token_ids != nullptr) { ++ skippy_record_tokens(session, token_ids, token_count); ++ } + return status; + } + +@@ -5078,6 +5070,7 @@ enum skippy_status skippy_decode_batch_sampled( + } + + const int32_t n_tokens = static_cast(request_count); ++ GGML_ASSERT(static_cast(n_tokens) == request_count); + llama_batch batch = llama_batch_init(n_tokens, 0, 1); + batch.n_tokens = n_tokens; + for (int32_t i = 0; i < n_tokens; ++i) { +@@ -5586,7 +5579,7 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( + if (activation_input) { + embd_storage.resize(static_cast(n_tokens)*n_embd_inp); + if ((batch_input_flags & SKIPPY_ACTIVATION_FLAG_GLM_DSA_TOP_K) != 0) { +- glm_dsa_input_top_k_storage.resize(static_cast(n_tokens)*batch_input_top_k); ++ glm_dsa_input_top_k_storage.resize(request_count*batch_input_top_k); + } + for (int32_t i = 0; i < n_tokens; ++i) { + const skippy_activation_desc * input_desc = input_descs[i]; +@@ -5648,7 +5641,7 @@ enum skippy_status skippy_decode_step_frame_batch_sampled( + const llama_pos batch_pos_start = pos_storage.empty() ? first->n_past : pos_storage[0]; + skippy_glm_dsa_top_k_scope glm_dsa_top_k_scope( + glm_dsa_input_top_k_storage.empty() ? nullptr : glm_dsa_input_top_k_storage.data(), +- static_cast(n_tokens), ++ static_cast(request_count), + batch_input_top_k, + 1, + batch_pos_start); +@@ -6384,7 +6377,9 @@ enum skippy_status skippy_import_kv_page( + // KV residency determines graph topology for paths such as GLM-DSA, where + // the selected-index width is capped by the visible KV length. Do not reuse + // a graph that was built before this page was imported. +- session->ctx->get_gf_res_prev()->reset(); ++ if (llm_graph_result * graph = session->ctx->get_gf_res_prev()) { ++ graph->reset(); ++ } + session->n_past = std::max( + session->n_past, + static_cast(std::min( +-- +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/patches/0043-Correct-GLM-DSA-staged-graph-contracts.patch b/third_party/llama.cpp/patches/0043-Correct-GLM-DSA-staged-graph-contracts.patch new file mode 100644 index 0000000000..96f20bd2b6 --- /dev/null +++ b/third_party/llama.cpp/patches/0043-Correct-GLM-DSA-staged-graph-contracts.patch @@ -0,0 +1,82 @@ +From abb3456459c7535ec1e09286c8013a60c1170efe Mon Sep 17 00:00:00 2001 +From: Mesh-LLM CI +Date: Sun, 26 Jul 2026 21:14:11 +1000 +Subject: [PATCH 43/45] Correct GLM-DSA staged graph contracts + +--- + src/models/glm-dsa.cpp | 21 ++++++++------------- + 1 file changed, 8 insertions(+), 13 deletions(-) + +diff --git a/src/models/glm-dsa.cpp b/src/models/glm-dsa.cpp +index 32996613..2ddbfbe1 100644 +--- a/src/models/glm-dsa.cpp ++++ b/src/models/glm-dsa.cpp +@@ -122,13 +122,8 @@ void llama_model_glm_dsa::load_arch_tensors(llama_model_loader &) { + + layer.wkv_a_mqa = create_tensor(tn(LLM_TENSOR_ATTN_KV_A_MQA, "weight", i), {n_embd, kv_lora_rank + n_embd_head_qk_rope}, flags); + +- layer.wk_b = create_tensor(tn(LLM_TENSOR_ATTN_K_B, "weight", i), {n_embd_head_qk_nope, kv_lora_rank, n_head}, flags | TENSOR_NOT_REQUIRED); +- layer.wv_b = create_tensor(tn(LLM_TENSOR_ATTN_V_B, "weight", i), {kv_lora_rank, n_embd_head_v_mla, n_head}, flags | TENSOR_NOT_REQUIRED); +- if (!layer.wk_b || !layer.wv_b) { +- layer.wk_b = nullptr; +- layer.wv_b = nullptr; +- layer.wkv_b = create_tensor(tn(LLM_TENSOR_ATTN_KV_B, "weight", i), {kv_lora_rank, n_head * (n_embd_head_qk_nope + n_embd_head_v_mla)}, flags); +- } ++ layer.wk_b = create_tensor(tn(LLM_TENSOR_ATTN_K_B, "weight", i), {n_embd_head_qk_nope, kv_lora_rank, n_head}, flags); ++ layer.wv_b = create_tensor(tn(LLM_TENSOR_ATTN_V_B, "weight", i), {kv_lora_rank, n_embd_head_v_mla, n_head}, flags); + + layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_head * n_embd_head_v_mla, n_embd}, flags); + +@@ -285,7 +280,7 @@ llama_model_glm_dsa::graph_mtp::graph_mtp(const llama_model & model, const llm_g + ggml_view_3d(ctx0, indexer_q, n_embd_indexer_head_nope, n_indexer_head, n_tokens, + ggml_row_size(indexer_q->type, n_embd_indexer_head), + ggml_row_size(indexer_q->type, n_embd_indexer_head) * n_indexer_head, +- ggml_row_size(indexer_q->type, n_embd_indexer_head_nope)); ++ ggml_row_size(indexer_q->type, n_embd_indexer_head_rope)); + cb(indexer_q_nope, "mtp_indexer_q_nope", il); + + indexer_q_pe = ggml_rope_ext(ctx0, indexer_q_pe, inp_pos, nullptr, n_rot, +@@ -312,7 +307,7 @@ llama_model_glm_dsa::graph_mtp::graph_mtp(const llama_model & model, const llm_g + ggml_view_3d(ctx0, indexer_k, n_embd_indexer_head_nope, 1, n_tokens, + ggml_row_size(indexer_k->type, n_embd_indexer_head), + ggml_row_size(indexer_k->type, n_embd_indexer_head), +- ggml_row_size(indexer_k->type, n_embd_indexer_head_nope)); ++ ggml_row_size(indexer_k->type, n_embd_indexer_head_rope)); + cb(indexer_k_nope, "mtp_indexer_k_nope", il); + + indexer_k_pe = ggml_rope_ext(ctx0, indexer_k_pe, inp_pos, nullptr, n_rot, +@@ -610,7 +605,7 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par + ggml_view_3d(ctx0, indexer_q, n_embd_indexer_head_nope, n_indexer_head, n_tokens, + ggml_row_size(indexer_q->type, n_embd_indexer_head), + ggml_row_size(indexer_q->type, n_embd_indexer_head) * n_indexer_head, +- ggml_row_size(indexer_q->type, n_embd_indexer_head_nope)); ++ ggml_row_size(indexer_q->type, n_embd_indexer_head_rope)); + cb(indexer_q_nope, "indexer_q_nope", il); + + indexer_q_pe = ggml_rope_ext(ctx0, indexer_q_pe, inp_pos, nullptr, n_rot, +@@ -640,7 +635,7 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par + ggml_view_3d(ctx0, indexer_k, n_embd_indexer_head_nope, 1, n_tokens, + ggml_row_size(indexer_k->type, n_embd_indexer_head), + ggml_row_size(indexer_k->type, n_embd_indexer_head) * 1, +- ggml_row_size(indexer_k->type, n_embd_indexer_head_nope)); ++ ggml_row_size(indexer_k->type, n_embd_indexer_head_rope)); + cb(indexer_k_nope, "indexer_k_nope", il); + + indexer_k_pe = ggml_rope_ext(ctx0, indexer_k_pe, inp_pos, nullptr, n_rot, +@@ -882,11 +877,11 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par + return; + } + +- cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1); +- + cb(cur, "h_nextn", -1); + res->t_h_nextn = cur; + ++ cur = build_norm(cur, model.output_norm, NULL, LLM_NORM_RMS, -1); ++ + if (!cparams.embeddings_nextn_masked && inp_out_ids) { + cur = ggml_get_rows(ctx0, cur, inp_out_ids); + } +-- +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/patches/0044-Fix-GLM-DSA-Metal-dispatch-safety.patch b/third_party/llama.cpp/patches/0044-Fix-GLM-DSA-Metal-dispatch-safety.patch new file mode 100644 index 0000000000..6ed3d91888 --- /dev/null +++ b/third_party/llama.cpp/patches/0044-Fix-GLM-DSA-Metal-dispatch-safety.patch @@ -0,0 +1,213 @@ +From e47c41e7d0e39fbc95f059678b01653d813f0e6a Mon Sep 17 00:00:00 2001 +From: Mesh-LLM CI +Date: Sun, 26 Jul 2026 21:15:27 +1000 +Subject: [PATCH 44/45] Fix GLM-DSA Metal dispatch safety + +--- + ggml/src/ggml-metal/ggml-metal-device.cpp | 13 +++--- + ggml/src/ggml-metal/ggml-metal-impl.h | 2 +- + ggml/src/ggml-metal/ggml-metal-ops.cpp | 49 +++++++++++------------ + ggml/src/ggml-metal/ggml-metal.metal | 6 +-- + 4 files changed, 33 insertions(+), 37 deletions(-) + +diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp +index 870d5a92..9ad67837 100644 +--- a/ggml/src/ggml-metal/ggml-metal-device.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-device.cpp +@@ -2779,15 +2779,14 @@ static ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id + + constexpr int nr0 = 8; + const int nsg = active_slots; ++ const char * base = ++ active_slots == 6 ? "kernel_mul_mv_id_q3_K_wr_slots_r8_a6" : ++ active_slots == 4 ? "kernel_mul_mv_id_q3_K_wr_slots_r8_a4" : ++ "kernel_mul_mv_id_q3_K_wr_slots_r8_a2"; + +- char base[256]; +- char name[256]; +- snprintf(base, 256, "kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_r8_nb8_w0"); +- snprintf(name, 256, "%s_nsg=%d_nr0=%d", base, nsg, nr0); +- +- ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); ++ ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, base); + if (!res.pipeline) { +- res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); ++ res = ggml_metal_library_compile_pipeline(lib, base, base, nullptr); + } + + res.nr0 = nr0; +diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h +index f7112ee5..f8431715 100644 +--- a/ggml/src/ggml-metal/ggml-metal-impl.h ++++ b/ggml/src/ggml-metal/ggml-metal-impl.h +@@ -575,7 +575,7 @@ typedef struct { + uint64_t weights_nb1; + uint64_t weights_nb2; + int32_t weighted; +- int32_t _pad0; ++ int32_t n_active_slots; + } ggml_metal_kargs_mul_mv_id_gate_up_swiglu; + + typedef struct { +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index 646c983a..f6a287b1 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -230,6 +230,11 @@ static int ggml_metal_glm_dsa_moe_max_active_experts() { + return requested >= 1 && requested <= 8 ? requested : 0; + } + ++static int ggml_metal_glm_dsa_moe_active_slots(int total_slots) { ++ const int requested = ggml_metal_glm_dsa_moe_max_active_experts(); ++ return requested > 0 && requested < total_slots ? requested : total_slots; ++} ++ + static int ggml_metal_glm_dsa_moe_q2_weight_roofline_chunks() { + const char * value = getenv("GGML_METAL_EXPERIMENTAL_GLM_MOE_Q2_WEIGHT_ROOFLINE_CHUNKS"); + if (value == nullptr) { +@@ -678,18 +683,11 @@ static int ggml_metal_glm_dsa_sparse_attn_threads_requested() { + } + } + +-static int ggml_metal_glm_dsa_sparse_attn_threads_for_shape(int requested, int n_batch, int n_top_k) { +- // Large-top-k prefill shapes are sensitive to the 256-thread sparse-attn +- // kernel on Apple Metal and can leave rows unwritten. Keep decode and +- // small-top-k prefill on the requested path, but cap large prefill rows to +- // the shape that passes backend parity. +- if (n_batch > 1 && n_top_k >= 64) { +- return std::min(requested, 32); +- } +- if (n_top_k > 512) { +- return std::min(requested, 128); +- } +- return requested; ++static int ggml_metal_glm_dsa_sparse_attn_threads(int requested) { ++ // The default sparse-attention kernel is parity-certified with one ++ // simdgroup. Wider threadgroups can leave output rows unwritten, so keep ++ // every shape on the safe dispatch until the kernel is made width-agnostic. ++ return std::min(requested, 32); + } + + static bool ggml_metal_glm_dsa_sparse_attn_cache_topk_enabled() { +@@ -3169,7 +3167,7 @@ static int ggml_metal_op_glm_moe_route_gate_up_swiglu( + /*.weights_nb1 =*/ nb_weights1, + /*.weights_nb2 =*/ nb_weights2, + /*.weighted =*/ motif.weighted_down_input != nullptr ? 1 : 0, +- /*._pad0 =*/ ggml_metal_glm_dsa_moe_max_active_experts(), ++ /*.n_active_slots =*/ ggml_metal_glm_dsa_moe_active_slots(ne20), + }; + + auto pipeline = ggml_metal_library_get_pipeline_glm_moe_route_q2_gate_up_swiglu_pair_sg_slot8(lib, dst); +@@ -3337,7 +3335,7 @@ static int ggml_metal_op_glm_moe_weights_gate_up_swiglu( + /*.weights_nb1 =*/ nb_route_weights1, + /*.weights_nb2 =*/ nb_route_weights2, + /*.weighted =*/ 0, +- /*._pad0 =*/ ggml_metal_glm_dsa_moe_max_active_experts(), ++ /*.n_active_slots =*/ ggml_metal_glm_dsa_moe_active_slots(ne20), + }; + + auto pipeline = ggml_metal_library_get_pipeline_glm_moe_weights_q2_gate_up_swiglu_pair_sg_slot1( +@@ -3623,7 +3621,7 @@ static int ggml_metal_op_glm_moe_two_phase( + /*.weights_nb1 =*/ motif.weights->nb[1], + /*.weights_nb2 =*/ motif.weights->nb[2], + /*.weighted =*/ 0, +- /*._pad0 =*/ 0, ++ /*.n_active_slots =*/ (int32_t) motif.ids->ne[0], + }; + ggml_metal_kargs_glm_moe_two_phase phase_args = { + /*.n_embd =*/ 6144, +@@ -4522,7 +4520,7 @@ static int ggml_metal_op_mul_mv_id_gate_up_swiglu( + /*.weights_nb1 =*/ weights_nb1, + /*.weights_nb2 =*/ weights_nb2, + /*.weighted =*/ weights != nullptr ? 1 : 0, +- /*._pad0 =*/ ggml_metal_glm_dsa_moe_max_active_experts(), ++ /*.n_active_slots =*/ ggml_metal_glm_dsa_moe_active_slots(ne20), + }; + + ggml_metal_encoder_set_pipeline(enc, pipeline); +@@ -4556,8 +4554,7 @@ static int ggml_metal_op_mul_mv_id_gate_up_swiglu( + slots_per_threadgroup = 4; + } + } +- const int active_slots = +- args._pad0 > 0 && args._pad0 < ne20 ? args._pad0 : ne20; ++ const int active_slots = args.n_active_slots; + const int64_t ne123 = pair_sg ? + ((active_slots + slots_per_threadgroup - 1)/slots_per_threadgroup)*ne21 : + active_slots*ne21; +@@ -5260,24 +5257,24 @@ static int ggml_metal_op_mul_mv_id_weighted_reduce( + } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active6_enabled() && + ne00 == 2048 && extra.already_weighted == 0) { + pipeline_slot_parallel = +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0(lib, weighted_sum); ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active6(lib, weighted_sum); + } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active4_enabled() && + ne00 == 2048 && extra.already_weighted == 0) { + pipeline_slot_parallel = +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0(lib, weighted_sum); ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active4(lib, weighted_sum); + } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active2_enabled() && + ne00 == 2048 && extra.already_weighted == 0) { + pipeline_slot_parallel = +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0(lib, weighted_sum); ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active2(lib, weighted_sum); + } else if (q3_active_count_shape && q3_active_count == 6) { + pipeline_slot_parallel = +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0(lib, weighted_sum); ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active6(lib, weighted_sum); + } else if (q3_active_count_shape && q3_active_count == 4) { + pipeline_slot_parallel = +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0(lib, weighted_sum); ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active4(lib, weighted_sum); + } else if (q3_active_count_shape && q3_active_count == 2) { + pipeline_slot_parallel = +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0(lib, weighted_sum); ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active2(lib, weighted_sum); + } else if (q3_f16_w0_shape) { + pipeline_slot_parallel = + ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_f16(lib, weighted_sum); +@@ -11666,8 +11663,8 @@ int ggml_metal_op_dsa_sparse_attn(ggml_metal_op_t ctx, int idx) { + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[4]), ida++); // top_k + ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), ida++); // dst + +- const int nth_requested = ggml_metal_glm_dsa_sparse_attn_threads_for_shape( +- ggml_metal_glm_dsa_sparse_attn_threads_requested(), ne1, ne40); ++ const int nth_requested = ++ ggml_metal_glm_dsa_sparse_attn_threads(ggml_metal_glm_dsa_sparse_attn_threads_requested()); + const int head_group = use_decode_grouped ? requested_head_group : 1; + const int max_threads_per_group = std::max(1, ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)/head_group); + const int nth = std::min(nth_requested, max_threads_per_group); +diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal +index b098231d..726c02ad 100644 +--- a/ggml/src/ggml-metal/ggml-metal.metal ++++ b/ggml/src/ggml-metal/ggml-metal.metal +@@ -10331,7 +10331,7 @@ void kernel_mul_mv_id_q2_K_gate_up_swiglu_pair_sg_impl( + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + const int nb = args.ne00/QK_K; + +- const int active_slots = args._pad0 > 0 && args._pad0 < args.nei0 ? args._pad0 : args.nei0; ++ const int active_slots = args.n_active_slots; + const int slot_groups = (active_slots + slots_per_tg - 1) / slots_per_tg; + const int r0 = tgpig.x; + const int token = tgpig.z / slot_groups; +@@ -10344,7 +10344,7 @@ void kernel_mul_mv_id_q2_K_gate_up_swiglu_pair_sg_impl( + + if (slot < args.nei0 && token < args.nei1) { + const int32_t expert_id = ((device const int32_t *) (ids + token*args.nbi1))[slot]; +- const bool active_slot = args._pad0 <= 0 || slot < args._pad0; ++ const bool active_slot = slot < args.n_active_slots; + if (active_slot && expert_id >= 0 && expert_id < args.ne02) { + device const char * src0_all = is_gate ? src0_gate_all : src0_up_all; + const uint64_t offset0 = uint64_t(expert_id)*args.nb02 + uint64_t(first_row)*args.nb01; +@@ -10782,7 +10782,7 @@ void kernel_mul_mv_id_q2_K_gate_up_swiglu_pair_sg_dual_impl( + + if (valid_output) { + const int32_t expert_id = ((device const int32_t *) (ids + token*args.nbi1))[slot]; +- const bool active_slot = (args._pad0 <= 0 || slot < args._pad0) && ++ const bool active_slot = slot < args.n_active_slots && + (args.weighted == 0 || route_weight != 0.0f); + if (active_slot && expert_id >= 0 && expert_id < args.ne02) { + const uint64_t offset0 = uint64_t(expert_id)*args.nb02 + uint64_t(first_row)*args.nb01; +-- +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/patches/0045-Scope-GLM-DSA-backend-test-environment.patch b/third_party/llama.cpp/patches/0045-Scope-GLM-DSA-backend-test-environment.patch new file mode 100644 index 0000000000..df2d2f2df9 --- /dev/null +++ b/third_party/llama.cpp/patches/0045-Scope-GLM-DSA-backend-test-environment.patch @@ -0,0 +1,189 @@ +From 99a0fe861d955acc921d0f328f0b1e8ba95b7699 Mon Sep 17 00:00:00 2001 +From: Mesh-LLM CI +Date: Sun, 26 Jul 2026 21:16:20 +1000 +Subject: [PATCH 45/45] Scope GLM-DSA backend test environment + +--- + tests/test-backend-ops.cpp | 87 +++++++++++++++++++++++++++++--------- + 1 file changed, 66 insertions(+), 21 deletions(-) + +diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp +index 43ce1a0e..7d47a4e0 100644 +--- a/tests/test-backend-ops.cpp ++++ b/tests/test-backend-ops.cpp +@@ -42,6 +42,7 @@ + #include + #include + #include ++#include + #include + #include + +@@ -11625,7 +11626,36 @@ private: + } + }; + +-struct test_glm_dsa_selected_row_tiled_parity : public test_case { ++struct test_case_with_scoped_env : public test_case { ++ using env_setting = std::pair; ++ ++ virtual std::vector env_settings() const = 0; ++ ++ test_status_t eval( ++ ggml_backend_t backend1, ++ ggml_backend_t backend2, ++ const char * op_names_filter, ++ printer * output_printer) override { ++ [[maybe_unused]] auto env = scoped_env(); ++ return test_case::eval(backend1, backend2, op_names_filter, output_printer); ++ } ++ ++ bool eval_perf(ggml_backend_t backend, const char * op_names_filter, printer * output_printer) override { ++ [[maybe_unused]] auto env = scoped_env(); ++ return test_case::eval_perf(backend, op_names_filter, output_printer); ++ } ++ ++private: ++ std::vector> scoped_env() const { ++ std::vector> env; ++ for (const auto & setting : env_settings()) { ++ env.emplace_back(std::make_unique(setting.first, setting.second)); ++ } ++ return env; ++ } ++}; ++ ++struct test_glm_dsa_selected_row_tiled_parity : public test_case_with_scoped_env { + static constexpr int64_t dk = 576; + static constexpr int64_t dv = 512; + static constexpr int64_t n_kv = 2304; +@@ -11642,11 +11672,15 @@ struct test_glm_dsa_selected_row_tiled_parity : public test_case { + return "GLM_DSA_SELECTED_ROW_TILED_PARITY"; + } + +- ggml_tensor * build_graph(ggml_context * ctx) override { +- set_test_env("SKIPPY_GLM_DSA_EXPERIMENTAL_SELECTED_ROW_FLASH", "1"); +- set_test_env("LLAMA_GLM_DSA_EXPERIMENTAL_SELECTED_ROW_FLASH_TILED", "1"); +- set_test_env("LLAMA_GLM_DSA_SELECTED_ROW_FLASH_TILED_TENSOR", "dsa_tiled"); ++ std::vector env_settings() const override { ++ return { ++ {"SKIPPY_GLM_DSA_EXPERIMENTAL_SELECTED_ROW_FLASH", "1"}, ++ {"LLAMA_GLM_DSA_EXPERIMENTAL_SELECTED_ROW_FLASH_TILED", "1"}, ++ {"LLAMA_GLM_DSA_SELECTED_ROW_FLASH_TILED_TENSOR", "dsa_tiled"}, ++ }; ++ } + ++ ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * q = ggml_new_tensor_4d(ctx, GGML_TYPE_F32, dk, n_batch, n_head, 1); + ggml_tensor * k = ggml_new_tensor_4d(ctx, GGML_TYPE_F16, dk, n_kv, 1, 1); + ggml_tensor * v = ggml_new_tensor_4d(ctx, GGML_TYPE_F16, dv, n_kv, 1, 1); +@@ -11869,7 +11903,7 @@ struct test_flash_attn_ext : public test_case { + } + }; + +-struct test_glm_dsa_selected_row_flash : public test_case { ++struct test_glm_dsa_selected_row_flash : public test_case_with_scoped_env { + const int64_t kv; + const int64_t top_k; + const int64_t heads; +@@ -11886,13 +11920,16 @@ struct test_glm_dsa_selected_row_flash : public test_case { + test_glm_dsa_selected_row_flash(int64_t kv = 257, int64_t top_k = 64, int64_t heads = 1) + : kv(kv), top_k(top_k), heads(heads) {} + +- ggml_tensor * build_graph(ggml_context * ctx) override { ++ std::vector env_settings() const override { + const char * disabled = getenv("SKIPPY_GLM_DSA_DISABLE_SELECTED_ROW_FLASH"); + const char * explicit_selected_row = getenv("SKIPPY_GLM_DSA_EXPERIMENTAL_SELECTED_ROW_FLASH"); + if (!(disabled && atoi(disabled) != 0) && explicit_selected_row == nullptr) { +- set_test_env("SKIPPY_GLM_DSA_EXPERIMENTAL_SELECTED_ROW_FLASH", "1"); ++ return {{"SKIPPY_GLM_DSA_EXPERIMENTAL_SELECTED_ROW_FLASH", "1"}}; + } ++ return {}; ++ } + ++ ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * q = ggml_new_tensor_4d(ctx, GGML_TYPE_F32, 576, 1, heads, 1); + ggml_tensor * packed_kv = ggml_new_tensor_4d(ctx, GGML_TYPE_F16, 576, kv, 1, 1); + ggml_tensor * top_k_ids = ggml_new_tensor_4d(ctx, GGML_TYPE_I32, top_k, 1, 1, 1); +@@ -12007,7 +12044,7 @@ struct test_glm_dsa_compact_multihead_flash : public test_case { + } + }; + +-struct test_glm_dsa_compact_multihead_parity : public test_case { ++struct test_glm_dsa_compact_multihead_parity : public test_case_with_scoped_env { + std::string vars() override { + return "kv=2048,heads=64,stock_vs_multihead"; + } +@@ -12017,12 +12054,14 @@ struct test_glm_dsa_compact_multihead_parity : public test_case { + return "GLM_DSA_COMPACT_MULTIHEAD_PARITY"; + } + +- ggml_tensor * build_graph(ggml_context * ctx) override { +- set_test_env("GGML_METAL_EXPERIMENTAL_GLM_COMPACT_MULTIHEAD_FLASH", "1"); +- set_test_env( +- "GGML_METAL_EXPERIMENTAL_GLM_COMPACT_MULTIHEAD_FLASH_TENSOR", +- "compact_multihead_candidate"); ++ std::vector env_settings() const override { ++ return { ++ {"GGML_METAL_EXPERIMENTAL_GLM_COMPACT_MULTIHEAD_FLASH", "1"}, ++ {"GGML_METAL_EXPERIMENTAL_GLM_COMPACT_MULTIHEAD_FLASH_TENSOR", "compact_multihead_candidate"}, ++ }; ++ } + ++ ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * q = ggml_new_tensor_4d(ctx, GGML_TYPE_F32, 576, 1, 64, 1); + ggml_tensor * k = ggml_new_tensor_4d(ctx, GGML_TYPE_F16, 576, 2048, 1, 1); + ggml_tensor * v = ggml_new_tensor_4d(ctx, GGML_TYPE_F16, 512, 2048, 1, 1); +@@ -12096,7 +12135,7 @@ struct test_glm_dsa_compact_multihead_parity : public test_case { + } + }; + +-struct test_glm_dsa_compact_nwg_parity : public test_case { ++struct test_glm_dsa_compact_nwg_parity : public test_case_with_scoped_env { + std::string vars() override { + return "kv=2048,heads=64,nwg=4_vs_8"; + } +@@ -12106,9 +12145,11 @@ struct test_glm_dsa_compact_nwg_parity : public test_case { + return "GLM_DSA_COMPACT_NWG_PARITY"; + } + +- ggml_tensor * build_graph(ggml_context * ctx) override { +- set_test_env("GGML_METAL_EXPERIMENTAL_GLM_COMPACT_NWG8_TENSOR", "compact_nwg8_candidate"); ++ std::vector env_settings() const override { ++ return {{"GGML_METAL_EXPERIMENTAL_GLM_COMPACT_NWG8_TENSOR", "compact_nwg8_candidate"}}; ++ } + ++ ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * q = ggml_new_tensor_4d(ctx, GGML_TYPE_F32, 576, 1, 64, 1); + ggml_tensor * k = ggml_new_tensor_4d(ctx, GGML_TYPE_F16, 576, 2048, 1, 1); + ggml_tensor * v = ggml_view_4d( +@@ -12172,7 +12213,7 @@ struct test_glm_dsa_compact_nwg_parity : public test_case { + } + }; + +-struct test_glm_dsa_compact_split_parity : public test_case { ++struct test_glm_dsa_compact_split_parity : public test_case_with_scoped_env { + std::string vars() override { + return "kv=2048,heads=64,split=control_vs_exact"; + } +@@ -12182,10 +12223,14 @@ struct test_glm_dsa_compact_split_parity : public test_case { + return "GLM_DSA_COMPACT_SPLIT_PARITY"; + } + +- ggml_tensor * build_graph(ggml_context * ctx) override { +- set_test_env("GGML_METAL_EXPERIMENTAL_GLM_COMPACT_SPLIT_EXACT", "1"); +- set_test_env("GGML_METAL_EXPERIMENTAL_GLM_COMPACT_SPLIT_EXACT_TENSOR", "compact_split_candidate"); ++ std::vector env_settings() const override { ++ return { ++ {"GGML_METAL_EXPERIMENTAL_GLM_COMPACT_SPLIT_EXACT", "1"}, ++ {"GGML_METAL_EXPERIMENTAL_GLM_COMPACT_SPLIT_EXACT_TENSOR", "compact_split_candidate"}, ++ }; ++ } + ++ ggml_tensor * build_graph(ggml_context * ctx) override { + ggml_tensor * q = ggml_new_tensor_4d(ctx, GGML_TYPE_F32, 576, 1, 64, 1); + ggml_tensor * k = ggml_new_tensor_4d(ctx, GGML_TYPE_F16, 576, 2048, 1, 1); + ggml_tensor * v = ggml_view_4d( +-- +2.50.1 (Apple Git-155) + diff --git a/third_party/llama.cpp/upstream.txt b/third_party/llama.cpp/upstream.txt index d3a005b7f5..87b944fcc7 100644 --- a/third_party/llama.cpp/upstream.txt +++ b/third_party/llama.cpp/upstream.txt @@ -1 +1 @@ -76f46ad29d61fd8c1401e8221842934bf62a6064 +ff067f76dd8e9e05f0528056f1274adf01a54d70