diff --git a/third_party/llama.cpp/patches/0030-ggml-default-GLM-MoE-two-phase-Metal-path.patch b/third_party/llama.cpp/patches/0030-ggml-default-GLM-MoE-two-phase-Metal-path.patch new file mode 100644 index 0000000000..e5d3caa626 --- /dev/null +++ b/third_party/llama.cpp/patches/0030-ggml-default-GLM-MoE-two-phase-Metal-path.patch @@ -0,0 +1,63 @@ +From 11fb71bbe0c6a1b525097a455d75810f0a4e43e4 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 + +--- + ggml/src/ggml-metal/ggml-metal-ops.cpp | 9 ++++++++- + tests/test-backend-ops.cpp | 12 ++++++++++++ + 2 files changed, 20 insertions(+), 1 deletion(-) + +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index 36378f8e..684efdc6 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -127,7 +127,14 @@ static bool ggml_metal_glm_dsa_moe_private_scratch_enabled() { + + static bool ggml_metal_glm_dsa_moe_two_phase_enabled() { + const char * value = getenv("GGML_METAL_EXPERIMENTAL_GLM_MOE_TWO_PHASE"); +- return value && atoi(value) != 0; ++ if (value) { ++ return atoi(value) != 0; ++ } ++ value = getenv("GGML_METAL_DISABLE_GLM_MOE_TWO_PHASE"); ++ if (value) { ++ return atoi(value) == 0; ++ } ++ return true; + } + + 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 +--- 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 { + return std::string("GLM_FFN_TWO_PHASE_") + variant; + } + ++ test_status_t eval( ++ ggml_backend_t backend1, ++ ggml_backend_t backend2, ++ 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( ++ "GGML_METAL_EXPERIMENTAL_GLM_MOE_TWO_PHASE", two_phase ? "1" : "0"); ++ 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 { + 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() { + + 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")); + test_cases.emplace_back(new test_glm_ffn_dual_lane_consumer(true, "CANDIDATE")); + + return test_cases; +-- +2.54.0 (Apple Git-157) + diff --git a/third_party/llama.cpp/patches/0031-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch b/third_party/llama.cpp/patches/0031-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch new file mode 100644 index 0000000000..cd3effc48c --- /dev/null +++ b/third_party/llama.cpp/patches/0031-ggml-add-GLM-MoE-Metal-selector-diagnostics.patch @@ -0,0 +1,221 @@ +From 8797805107cd1e49bc018a498ce9d008493000f8 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 + +--- + ggml/src/ggml-metal/ggml-metal-ops.cpp | 75 +++++++++++++++++++++----- + tests/test-backend-ops.cpp | 39 ++++++++++++-- + 2 files changed, 98 insertions(+), 16 deletions(-) + +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index 684efdc6..ab41e1ec 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -147,6 +147,11 @@ static bool ggml_metal_glm_dsa_moe_dual_lane_gate_slot4_enabled() { + return value && atoi(value) != 0; + } + ++static bool ggml_metal_glm_dsa_moe_debug_terminal_outputs_enabled() { ++ const char * value = getenv("GGML_METAL_DEBUG_GLM_MOE_TERMINAL_OUTPUTS"); ++ return value && atoi(value) != 0; ++} ++ + static bool ggml_metal_glm_dsa_absorbed_qkv_phases_enabled() { + const char * value = getenv("GGML_METAL_EXPERIMENTAL_GLM_ABSORBED_QKV_PHASES"); + return value && atoi(value) != 0; +@@ -3731,6 +3736,17 @@ static int ggml_metal_op_glm_moe_two_phase( + ctx->enc, (phase_args.n_out + 255)/256, 1, 1, 256, 1, 1); + ctx->set_fused_range_outputs(motif.final_out_offset); + } else { ++ const bool debug_terminal_outputs = ++ ggml_metal_glm_dsa_moe_debug_terminal_outputs_enabled(); ++ if (debug_terminal_outputs) { ++ ggml_metal_encode_glm_moe_shared_mul_mv( ++ ctx, ++ motif.shared_down, ++ ggml_metal_get_buffer_id(motif.shared_glu), ++ ggml_metal_get_buffer_id(motif.shared_down)); ++ ggml_metal_op_concurrency_reset(ctx); ++ } ++ + const char * down_kernel = "kernel_glm_moe_two_phase_down"; + auto down_pipeline = ggml_metal_library_get_pipeline(ctx->lib, down_kernel); + if (!down_pipeline.pipeline) { +@@ -3753,7 +3769,12 @@ static int ggml_metal_op_glm_moe_two_phase( + ggml_metal_encoder_set_threadgroup_memory_size(ctx->enc, 8*8*sizeof(float), 0); + ggml_metal_encoder_dispatch_threadgroups( + ctx->enc, (phase_args.n_out + 7)/8, 1, 1, 32, 8, 1); +- ctx->set_fused_range_outputs(motif.out_offset, motif.final_out_offset); ++ if (debug_terminal_outputs) { ++ ctx->set_fused_range_outputs( ++ motif.out_offset, motif.shared_down_offset, motif.final_out_offset); ++ } else { ++ ctx->set_fused_range_outputs(motif.out_offset, motif.final_out_offset); ++ } + } + ggml_metal_op_concurrency_reset(ctx); + +@@ -4288,25 +4309,52 @@ static int ggml_metal_op_mul_mv_id_gate_up_swiglu( + if (dst->type == GGML_TYPE_F16 && !pair_sg) { + return 0; + } +- const bool pair_sg_slot8 = pair_sg && dst->type == GGML_TYPE_F32 && ++ const bool pair_sg_slot8_requested = + ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot8_enabled(); +- const bool pair_sg_slot2 = pair_sg && dst->type == GGML_TYPE_F32 && ++ const bool pair_sg_slot2_requested = + ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot2_enabled(); +- const bool pair_sg_slot4_dual_default = pair_sg_default_shape && src1_override != nullptr; ++ const bool pair_sg_slot4_dual_requested = ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot4_dual_enabled(); ++ const bool pair_sg_slot4_dual_r12_requested = ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot4_dual_r12_enabled(); ++ const bool pair_sg_slot4_dual_r16_requested = ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot4_dual_r16_enabled(); ++ const bool pair_sg_slot2_dual_requested = ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot2_dual_enabled(); ++ const bool pair_sg_slot8_split_requested = ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot8_split_enabled(); ++ const bool pair_sg_rowtile_requested = ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_rowtile_enabled(); ++ const bool pair_sg_slot1_dual_default = ++ pair_sg_default_shape && ++ !pair_sg_slot8_requested && ++ !pair_sg_slot2_requested && ++ !pair_sg_slot4_dual_requested && ++ !pair_sg_slot4_dual_r12_requested && ++ !pair_sg_slot4_dual_r16_requested && ++ !pair_sg_slot2_dual_requested && ++ !pair_sg_slot8_split_requested && ++ !pair_sg_rowtile_requested; ++ const bool pair_sg_slot8 = pair_sg && dst->type == GGML_TYPE_F32 && ++ pair_sg_slot8_requested; ++ const bool pair_sg_slot2 = pair_sg && dst->type == GGML_TYPE_F32 && ++ pair_sg_slot2_requested; ++ const bool pair_sg_slot4_dual_default = pair_sg_default_shape && ++ src1_override != nullptr && ++ pair_sg_slot1_dual_default; + const bool pair_sg_slot4_dual = pair_sg && dst->type == GGML_TYPE_F32 && +- (ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot4_dual_enabled() || +- pair_sg_slot4_dual_default); ++ (pair_sg_slot4_dual_requested || pair_sg_slot4_dual_default); + const bool pair_sg_slot4_dual_r12 = pair_sg && dst->type == GGML_TYPE_F32 && +- ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot4_dual_r12_enabled(); ++ pair_sg_slot4_dual_r12_requested; + const bool pair_sg_slot4_dual_r16 = pair_sg && dst->type == GGML_TYPE_F32 && +- ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot4_dual_r16_enabled(); ++ pair_sg_slot4_dual_r16_requested; + const bool pair_sg_slot1_dual = pair_sg && dst->type == GGML_TYPE_F32 && + (ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot1_dual_enabled() || +- (pair_sg_default_shape && !pair_sg_slot4_dual_default)); ++ (pair_sg_slot1_dual_default && !pair_sg_slot4_dual_default)); + const bool pair_sg_slot2_dual = pair_sg && dst->type == GGML_TYPE_F32 && +- ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot2_dual_enabled(); ++ pair_sg_slot2_dual_requested; + const bool pair_sg_slot8_split = pair_sg && dst->type == GGML_TYPE_F32 && +- ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot8_split_enabled(); ++ pair_sg_slot8_split_requested; + const bool pair_sg_share_y = pair_sg && dst->type == GGML_TYPE_F32 && + ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_share_y_enabled() && + ne00 == 6144 && +@@ -4351,7 +4399,7 @@ static int ggml_metal_op_mul_mv_id_gate_up_swiglu( + ne21 == 1 && + src1->type == GGML_TYPE_F32; + const bool pair_sg_rowtile = pair_sg && dst->type == GGML_TYPE_F32 && +- ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_rowtile_enabled() && ++ pair_sg_rowtile_requested && + ne00 == 6144 && + ne01 == 2048 && + ne20 == 8 && +@@ -4893,7 +4941,8 @@ static int ggml_metal_op_mul_mv_id_weighted_reduce( + enc, ggml_metal_get_buffer_id(down->src[2])); + + const bool use_output_scratch = +- down->src[0]->type == GGML_TYPE_Q3_K && !preserve_routed_input; ++ down->src[0]->type == GGML_TYPE_Q3_K && ++ !preserve_routed_input; + const ggml_metal_buffer_id weighted_reduce_dst = use_output_scratch ? + 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 +--- a/tests/test-backend-ops.cpp ++++ b/tests/test-backend-ops.cpp +@@ -9575,6 +9575,7 @@ struct test_glm_ffn_consumer : public test_case { + ggml_tensor * selected_experts = nullptr; + ggml_tensor * weights = nullptr; + ggml_tensor * routed_out = nullptr; ++ ggml_tensor * shexp_out = nullptr; + ggml_tensor * out = nullptr; + + test_glm_ffn_consumer( +@@ -9762,9 +9763,9 @@ struct test_glm_ffn_consumer : public test_case { + } + } + +- ggml_tensor * shexp = ggml_mul_mat(ctx, downs, act); +- ggml_set_name(shexp, "ffn_shexp"); +- return shexp; ++ shexp_out = ggml_mul_mat(ctx, downs, act); ++ ggml_set_name(shexp_out, "ffn_shexp"); ++ return shexp_out; + } + + ggml_tensor * build_graph(ggml_context * ctx) override { +@@ -9851,6 +9852,36 @@ struct test_glm_ffn_two_phase_consumer : public test_glm_ffn_consumer { + } + }; + ++struct test_glm_ffn_terminal_diagnostic : public test_glm_ffn_two_phase_consumer { ++ explicit test_glm_ffn_terminal_diagnostic(bool two_phase, const char * variant) ++ : test_glm_ffn_two_phase_consumer(two_phase, variant) {} ++ ++ std::string op_desc(ggml_tensor * t) override { ++ GGML_UNUSED(t); ++ return std::string("GLM_FFN_TERMINAL_DIAGNOSTIC_") + variant; ++ } ++ ++ std::vector fusion_test_nodes() override { ++ return { selected_experts, weights, routed_out, shexp_out, out }; ++ } ++ ++ test_status_t eval( ++ ggml_backend_t backend1, ++ ggml_backend_t backend2, ++ const char * op_names_filter, ++ printer * output_printer) override { ++ scoped_test_env terminal_outputs_env("GGML_METAL_DEBUG_GLM_MOE_TERMINAL_OUTPUTS", "1"); ++ return test_glm_ffn_two_phase_consumer::eval( ++ backend1, backend2, op_names_filter, output_printer); ++ } ++ ++ bool eval_perf(ggml_backend_t backend, const char * op_names_filter, printer * output_printer) override { ++ scoped_test_env terminal_outputs_env("GGML_METAL_DEBUG_GLM_MOE_TERMINAL_OUTPUTS", "1"); ++ return test_glm_ffn_two_phase_consumer::eval_perf( ++ backend, op_names_filter, output_printer); ++ } ++}; ++ + 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() { + 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")); ++ test_cases.emplace_back(new test_glm_ffn_terminal_diagnostic(true, "CANDIDATE")); + 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() { + 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)); ++ test_cases.emplace_back(new test_glm_ffn_terminal_diagnostic(true, "CANDIDATE")); + test_cases.emplace_back(new test_glm_ffn_consumer(GGML_TYPE_TQ2_0, GGML_TYPE_TQ2_0, GGML_TYPE_TQ2_0, + 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) + 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/0032-ggml-skip-zero-weight-GLM-MoE-gate-up-slots.patch new file mode 100644 index 0000000000..2ed38ec876 --- /dev/null +++ b/third_party/llama.cpp/patches/0032-ggml-skip-zero-weight-GLM-MoE-gate-up-slots.patch @@ -0,0 +1,45 @@ +From 4523e19eecd43c47ca650939bb0c2dfc87c43920 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 + +--- + 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 +--- 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( + return; + } + +- if (slot < args.nei0 && token < args.nei1) { ++ const bool valid_output = slot < args.nei0 && token < args.nei1; ++ const float route_weight = valid_output && args.weighted != 0 ? ++ ((device const float *) (weights + uint64_t(slot)*args.weights_nb1 + uint64_t(token)*args.weights_nb2))[0] : ++ 1.0f; ++ ++ 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 = (args._pad0 <= 0 || slot < args._pad0) && ++ (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; + const uint64_t offset1 = uint64_t(slot % args.ne11)*args.nb11 + uint64_t(token)*args.nb12; +@@ -10936,10 +10942,6 @@ void kernel_mul_mv_id_q2_K_gate_up_swiglu_pair_sg_dual_impl( + } + } + +- const bool valid_output = slot < args.nei0 && token < args.nei1; +- const float route_weight = valid_output && args.weighted != 0 ? +- ((device const float *) (weights + uint64_t(slot)*args.weights_nb1 + uint64_t(token)*args.weights_nb2))[0] : +- 1.0f; + for (int row = 0; row < nr0; ++row) { + const float up_value = simd_sum(sum_up[row]); + const float gate_value = simd_sum(sum_gate[row]); +-- +2.54.0 (Apple Git-157) + diff --git a/third_party/llama.cpp/patches/0033-tests-cover-full-GLM-MoE-selected-chain.patch b/third_party/llama.cpp/patches/0033-tests-cover-full-GLM-MoE-selected-chain.patch new file mode 100644 index 0000000000..15808800a4 --- /dev/null +++ b/third_party/llama.cpp/patches/0033-tests-cover-full-GLM-MoE-selected-chain.patch @@ -0,0 +1,35 @@ +From 37ee00411144af6570074a352ac8584ab8527ad6 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 + +--- + 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 +--- 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 { + } + } + +- double max_nmse_err() override { return 5e-4; } ++ double max_nmse_err() override { ++ return n_repeats > 1 ? 5e-3 : 5e-4; ++ } + + 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() { + 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")); ++ 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}) { + test_cases.emplace_back(new test_glm_attention_projection_chain(type, 2)); + } +-- +2.54.0 (Apple Git-157) + diff --git a/third_party/llama.cpp/patches/0034-tests-add-GLM-Q2Q3-selected-weight-roofline.patch b/third_party/llama.cpp/patches/0034-tests-add-GLM-Q2Q3-selected-weight-roofline.patch new file mode 100644 index 0000000000..71db34b48a --- /dev/null +++ b/third_party/llama.cpp/patches/0034-tests-add-GLM-Q2Q3-selected-weight-roofline.patch @@ -0,0 +1,67 @@ +From e9f260788b857c0f86f41229b3bb2da52f3d2187 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 + +--- + 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 +--- 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 + } + }; + ++struct test_glm_moe_q2q3_weight_roofline : public test_glm_moe_decode_post_weight_component { ++ const bool bypass_route; ++ const char * variant; ++ ++ test_glm_moe_q2q3_weight_roofline(bool bypass_route, const char * variant) ++ : test_glm_moe_decode_post_weight_component(6), ++ bypass_route(bypass_route), ++ variant(variant) {} ++ ++ std::string op_desc(ggml_tensor * t) override { ++ GGML_UNUSED(t); ++ return std::string("GLM_MOE_Q2Q3_WEIGHT_ROOFLINE_") + variant; ++ } ++ ++ 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"); ++ scoped_test_env chunks_env("GGML_METAL_EXPERIMENTAL_GLM_MOE_Q2_WEIGHT_ROOFLINE_CHUNKS", "16"); ++ scoped_test_env bypass_env( ++ "GGML_METAL_EXPERIMENTAL_GLM_MOE_Q2_WEIGHT_ROOFLINE_BYPASS_ROUTE", ++ bypass_route ? "1" : "0"); ++ return test_case::eval_perf(backend, op_names_filter, output_printer); ++ } ++}; ++ + 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() { + 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")); ++ test_cases.emplace_back(new test_glm_moe_q2q3_weight_roofline(false, "ROUTE_A")); ++ test_cases.emplace_back(new test_glm_moe_q2q3_weight_roofline(true, "FIXED_IDS")); ++ test_cases.emplace_back(new test_glm_moe_q2q3_weight_roofline(false, "ROUTE_B")); + 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() { + 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")); ++ test_cases.emplace_back(new test_glm_moe_q2q3_weight_roofline(false, "ROUTE_A")); ++ test_cases.emplace_back(new test_glm_moe_q2q3_weight_roofline(true, "FIXED_IDS")); ++ test_cases.emplace_back(new test_glm_moe_q2q3_weight_roofline(false, "ROUTE_B")); + test_cases.emplace_back(new test_glm_moe_selected_chain(3, false, "USEFUL_A")); + 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) + 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/0035-ggml-add-active-count-Q3-GLM-MoE-down-kernels.patch new file mode 100644 index 0000000000..c908a18f9c --- /dev/null +++ b/third_party/llama.cpp/patches/0035-ggml-add-active-count-Q3-GLM-MoE-down-kernels.patch @@ -0,0 +1,432 @@ +From 66c8184a2a33a5fb60bfcadc47d4f8088669a5fe 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 + +--- + ggml/src/ggml-metal/ggml-metal-device.cpp | 49 +++++++++++ + ggml/src/ggml-metal/ggml-metal-device.h | 3 + + ggml/src/ggml-metal/ggml-metal-ops.cpp | 102 +++++++++++++++++----- + ggml/src/ggml-metal/ggml-metal.metal | 54 +++++++++++- + tests/test-backend-ops.cpp | 44 +++++++++- + 5 files changed, 222 insertions(+), 30 deletions(-) + +diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp +index 088089aa..870d5a92 100644 +--- a/ggml/src/ggml-metal/ggml-metal-device.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-device.cpp +@@ -2761,6 +2761,55 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_wei + return res; + } + ++static ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active( ++ ggml_metal_library_t lib, ++ const ggml_tensor * op, ++ int active_slots) { ++ GGML_ASSERT(op->op == GGML_OP_MOE_WEIGHTED_SUM); ++ GGML_ASSERT(op->src[0] != nullptr); ++ GGML_ASSERT(op->src[0]->op == GGML_OP_MUL_MAT_ID); ++ GGML_ASSERT(op->src[0]->src[0]->type == GGML_TYPE_Q3_K); ++ GGML_ASSERT(op->src[0]->src[0]->ne[0] == 2048); ++ GGML_ASSERT(op->src[0]->src[1]->type == GGML_TYPE_F32); ++ GGML_ASSERT(op->src[0]->src[2]->ne[0] == 8); ++ GGML_ASSERT(op->src[1]->type == GGML_TYPE_F32); ++ GGML_ASSERT(op->type == GGML_TYPE_F32); ++ GGML_ASSERT(ggml_get_op_params_i32(op, 0) == 0); ++ GGML_ASSERT(active_slots == 2 || active_slots == 4 || active_slots == 6); ++ ++ constexpr int nr0 = 8; ++ const int nsg = active_slots; ++ ++ 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); ++ if (!res.pipeline) { ++ res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr); ++ } ++ ++ res.nr0 = nr0; ++ res.nr1 = 1; ++ res.nsg = nsg; ++ res.smem = nsg*nr0*sizeof(float); ++ ++ return res; ++} ++ ++ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active6(ggml_metal_library_t lib, const ggml_tensor * op) { ++ return ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active(lib, op, 6); ++} ++ ++ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active4(ggml_metal_library_t lib, const ggml_tensor * op) { ++ return ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active(lib, op, 4); ++} ++ ++ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active2(ggml_metal_library_t lib, const ggml_tensor * op) { ++ return ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active(lib, op, 2); ++} ++ + ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_f16(ggml_metal_library_t lib, const ggml_tensor * op) { + assert(op->op == GGML_OP_MOE_WEIGHTED_SUM); + GGML_ASSERT(op->src[0] != nullptr); +diff --git a/ggml/src/ggml-metal/ggml-metal-device.h b/ggml/src/ggml-metal/ggml-metal-device.h +index 03567502..b8a48d13 100644 +--- a/ggml/src/ggml-metal/ggml-metal-device.h ++++ b/ggml/src/ggml-metal/ggml-metal-device.h +@@ -190,6 +190,9 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8(ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8(ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0(ggml_metal_library_t lib, const struct ggml_tensor * op); ++struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active6(ggml_metal_library_t lib, const struct ggml_tensor * op); ++struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active4(ggml_metal_library_t lib, const struct ggml_tensor * op); ++struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_active2(ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_f16(ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r6_nb8_w0(ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r10_nb8_w0(ggml_metal_library_t lib, const struct ggml_tensor * op); +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index ab41e1ec..b4fc441a 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -507,6 +507,21 @@ static bool ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_w0_enabled() + return value && atoi(value) != 0; + } + ++static bool ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active6_enabled() { ++ const char * value = getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_ACTIVE6"); ++ return value && atoi(value) != 0; ++} ++ ++static bool ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active4_enabled() { ++ const char * value = getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_ACTIVE4"); ++ return value && atoi(value) != 0; ++} ++ ++static bool ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active2_enabled() { ++ const char * value = getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_ACTIVE2"); ++ return value && atoi(value) != 0; ++} ++ + static bool ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r6_nb8_w0_enabled() { + const char * value = getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_R6_NB8_W0"); + return value && atoi(value) != 0; +@@ -4773,6 +4788,9 @@ static bool ggml_metal_match_mul_mv_id_weighted_reduce( + !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_enabled() && + !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_enabled() && + !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_w0_enabled() && ++ !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active6_enabled() && ++ !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active4_enabled() && ++ !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active2_enabled() && + !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r6_nb8_w0_enabled() && + !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r10_nb8_w0_enabled() && + !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_glm52_w0_enabled() && +@@ -5162,6 +5180,9 @@ static int ggml_metal_op_mul_mv_id_weighted_reduce( + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_w0_enabled() || ++ ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active6_enabled() || ++ ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active4_enabled() || ++ ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active2_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r6_nb8_w0_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r10_nb8_w0_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_glm52_w0_enabled() || +@@ -5171,33 +5192,68 @@ static int ggml_metal_op_mul_mv_id_weighted_reduce( + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r12_nb8_w0_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r16_enabled() || + ggml_metal_glm_dsa_q3_down_slot_split2_reduce_enabled())) { +- auto pipeline_slot_parallel = ggml_metal_glm_dsa_q3_down_slot_split2_reduce_enabled() ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_split2(lib, weighted_sum) : +- (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r16_enabled() ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r16(lib, weighted_sum) : +- (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_glm52_w0_enabled() && ne00 == 2048 && ne01 == 6144 && ne20 == 8 && ne21 == 1 && extra.already_weighted == 0 ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_glm52_w0(lib, weighted_sum) : +- (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r10_nb8_w0_enabled() && ne00 == 2048 && extra.already_weighted == 0 ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r10_nb8_w0(lib, weighted_sum) : +- (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r12_nb8_w0_enabled() && ne00 == 2048 && extra.already_weighted == 0 ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r12_nb8_w0(lib, weighted_sum) : +- (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r6_nb8_w0_enabled() && ne00 == 2048 && extra.already_weighted == 0 ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r6_nb8_w0(lib, weighted_sum) : +- (q3_f16_w0_shape ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0_f16(lib, weighted_sum) : +- ((ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_w0_enabled() || q3_default_w0_shape) && ne00 == 2048 && extra.already_weighted == 0 ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8_w0(lib, weighted_sum) : +- (q3_default_w1_shape && ne00 == 2048 && extra.already_weighted != 0 ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8(lib, weighted_sum) : +- (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_enabled() && ne00 == 2048 ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8(lib, weighted_sum) : +- (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_enabled() ? +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8(lib, weighted_sum) : +- ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg(lib, weighted_sum))))))))))); ++ ggml_metal_pipeline_with_params pipeline_slot_parallel; ++ if (ggml_metal_glm_dsa_q3_down_slot_split2_reduce_enabled()) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_split2(lib, weighted_sum); ++ } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r16_enabled()) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r16(lib, weighted_sum); ++ } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_glm52_w0_enabled() && ++ ne00 == 2048 && ne01 == 6144 && ne20 == 8 && ne21 == 1 && extra.already_weighted == 0) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_glm52_w0(lib, weighted_sum); ++ } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r10_nb8_w0_enabled() && ++ ne00 == 2048 && extra.already_weighted == 0) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r10_nb8_w0(lib, weighted_sum); ++ } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r12_nb8_w0_enabled() && ++ ne00 == 2048 && extra.already_weighted == 0) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r12_nb8_w0(lib, weighted_sum); ++ } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r6_nb8_w0_enabled() && ++ ne00 == 2048 && extra.already_weighted == 0) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r6_nb8_w0(lib, weighted_sum); ++ } 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); ++ } 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); ++ } 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); ++ } 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); ++ } else if ((ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_w0_enabled() || q3_default_w0_shape) && ++ 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); ++ } else if (q3_default_w1_shape && ne00 == 2048 && extra.already_weighted != 0) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8(lib, weighted_sum); ++ } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_nb8_enabled() && ne00 == 2048) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8_nb8(lib, weighted_sum); ++ } else if (ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_enabled()) { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg_r8(lib, weighted_sum); ++ } else { ++ pipeline_slot_parallel = ++ ggml_metal_library_get_pipeline_mul_mv_id_q3_weighted_reduce_slots_sg(lib, weighted_sum); ++ } + const int nr0_slot_parallel = pipeline_slot_parallel.nr0; + const int nsg_slot_parallel = pipeline_slot_parallel.nsg; + const int grid_x_slot_parallel = (ne01 + nr0_slot_parallel - 1)/nr0_slot_parallel; + const int grid_y_slot_parallel = ne21; ++ if (nsg_slot_parallel < 8) { ++ extra._pad0 = nsg_slot_parallel; ++ } + + 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 +--- 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( + } + } + +-template ++template + void kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_impl( + constant ggml_metal_kargs_mul_mv_id & args, + device const char * src0s, +@@ -11767,7 +11767,7 @@ void kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_impl( + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { +- const int n_slots = 8; ++ const int n_slots = extra._pad0 > 0 ? extra._pad0 : ACTIVE_SLOTS; + const int nb = NB_FIXED == 0 ? args.ne00/QK_K : NB_FIXED; + + const int r0 = tgpig.x; +@@ -11983,6 +11983,54 @@ kernel void kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_r8_nb8_w0( + args, src0s, src1, dst, ids, weights, extra, partials, tgpig, tiisg, sgitg); + } + ++kernel void kernel_mul_mv_id_q3_K_wr_slots_r8_a6( ++ constant ggml_metal_kargs_mul_mv_id & args, ++ device const char * src0s, ++ device const char * src1, ++ device char * dst, ++ device const char * ids, ++ device const char * weights, ++ constant ggml_metal_kargs_mul_mv_id_weighted_reduce_extra & extra, ++ threadgroup float * partials [[threadgroup(0)]], ++ uint3 tgpig[[threadgroup_position_in_grid]], ++ ushort tiisg[[thread_index_in_simdgroup]], ++ ushort sgitg[[simdgroup_index_in_threadgroup]]) { ++ kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_impl<8, 8, 0, 6>( ++ args, src0s, src1, dst, ids, weights, extra, partials, tgpig, tiisg, sgitg); ++} ++ ++kernel void kernel_mul_mv_id_q3_K_wr_slots_r8_a4( ++ constant ggml_metal_kargs_mul_mv_id & args, ++ device const char * src0s, ++ device const char * src1, ++ device char * dst, ++ device const char * ids, ++ device const char * weights, ++ constant ggml_metal_kargs_mul_mv_id_weighted_reduce_extra & extra, ++ threadgroup float * partials [[threadgroup(0)]], ++ uint3 tgpig[[threadgroup_position_in_grid]], ++ ushort tiisg[[thread_index_in_simdgroup]], ++ ushort sgitg[[simdgroup_index_in_threadgroup]]) { ++ kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_impl<8, 8, 0, 4>( ++ args, src0s, src1, dst, ids, weights, extra, partials, tgpig, tiisg, sgitg); ++} ++ ++kernel void kernel_mul_mv_id_q3_K_wr_slots_r8_a2( ++ constant ggml_metal_kargs_mul_mv_id & args, ++ device const char * src0s, ++ device const char * src1, ++ device char * dst, ++ device const char * ids, ++ device const char * weights, ++ constant ggml_metal_kargs_mul_mv_id_weighted_reduce_extra & extra, ++ threadgroup float * partials [[threadgroup(0)]], ++ uint3 tgpig[[threadgroup_position_in_grid]], ++ ushort tiisg[[thread_index_in_simdgroup]], ++ ushort sgitg[[simdgroup_index_in_threadgroup]]) { ++ kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_impl<8, 8, 0, 2>( ++ args, src0s, src1, dst, ids, weights, extra, partials, tgpig, tiisg, sgitg); ++} ++ + kernel void kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_r8_nb8_w0_f16( + constant ggml_metal_kargs_mul_mv_id & args, + device const char * src0s, +@@ -11995,7 +12043,7 @@ kernel void kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_r8_nb8_w0_f16( + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { +- kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_impl<8, 8, 0, half>( ++ kernel_mul_mv_id_q3_K_weighted_reduce_slots_sg_impl<8, 8, 0, 8, half>( + args, src0s, src1, dst, ids, weights, extra, partials, tgpig, tiisg, sgitg); + } + +diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp +index c80ddaf9..3a2448a8 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 { + } + return weighted_down ? "GLM_MOE_STATIC_MOTIF_Q2_WEIGHTED" : "GLM_MOE_STATIC_MOTIF_Q2"; + } ++ if (gate_type == GGML_TYPE_Q2_K && up_type == GGML_TYPE_Q2_K && down_type == GGML_TYPE_Q3_K && ++ n_embd == 6144 && n_ff == 2048) { ++ if (active_experts == 6 && packed_selected_experts && !weighted_down) { ++ return "GLM_MOE_STATIC_MOTIF_Q2Q3_ACTIVE6_POSTWEIGHT_GLM"; ++ } ++ if (active_experts == 8 && !weighted_down) { ++ return "GLM_MOE_STATIC_MOTIF_Q2Q3_POSTWEIGHT_GLM"; ++ } ++ } + 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 { + const int64_t n_ff; + const bool already_weighted; + const bool f16_act; ++ const int64_t active_experts; + + ggml_tensor * selected_experts = nullptr; + ggml_tensor * weights = nullptr; +@@ -9097,23 +9107,38 @@ struct test_glm_moe_down_reduce : public test_case { + int64_t n_embd = 6144, + int64_t n_ff = 2048, + bool already_weighted = false, +- bool f16_act = false) ++ bool f16_act = false, ++ int64_t active_experts = n_expert_used) + : down_type(down_type), + n_embd(n_embd), + n_ff(n_ff), + already_weighted(already_weighted), +- f16_act(f16_act) { ++ f16_act(f16_act), ++ active_experts(active_experts) { + GGML_ASSERT(n_ff % ggml_blck_size(down_type) == 0); + GGML_ASSERT(!f16_act || !already_weighted); ++ GGML_ASSERT(active_experts > 0 && active_experts <= n_expert_used); + } + + std::string vars() override { + return VARS_TO_STR6(down_type, n_expert, n_expert_used, n_tokens, n_embd, n_ff) + "," + +- VARS_TO_STR2(already_weighted, f16_act); ++ VARS_TO_STR3(already_weighted, f16_act, active_experts); + } + + std::string op_desc(ggml_tensor * t) override { + GGML_UNUSED(t); ++ if (down_type == GGML_TYPE_Q3_K && n_embd == 6144 && n_ff == 2048 && !already_weighted && !f16_act) { ++ if (active_experts == 2) { ++ return "GLM_MOE_DOWN_REDUCE_Q3_ACTIVE2"; ++ } ++ if (active_experts == 4) { ++ return "GLM_MOE_DOWN_REDUCE_Q3_ACTIVE4"; ++ } ++ if (active_experts == 6) { ++ return "GLM_MOE_DOWN_REDUCE_Q3_ACTIVE6"; ++ } ++ return "GLM_MOE_DOWN_REDUCE_Q3_ACTIVE8"; ++ } + if (f16_act) { + return "GLM_MOE_DOWN_REDUCE_F16_ACT"; + } +@@ -9197,7 +9222,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) { +- const float value = 0.5f + 0.125f * float((slot * 7 + token * 3) % 11); ++ const float value = slot < active_experts ? ++ 0.5f + 0.125f * float((slot * 7 + token * 3) % 11) : ++ 0.0f; + data[token*n_expert_used + slot] = value; + sum += value; + } +@@ -15033,6 +15060,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)); ++ 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, false, true, false, 6)); + 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() { + 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)); ++ test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048, false, false, 6)); ++ test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048, false, false, 4)); ++ test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048, false, false, 2)); + 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() { + 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)); ++ 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, false, 6)); + 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() { + 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)); ++ test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048, false, false, 6)); ++ test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048, false, false, 4)); ++ test_cases.emplace_back(new test_glm_moe_down_reduce(GGML_TYPE_Q3_K, 6144, 2048, false, false, 2)); + 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_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) + 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/0036-ggml-use-GLM-max-active-policy-for-Q3-down-kernels.patch new file mode 100644 index 0000000000..24832c6f74 --- /dev/null +++ b/third_party/llama.cpp/patches/0036-ggml-use-GLM-max-active-policy-for-Q3-down-kernels.patch @@ -0,0 +1,109 @@ +From 3fc13c2ed2cecf50726e7d10927f5070539c9aa1 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 + +--- + ggml/src/ggml-metal/ggml-metal-ops.cpp | 30 ++++++++++++++++++++++++++ + tests/test-backend-ops.cpp | 2 ++ + 2 files changed, 32 insertions(+) + +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index b4fc441a..2e42a168 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -522,6 +522,11 @@ static bool ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active2_enabled() { + return value && atoi(value) != 0; + } + ++static int ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active_count() { ++ const int active_experts = ggml_metal_glm_dsa_moe_max_active_experts(); ++ return active_experts == 2 || active_experts == 4 || active_experts == 6 ? active_experts : 0; ++} ++ + static bool ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r6_nb8_w0_enabled() { + const char * value = getenv("GGML_METAL_ENABLE_Q3_DOWN_SLOT_PARALLEL_REDUCE_R6_NB8_W0"); + return value && atoi(value) != 0; +@@ -4776,6 +4781,12 @@ static bool ggml_metal_match_mul_mv_id_weighted_reduce( + down->src[0]->ne[0] == 2048 && + down->src[2]->ne[1] == 1 && + ggml_get_op_params_i32(weighted_sum, 0) != 0; ++ const bool q3_active_count_shape = ++ down_type == GGML_TYPE_Q3_K && ++ ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active_count() != 0 && ++ down->src[0]->ne[0] == 2048 && ++ down->src[2]->ne[1] == 1 && ++ ggml_get_op_params_i32(weighted_sum, 0) == 0; + if (down_type == GGML_TYPE_Q3_K && + !ggml_metal_glm_dsa_q3_down_weighted_reduce_tensor_selected(weighted_sum)) { + return reject("q3_tensor"); +@@ -4783,6 +4794,7 @@ static bool ggml_metal_match_mul_mv_id_weighted_reduce( + if (down_type == GGML_TYPE_Q3_K && + !q3_default_w0_shape && + !q3_default_w1_shape && ++ !q3_active_count_shape && + !ggml_metal_glm_dsa_q3_down_weighted_reduce_enabled() && + !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_enabled() && + !ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r8_enabled() && +@@ -5068,6 +5080,14 @@ static int ggml_metal_op_mul_mv_id_weighted_reduce( + ggml_metal_glm_dsa_q3_down_f16_act_enabled() && + ne00 == 2048 && + extra.already_weighted == 0; ++ const int q3_active_count = ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active_count(); ++ const bool q3_active_count_shape = ++ down->src[0]->type == GGML_TYPE_Q3_K && ++ q3_active_count != 0 && ++ ne00 == 2048 && ++ ne20 == 8 && ++ ne21 == 1 && ++ extra.already_weighted == 0; + const bool q2_slot_parallel_requested = ggml_metal_glm_dsa_q2_down_slot_parallel_reduce_enabled(); + const bool q2_slot_parallel_r4 = ggml_metal_glm_dsa_q2_down_slot_parallel_reduce_r4_enabled(); + const bool q2_slot_parallel_r16 = ggml_metal_glm_dsa_q2_down_slot_parallel_reduce_r16_enabled(); +@@ -5183,6 +5203,7 @@ static int ggml_metal_op_mul_mv_id_weighted_reduce( + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active6_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active4_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_active2_enabled() || ++ q3_active_count_shape || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r6_nb8_w0_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_r10_nb8_w0_enabled() || + ggml_metal_glm_dsa_q3_down_slot_parallel_reduce_glm52_w0_enabled() || +@@ -5227,6 +5248,15 @@ static int ggml_metal_op_mul_mv_id_weighted_reduce( + 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); ++ } 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); ++ } 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); ++ } 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); + } 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); +diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp +index 3a2448a8..44f36f73 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 { + 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 || ++ getenv("GGML_METAL_EXPERIMENTAL_GLM_MOE_MAX_ACTIVE_EXPERTS") != nullptr || + 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 { + 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 || ++ getenv("GGML_METAL_EXPERIMENTAL_GLM_MOE_MAX_ACTIVE_EXPERTS") != nullptr || + 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 || +-- +2.54.0 (Apple Git-157) + 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/0037-ggml-honor-explicit-Q2-gate-up-Metal-variant-flags.patch new file mode 100644 index 0000000000..c2d92b5d5b --- /dev/null +++ b/third_party/llama.cpp/patches/0037-ggml-honor-explicit-Q2-gate-up-Metal-variant-flags.patch @@ -0,0 +1,42 @@ +From 8fec61c2696f179dbe598552bef7da0237ebedbb 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 + +--- + ggml/src/ggml-metal/ggml-metal-ops.cpp | 12 +++++++++++- + 1 file changed, 11 insertions(+), 1 deletion(-) + +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index 2e42a168..9469192d 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -4345,6 +4345,15 @@ static int ggml_metal_op_mul_mv_id_gate_up_swiglu( + ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_slot8_split_enabled(); + const bool pair_sg_rowtile_requested = + ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_rowtile_enabled(); ++ const bool pair_sg_special_requested = ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_share_y_enabled() || ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_vecscale_enabled() || ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_q8_act_enabled() || ++ ggml_metal_glm_dsa_q2_gate_up_prequant_q8_enabled() || ++ ggml_metal_glm_dsa_q2_gate_up_inblock_repack_enabled() || ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_half_y_enabled() || ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_r16_enabled() || ++ ggml_metal_glm_dsa_q2_gate_up_swiglu_pair_sg_r12_enabled(); + const bool pair_sg_slot1_dual_default = + pair_sg_default_shape && + !pair_sg_slot8_requested && +@@ -4354,7 +4363,8 @@ static int ggml_metal_op_mul_mv_id_gate_up_swiglu( + !pair_sg_slot4_dual_r16_requested && + !pair_sg_slot2_dual_requested && + !pair_sg_slot8_split_requested && +- !pair_sg_rowtile_requested; ++ !pair_sg_rowtile_requested && ++ !pair_sg_special_requested; + const bool pair_sg_slot8 = pair_sg && dst->type == GGML_TYPE_F32 && + pair_sg_slot8_requested; + const bool pair_sg_slot2 = pair_sg && dst->type == GGML_TYPE_F32 && +-- +2.54.0 (Apple Git-157) + diff --git a/third_party/llama.cpp/patches/0038-ggml-make-GLM-MoE-roofline-honor-active-experts.patch b/third_party/llama.cpp/patches/0038-ggml-make-GLM-MoE-roofline-honor-active-experts.patch new file mode 100644 index 0000000000..85ef7a71ae --- /dev/null +++ b/third_party/llama.cpp/patches/0038-ggml-make-GLM-MoE-roofline-honor-active-experts.patch @@ -0,0 +1,65 @@ +From 6208005ffe6bdc5b2d187f2d6c31138773bfdbc9 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 + +--- + ggml/src/ggml-metal/ggml-metal-ops.cpp | 18 ++++++++++++------ + 1 file changed, 12 insertions(+), 6 deletions(-) + +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index 9469192d..2121075e 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -2026,7 +2026,12 @@ static int ggml_metal_op_glm_moe_q2_selected_weight_scan( + ggml_metal_routed_moe_decode_contract contract; + GGML_ASSERT(ggml_metal_make_routed_moe_decode_contract("glm_dsa", motif, contract)); + ++ const int max_active_experts = ggml_metal_glm_dsa_moe_max_active_experts(); ++ const int active_top_k = ++ max_active_experts > 0 && max_active_experts < contract.top_k ? ++ max_active_experts : int(contract.top_k); + const uint64_t top_k = uint64_t(contract.top_k); ++ const uint64_t active_top_k_u64 = uint64_t(active_top_k); + const uint64_t q2_k_block_bytes = ggml_type_size(GGML_TYPE_Q2_K); + const auto stored_bytes = [storage_block_bytes, q2_k_block_bytes](uint64_t bytes, ggml_type type) { + if (type != GGML_TYPE_Q2_K) { +@@ -2043,7 +2048,7 @@ static int ggml_metal_op_glm_moe_q2_selected_weight_scan( + /*.down_expert_bytes =*/ stored_bytes(contract.selected_down_weight_bytes/top_k, contract.down_w->type), + /*.down_expert_stride =*/ contract.down_w->nb[2], + /*.ids_nb1 =*/ contract.ids->nb[1], +- /*.top_k =*/ int32_t(contract.top_k), ++ /*.top_k =*/ int32_t(active_top_k), + /*.n_tokens =*/ int32_t(contract.n_tokens), + /*.n_experts =*/ int32_t(contract.n_expert), + /*.chunks_per_expert =*/ chunks_per_expert, +@@ -2062,7 +2067,7 @@ static int ggml_metal_op_glm_moe_q2_selected_weight_scan( + ggml_metal_encoder_dispatch_threadgroups( + ctx->enc, + chunks_per_expert, +- int(contract.top_k), ++ active_top_k, + int(contract.n_tokens), + 256, + 1, +@@ -2073,12 +2078,13 @@ static int ggml_metal_op_glm_moe_q2_selected_weight_scan( + GGML_LOG_INFO( + "ggml_metal: moe_dispatch op=glm_moe_q2_selected_weight_roofline tensor=%s selected_weight_bytes=%llu storage_block_bytes=%d chunks_per_expert=%d threadgroups=%lld threads_per_threadgroup=256\n", + ggml_metal_tensor_name(contract.out), +- (unsigned long long) (stored_bytes(contract.selected_gate_weight_bytes, contract.gate_w->type) + +- stored_bytes(contract.selected_up_weight_bytes, contract.up_w->type) + +- stored_bytes(contract.selected_down_weight_bytes, contract.down_w->type)), ++ (unsigned long long) (active_top_k_u64 * ( ++ stored_bytes(contract.selected_gate_weight_bytes/top_k, contract.gate_w->type) + ++ stored_bytes(contract.selected_up_weight_bytes/top_k, contract.up_w->type) + ++ stored_bytes(contract.selected_down_weight_bytes/top_k, contract.down_w->type))), + storage_block_bytes, + chunks_per_expert, +- contract.top_k*contract.n_tokens*chunks_per_expert); ++ active_top_k*contract.n_tokens*chunks_per_expert); + } + + return motif.n_fuse; +-- +2.54.0 (Apple Git-157) + 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/0039-ggml-avoid-GLM-Q3-fused-tail-under-active-policy.patch new file mode 100644 index 0000000000..6bb6fadf3e --- /dev/null +++ b/third_party/llama.cpp/patches/0039-ggml-avoid-GLM-Q3-fused-tail-under-active-policy.patch @@ -0,0 +1,98 @@ +From dd76e903b54de937336fbea5f1758753a379109f 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 + +--- + ggml/src/ggml-metal/ggml-metal-ops.cpp | 3 ++ + tests/test-backend-ops.cpp | 46 ++++++++++++++++++++++++++ + 2 files changed, 49 insertions(+) + +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index 2121075e..91976954 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -2738,6 +2738,9 @@ static bool ggml_metal_glm_moe_can_fuse_swiglu_q3_down(const ggml_metal_glm_moe_ + if (!ggml_metal_glm_dsa_moe_swiglu_q3_down_fusion_enabled()) { + return false; + } ++ if (ggml_metal_glm_dsa_moe_max_active_experts() > 0) { ++ return false; ++ } + if (motif.has_weighted_down || + 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 +--- 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 + } + }; + ++struct test_glm_moe_routed_q2q3_final_output : public test_topk_moe_glm_component { ++ const bool use_native_route; ++ const bool use_direct_top_k; ++ ++ explicit test_glm_moe_routed_q2q3_final_output(bool native_route = false, bool direct_top_k = false) ++ : test_topk_moe_glm_component( ++ 6, ++ GGML_TYPE_Q2_K, ++ GGML_TYPE_Q2_K, ++ GGML_TYPE_Q3_K, ++ 6144, ++ 2048, ++ true, ++ false, ++ native_route, ++ direct_top_k), ++ use_native_route(native_route), ++ use_direct_top_k(direct_top_k) {} ++ ++ std::string op_desc(ggml_tensor * t) override { ++ GGML_UNUSED(t); ++ if (use_direct_top_k && use_native_route) { ++ return "GLM_MOE_ROUTED_MOTIF_Q2Q3_DIRECT_TOPK_NATIVE_ROUTE_GLM"; ++ } ++ if (use_direct_top_k) { ++ return "GLM_MOE_ROUTED_MOTIF_Q2Q3_DIRECT_TOPK_GLM"; ++ } ++ if (use_native_route) { ++ return "GLM_MOE_ROUTED_MOTIF_Q2Q3_NATIVE_ROUTE_GLM"; ++ } ++ return "GLM_MOE_ROUTED_MOTIF_Q2Q3_GLM"; ++ } ++ ++ std::vector fusion_test_nodes() override { ++ return { selected_experts, weights, ffn_out }; ++ } ++}; ++ + 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() { + 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)); ++ test_cases.emplace_back(new test_glm_moe_routed_q2q3_final_output(false, false)); ++ test_cases.emplace_back(new test_glm_moe_routed_q2q3_final_output(true, false)); ++ test_cases.emplace_back(new test_glm_moe_routed_q2q3_final_output(false, true)); ++ test_cases.emplace_back(new test_glm_moe_routed_q2q3_final_output(true, true)); + 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() { + 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)); ++ test_cases.emplace_back(new test_glm_moe_routed_q2q3_final_output(false, false)); ++ test_cases.emplace_back(new test_glm_moe_routed_q2q3_final_output(true, false)); ++ test_cases.emplace_back(new test_glm_moe_routed_q2q3_final_output(false, true)); ++ test_cases.emplace_back(new test_glm_moe_routed_q2q3_final_output(true, true)); + 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")); +-- +2.54.0 (Apple Git-157) + diff --git a/third_party/llama.cpp/patches/0040-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch b/third_party/llama.cpp/patches/0040-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch new file mode 100644 index 0000000000..088eff5174 --- /dev/null +++ b/third_party/llama.cpp/patches/0040-ggml-shrink-GLM-Q2-gate-up-active-dispatch.patch @@ -0,0 +1,49 @@ +From aacb3421872735fcd2f1a22e4560a60b44f53046 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 + +--- + ggml/src/ggml-metal/ggml-metal-ops.cpp | 6 ++++-- + ggml/src/ggml-metal/ggml-metal.metal | 6 ++++-- + 2 files changed, 8 insertions(+), 4 deletions(-) + +diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp +index 91976954..646c983a 100644 +--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp ++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp +@@ -4556,9 +4556,11 @@ 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 int64_t ne123 = pair_sg ? +- ((ne20 + slots_per_threadgroup - 1)/slots_per_threadgroup)*ne21 : +- ne20*ne21; ++ ((active_slots + slots_per_threadgroup - 1)/slots_per_threadgroup)*ne21 : ++ active_slots*ne21; + const int grid_x = pair_sg ? (ne01 + nr0 - 1)/nr0 : (ne01 + nr0*nsg - 1)/(nr0*nsg); + 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 +--- 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( + 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 slot_groups = (active_slots + slots_per_tg - 1) / slots_per_tg; + const int r0 = tgpig.x; +- const int token = tgpig.z / ((args.nei0 + slots_per_tg - 1) / slots_per_tg); +- const int slot_group = tgpig.z % ((args.nei0 + slots_per_tg - 1) / slots_per_tg); ++ const int token = tgpig.z / slot_groups; ++ const int slot_group = tgpig.z % slot_groups; + const int slot = slot_group * slots_per_tg + int(sgitg / 2); + const bool is_gate = (sgitg & 1) != 0; + const int first_row = r0 * nr0; +-- +2.54.0 (Apple Git-157) + diff --git a/third_party/llama.cpp/patches/0041-ggml-skip-inactive-GLM-Q3-down-slots.patch b/third_party/llama.cpp/patches/0041-ggml-skip-inactive-GLM-Q3-down-slots.patch new file mode 100644 index 0000000000..249be4fc9b --- /dev/null +++ b/third_party/llama.cpp/patches/0041-ggml-skip-inactive-GLM-Q3-down-slots.patch @@ -0,0 +1,25 @@ +From 00fa96c1d75aa4ee9885eb64d88cbebe861d9fe5 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 + +--- + 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 +--- 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( + float sumf1[nr0] = {0.f}; + float sumf2[nr0] = {0.f}; + +- if (slot < args.nei0 && token < args.nei1) { ++ if (slot < n_slots && slot < args.nei0 && token < args.nei1) { + const int32_t expert_id = ((device const int32_t *) (ids + token*args.nbi1))[slot]; + 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) +