models : fix incorrect uses of get_key_or_arr (#28868 )

models : fix mimo2 swa pattern load (#28865 )
tests(s390x): add non-vxe build to tests (#28776 )
2026-09-14 13:06:47 +02:00 · 2026-09-14 14:05:37 +03:00 · 2026-09-14 14:05:17 +03:00 · 2026-09-14 14:04:58 +03:00 · 2026-09-14 14:04:05 +03:00 · 2026-09-14 14:02:44 +03:00
43 changed files with 1085 additions and 311 deletions
@@ -34,10 +34,15 @@ env:
  LLAMA_ARG_LOG_TIMESTAMPS: 1

 jobs:
-
  ubuntu-24-s390x:
+    name: ubuntu-24-s390x (VXE ${{ matrix.vxe }})
    runs-on: ubuntu-24.04-s390x

+    strategy:
+      fail-fast: false
+      matrix:
+        vxe: ["ON", "OFF"]  # `-DGGML_VXE=ON/OFF`
+
    steps:
      - name: Clone
        id: checkout
@@ -77,7 +82,8 @@ jobs:
        run: |
          cmake -B build \
            -DLLAMA_FATAL_WARNINGS=ON \
-            -DGGML_RPC=ON
+            -DGGML_RPC=ON \
+            -DGGML_VXE=${{ matrix.vxe }}
          time cmake --build build --config Release -j $(nproc)

      - name: Test
@@ -7,6 +7,7 @@ on:
      - master
    paths: [
      '.github/workflows/build-self-hosted.yml',
+      'ci/run.sh',
      '**/CMakeLists.txt',
      '**/.cmake',
      '**/*.h',
@@ -27,6 +28,7 @@ on:
    types: [opened, synchronize, reopened]
    paths: [
      '.github/workflows/build-self-hosted.yml',
+      'ci/run.sh',
      '**/CMakeLists.txt',
      '**/.cmake',
      '**/*.h',
@@ -15,7 +15,7 @@ concurrency:

 jobs:
  model-naming:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim
    steps:
      - uses: actions/checkout@v6
      - name: Check model naming conventions
@@ -15,7 +15,7 @@ concurrency:

 jobs:
  editorconfig:
-    runs-on: [self-hosted, fast]
+    runs-on: ubuntu-slim
    steps:
      - uses: actions/checkout@v6
      - uses: editorconfig-checker/action-editorconfig-checker@840e866d93b8e032123c23bac69dece044d4d84c # v2.2.0
@@ -6,6 +6,7 @@ General:
 - PR and commit titles format: `<module> : <title>`. Lookup recents for examples
 - Don't try to build or run the code unless you are explicitly asked to do so
 - Use the `gh` CLI tool when querying PRs, issues, or other GitHub resources
+- When [MODEL] is needed, first try to get it from the `PI_MODEL_NAME` env var before asking the user

 Coding:
 - When in doubt, always refer to the CONTRIBUTING.md file of the project
@@ -20,7 +21,7 @@ Pull requests (PRs):
 - Don't explicitly wrap lines in the PR description (each paragraph and bullet is a single line)
 - When creating a pull request, look for the repository's PR template and follow it
 - For the AI usage disclosure section, write "YES. pi:llama.cpp/[MODEL]"
- Ask the user to tell you what model was used and write it in place of [MODEL]
+- If `PI_MODEL_NAME` env var is not set, ask the user to tell you what model was used and write it in place of [MODEL]
 - Always create the pull requests in draft mode

 Commits:
@@ -58,8 +58,6 @@ if [ -n "${GG_BUILD_ROCM}" ] && [ -n "${GITHUB_RUN_ID}" ]; then
 fi

 rm -f $OUT/*.log
-rm -f $OUT/*.exit
-rm -f $OUT/*.md

 sd=`dirname $0`
 cd $sd/../
@@ -211,10 +209,6 @@ function gg_wget {
    cd $cwd
 }

-function gg_printf {
-    printf -- "$@" >> $OUT/README.md
-}
-
 function gg_run {
    ci=$1

@@ -223,13 +217,10 @@ function gg_run {

    gg_run_$ci | tee $OUT/$ci.log
    cur=$?
-    echo "$cur" > $OUT/$ci.exit

    set +x
    set +o pipefail

-    gg_sum_$ci
-
    ret=$((ret | cur))
 }

@@ -255,17 +246,6 @@ function gg_run_ctest_debug {
    set +e
 }

-function gg_sum_ctest_debug {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs ctest in debug mode\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-ctest.log)"
-    gg_printf '```\n'
-    gg_printf '\n'
-}
-
 # ctest_release

 function gg_run_ctest_release {
@@ -290,16 +270,6 @@ function gg_run_ctest_release {
    set +e
 }

-function gg_sum_ctest_release {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs ctest in release mode\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-ctest.log)"
-    gg_printf '```\n'
-}
-
 # test_llama_archs_tensor_split

 function gg_run_test_llama_archs_tensor_split {
@@ -324,16 +294,6 @@ function gg_run_test_llama_archs_tensor_split {
    set +e
 }

-function gg_sum_test_llama_archs_tensor_split {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs test-llama-archs with 1 to 4 devices\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}.log)"
-    gg_printf '```\n'
-}
-
 # test_llama_archs_models

 function gg_run_test_llama_archs_models {
@@ -353,16 +313,6 @@ function gg_run_test_llama_archs_models {
    set +e
 }

-function gg_sum_test_llama_archs_models {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Generates the dummy models used by the model-dependent tests\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}.log)"
-    gg_printf '```\n'
-}
-
 # test_scripts

 function gg_run_test_scripts {
@@ -376,17 +326,6 @@ function gg_run_test_scripts {
    set +e
 }

-function gg_sum_test_scripts {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs test scripts\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-scripts.log)"
-    gg_printf '```\n'
-    gg_printf '\n'
-}
-
 function gg_get_model {
    #local gguf_0="$MNT/models/qwen3/0.6B/ggml-model-f16.gguf"
    local gguf_0="$MNT/models/qwen3/0.6B/ggml-model-q4_0.gguf"
@@ -430,26 +369,6 @@ function gg_run_ctest_with_model_release {
    cd ..
 }

-function gg_sum_ctest_with_model_debug {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs ctest with model files in debug mode\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-ctest.log)"
-    gg_printf '```\n'
-}
-
-function gg_sum_ctest_with_model_release {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs ctest with model files in release mode\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-ctest.log)"
-    gg_printf '```\n'
-}
-
 # qwen3_0_6b

 function gg_run_qwen3_0_6b {
@@ -554,50 +473,24 @@ function gg_run_qwen3_0_6b {
        return 0
    }

-    check_ppl "f16"  "$(cat $OUT/${ci}-tg-f16.log  | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
+    check_ppl "f16"  "$(cat $OUT/${ci}-tg-f16.log  | grep "^\[1\]")"
    if [ -z ${GG_BUILD_NO_BF16} ]; then
-        check_ppl "bf16" "$(cat $OUT/${ci}-tg-bf16.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
+        check_ppl "bf16" "$(cat $OUT/${ci}-tg-bf16.log | grep "^\[1\]")"
    fi
-    check_ppl "q8_0" "$(cat $OUT/${ci}-tg-q8_0.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q4_0" "$(cat $OUT/${ci}-tg-q4_0.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q4_1" "$(cat $OUT/${ci}-tg-q4_1.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q5_0" "$(cat $OUT/${ci}-tg-q5_0.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q5_1" "$(cat $OUT/${ci}-tg-q5_1.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-   #check_ppl "q2_k" "$(cat $OUT/${ci}-tg-q2_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log # note: ppl > 20.0 for this quant and model
-    check_ppl "q3_k" "$(cat $OUT/${ci}-tg-q3_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q4_k" "$(cat $OUT/${ci}-tg-q4_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q5_k" "$(cat $OUT/${ci}-tg-q5_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-    check_ppl "q6_k" "$(cat $OUT/${ci}-tg-q6_k.log | grep "^\[1\]")" | tee -a $OUT/${ci}-ppl.log
-
-    cat $OUT/${ci}-imatrix.log | grep "Final" >> $OUT/${ci}-imatrix-sum.log
+    check_ppl "q8_0" "$(cat $OUT/${ci}-tg-q8_0.log | grep "^\[1\]")"
+    check_ppl "q4_0" "$(cat $OUT/${ci}-tg-q4_0.log | grep "^\[1\]")"
+    check_ppl "q4_1" "$(cat $OUT/${ci}-tg-q4_1.log | grep "^\[1\]")"
+    check_ppl "q5_0" "$(cat $OUT/${ci}-tg-q5_0.log | grep "^\[1\]")"
+    check_ppl "q5_1" "$(cat $OUT/${ci}-tg-q5_1.log | grep "^\[1\]")"
+   #check_ppl "q2_k" "$(cat $OUT/${ci}-tg-q2_k.log | grep "^\[1\]")" # note: ppl > 20.0 for this quant and model
+    check_ppl "q3_k" "$(cat $OUT/${ci}-tg-q3_k.log | grep "^\[1\]")"
+    check_ppl "q4_k" "$(cat $OUT/${ci}-tg-q4_k.log | grep "^\[1\]")"
+    check_ppl "q5_k" "$(cat $OUT/${ci}-tg-q5_k.log | grep "^\[1\]")"
+    check_ppl "q6_k" "$(cat $OUT/${ci}-tg-q6_k.log | grep "^\[1\]")"

    set +e
 }

-function gg_sum_qwen3_0_6b {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Qwen3 0.6B:\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '- perplexity:\n%s\n' "$(cat $OUT/${ci}-ppl.log)"
-    gg_printf '- imatrix:\n```\n%s\n```\n' "$(cat $OUT/${ci}-imatrix-sum.log)"
-    gg_printf '- f16:\n```\n%s\n```\n'  "$(cat $OUT/${ci}-tg-f16.log)"
-    if [ -z ${GG_BUILD_NO_BF16} ]; then
-        gg_printf '- bf16:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-bf16.log)"
-    fi
-    gg_printf '- q8_0:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q8_0.log)"
-    gg_printf '- q4_0:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q4_0.log)"
-    gg_printf '- q4_1:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q4_1.log)"
-    gg_printf '- q5_0:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q5_0.log)"
-    gg_printf '- q5_1:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q5_1.log)"
-    gg_printf '- q2_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q2_k.log)"
-    gg_printf '- q3_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q3_k.log)"
-    gg_printf '- q4_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q4_k.log)"
-    gg_printf '- q5_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q5_k.log)"
-    gg_printf '- q6_k:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q6_k.log)"
-    gg_printf '- save-load-state: \n```\n%s\n```\n' "$(cat $OUT/${ci}-save-load-state.log)"
-}
-
 # bge-small

 function gg_run_embd_bge_small {
@@ -639,15 +532,6 @@ function gg_run_embd_bge_small {
    set +e
 }

-function gg_sum_embd_bge_small {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'BGE Small (BERT):\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '- f16: \n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-f16.log)"
-    gg_printf '- q8_0:\n```\n%s\n```\n' "$(cat $OUT/${ci}-tg-q8_0.log)"
-}
-
 # rerank_tiny

 function gg_run_rerank_tiny {
@@ -704,66 +588,58 @@ function gg_run_rerank_tiny {
    set +e
 }

-function gg_sum_rerank_tiny {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Rerank Tiny (Jina):\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '- f16: \n```\n%s\n```\n' "$(cat $OUT/${ci}-rk-f16.log)"
-}
-
 function gg_check_build_requirements {
    if ! command -v git &> /dev/null; then
-        gg_printf 'git not found, please install\n'
+        echo 'git not found, please install'
        exit 1
    fi

    if ! command -v git-lfs &> /dev/null; then
-        gg_printf 'git-lfs not found, please install\n'
+        echo 'git-lfs not found, please install'
        exit 1
    fi

    if ! git config --get filter.lfs.clean &> /dev/null; then
-        gg_printf 'git-lfs not initialized, please run `git lfs install`\n'
+        echo 'git-lfs not initialized, please run `git lfs install`'
        exit 1
    fi

    if ! command -v wget &> /dev/null; then
-        gg_printf 'wget not found, please install\n'
+        echo 'wget not found, please install'
        exit 1
    fi

    if ! command -v python3 &> /dev/null; then
-        gg_printf 'python3 not found, please install\n'
+        echo 'python3 not found, please install'
        exit 1
    fi

    if ! command -v pip3 &> /dev/null; then
-        gg_printf 'pip3 not found, please install\n'
+        echo 'pip3 not found, please install'
        exit 1
    fi

    if ! python3 -m ensurepip --help &> /dev/null; then
-        gg_printf 'ensurepip not found, please install python3-venv package\n'
+        echo 'ensurepip not found, please install python3-venv package'
        exit 1
    fi

    if ! command -v cmake &> /dev/null; then
-        gg_printf 'cmake not found, please install\n'
+        echo 'cmake not found, please install'
        exit 1
    fi

    if ! command -v ccache &> /dev/null; then
-        gg_printf 'ccache not found, please consider installing for faster builds\n'
+        echo 'ccache not found, please consider installing for faster builds'
    fi

    if ! command -v ctest &> /dev/null; then
-        gg_printf 'ctest not found, please install\n'
+        echo 'ctest not found, please install'
        exit 1
    fi

    if ! command -v unzip &> /dev/null; then
-        gg_printf 'unzip not found, please install\n'
+        echo 'unzip not found, please install'
        exit 1
    fi
 }
@@ -803,17 +679,6 @@ function gg_run_test_backend_ops {
    set +e
 }

-function gg_sum_test_backend_ops {
-    gg_printf '### %s\n\n' "${ci}"
-
-    gg_printf 'Runs test-backend-ops\n'
-    gg_printf '- status: %s\n' "$(cat $OUT/${ci}.exit)"
-    gg_printf '```\n'
-    gg_printf '%s\n' "$(cat $OUT/${ci}-test-backend-ops.log)"
-    gg_printf '```\n'
-    gg_printf '\n'
-}
-
 ## main

 export LLAMA_ARG_LOG_PREFIX=1
@@ -861,6 +726,4 @@ if [ -z ${GG_BUILD_LOW_PERF} ]; then
    test $ret -eq 0 && gg_run ctest_with_model_release
 fi

-cat $OUT/README.md
-
 exit $ret
@@ -1586,6 +1586,11 @@ common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) {
        return COMMON_CONTEXT_SEQ_RM_TYPE_NO;
    }

+    if (llama_n_rs_seq(ctx) > 0) {
+        COM_TRC("%s", "the context supports bounded partial sequence removal\n");
+        return COMMON_CONTEXT_SEQ_RM_TYPE_RS;
+    }
+
    common_context_seq_rm_type res = COMMON_CONTEXT_SEQ_RM_TYPE_PART;

    llama_memory_clear(mem, true);
@@ -1602,12 +1607,6 @@ common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) {
        goto done;
    }

-    if (llama_n_rs_seq(ctx) > 0) {
-        COM_TRC("%s", "the context supports bounded partial sequence removal\n");
-        res = COMMON_CONTEXT_SEQ_RM_TYPE_RS;
-        goto done;
-    }
-
    // try to remove the last tokens
    if (!llama_memory_seq_rm(mem, 0, 1, -1)) {
        COM_TRC("%s", "the context does not support partial sequence removal\n");
@@ -168,6 +168,7 @@ TEXT_MODEL_MAP: dict[str, str] = {
    "Mamba2ForCausalLM": "mamba",
    "MambaForCausalLM": "mamba",
    "MambaLMHeadModel": "mamba",
+    "MapleForCausalLM": "maple",
    "MellumForCausalLM": "mellum",
    "MiMoV2FlashForCausalLM": "mimo",
    "MiMoV2ForCausalLM": "mimo",
@@ -0,0 +1,87 @@
+from __future__ import annotations
+
+from typing import Iterable, TYPE_CHECKING, cast
+
+import torch
+
+if TYPE_CHECKING:
+    from torch import Tensor
+
+from .base import LazyTorchTensor, ModelBase, TextModel, gguf
+
+
+@ModelBase.register("MapleForCausalLM")
+@ModelBase.example("deepgrove/maple-preview")
+class MapleModel(TextModel):
+    model_arch = gguf.MODEL_ARCH.MAPLE
+
+    def set_gguf_parameters(self):
+        super().set_gguf_parameters()
+        hparams = self.hparams
+
+        assert hparams["hidden_act"] == "silu"
+        assert hparams.get("num_shared_experts", 0) == 0
+        assert hparams.get("norm_topk_prob", True)
+        assert hparams.get("nope_on_global_attention", False)
+
+        head_dim = hparams.get("head_dim", hparams["hidden_size"] // hparams["num_attention_heads"])
+        partial_rotary_factor = self.rope_parameters.get("partial_rotary_factor", 1.0)
+
+        self.gguf_writer.add_vocab_size(hparams["vocab_size"])
+        self.gguf_writer.add_rope_dimension_count(int(head_dim * partial_rotary_factor))
+        self.gguf_writer.add_sliding_window(hparams["sliding_window"])
+        self.gguf_writer.add_sliding_window_pattern([layer_type == "sliding_attention" for layer_type in hparams["layer_types"]])
+        self.gguf_writer.add_expert_feed_forward_length(hparams["moe_intermediate_size"])
+        # the reference clamps the MoE SwiGLU gate/up at 7.0 (modeling_maple.py)
+        self.gguf_writer.add_swiglu_clamp_exp([7.0] * self.block_count)
+
+    _experts: list[dict[str, Tensor]] | None = None
+
+    @staticmethod
+    def _stack_experts(tensors: list[Tensor]) -> Tensor:
+        shape = (len(tensors), *tensors[0].shape)
+        dtype = tensors[0].dtype
+        meta = LazyTorchTensor.meta_with_dtype_and_shape(dtype, shape)
+
+        # tensors goes through args, not the closure, so that `func` matches
+        # LazyBase's single-argument shape
+        def stack(ts: list[Tensor]) -> Tensor:
+            result = torch.empty(shape, dtype=dtype)
+            for expert_id, tensor in enumerate(ts):
+                result[expert_id].copy_(LazyTorchTensor.to_eager(tensor))
+            ts.clear()
+            return result
+
+        return cast(torch.Tensor, LazyTorchTensor(meta=meta, args=(tensors,), func=stack))
+
+    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
+        if "mlp.experts" in name:
+            n_experts = self.hparams["num_experts"]
+            assert bid is not None
+
+            if self._experts is None:
+                self._experts = [{} for _ in range(self.block_count)]
+
+            self._experts[bid][name] = data_torch
+
+            if len(self._experts[bid]) >= n_experts * 3:
+                for weight_name in ("down_proj", "gate_proj", "up_proj"):
+                    tensors = []
+
+                    for expert_id in range(n_experts):
+                        expert_name = f"model.layers.{bid}.mlp.experts.{expert_id}.{weight_name}.weight"
+                        tensors.append(self._experts[bid].pop(expert_name))
+
+                    merged_name = f"model.layers.{bid}.mlp.experts.{weight_name}.weight"
+                    yield from super().modify_tensors(self._stack_experts(tensors), merged_name, bid)
+            return
+
+        yield from super().modify_tensors(data_torch, name, bid)
+
+    def prepare_tensors(self):
+        super().prepare_tensors()
+
+        if self._experts is not None:
+            experts = [name for layer in self._experts for name in layer]
+            if experts:
+                raise ValueError(f"Unprocessed experts: {experts}")
@@ -790,14 +790,15 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm

 | Name              | Value            | Function                                                                                                                  |
 |-------------------|------------------|---------------------------------------------------------------------------------------------------------------------------|
-| GGML_SYCL_DEBUG   | 0 (default) or 1 | Enable log function by macro: GGML_SYCL_DEBUG                                                                             |
+| GGML_SYCL_DEBUG   | 0 (default) or 1 | Enable log function: GGML_SYCL_DEBUG() for common debug. |
+| GGML_SYCL_DEV_DEBUG   | 0 (default) or 1 | Enable log function: GGML_SYCL_DEV_DEBUG() for developmental purposes by replacing GGML_SYCL_DEBUG() in special codes. Restore to GGML_SYCL_DEBUG() before committing code.|
 | GGML_SYCL_DEV2DEV_MEMCPY | 0 (default), 1, 2 | Choose the method of dev2dev memory copy.<br>Value: <br>*  0: SYCL API (default), only support dGPUs.<br>* 1: L0 API -- Better performance, only support dGPUs, found to lead to abnormal crash in some case. <br>* 2: Host Forward -- Most stable method for all cases (including iGPU + dGPU*N), but with lower performance (-2% to -5%).<br>SYCL & L0 API are easy to be impacted by Intel GPU driver issue. When you meet the garbled output or crash issues in multiple GPUs case, try with this debug flag to work around or check the issue.|
 | GGML_SYCL_ENABLE_FLASH_ATTN | 1 (default) or 0| Enable Flash-Attention. It can reduce memory usage. The performance impact depends on the LLM.|
 | GGML_SYCL_ENABLE_OPT | 0 or 1 (default)| Enable optimize features for Intel GPUs. (Recommended to 0 for Intel devices older than Gen 10) |
 | GGML_SYCL_ENABLE_GRAPH | 0 (default) or 1 | Enable running computations through SYCL Graphs feature. Disabled by default because SYCL Graph is still on development, no better performance. |
 | GGML_SYCL_ENABLE_HOST_PINNED_MEM | 0 or 1 (default) | Enable host pinned memory to speed up copy data from host to device. When disable it, host memory will common malloc() on CPU. Disable it when use `--load-model mlock`.|
 | GGML_SYCL_HOST_PINNED_MEM_2G | 0 (default) or 1 | Limit the max memory allocation to be no more than 2GB when enable host pinned memory. USM allocations above 2 GiB take the relaxed/large-allocation path, which serializes H2D copies with compute and prevents copy/compute overlap. It will impact the startup time. Need more test. Depend on `GGML_SYCL_ENABLE_HOST_PINNED_MEM=1`.|
-| GGML_SYCL_GET_MEM_API | 0 (default) or 1  | Set to get memory info (free, total) by Level Zero or SYCL API:<br>0 - Level Zero API: support more GPUs, only run on Level Zero running time. When there is an error, fallback to call SYCL API. Depend on GGML_SYCL_SUPPORT_LEVEL_ZERO_API.<br>1 - SYCL API: legacy, support more running time, it can't get the free size of some GPUs (like Arc770). In such case, return total size for free size.|
+| GGML_SYCL_GET_MEM_API | 0 (default) or 1  | Set to get memory info (free, total) by Level Zero or SYCL API:<br>0 - Level Zero API: support more GPUs, only run on Level Zero running time. When there is an error, fallback to call SYCL API. Depend on GGML_SYCL_SUPPORT_LEVEL_ZERO_API.<br>1 - SYCL API: legacy, support more running time, it can't get the free size of some GPUs (like Arc770). In such case, return the free size as value of total size.|
 | GGML_SYCL_USE_LEVEL_ZERO_API | 1 (default) or 0 | Use Level Zero API for device memory allocation instead of SYCL. Reduces system RAM usage on Intel dGPUs by avoiding DMA-buf/TTM host memory staging. Requires GGML_SYCL_SUPPORT_LEVEL_ZERO_API=ON at build time. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).|
 | GGML_SYCL_ENABLE_DNN | 0 or 1 (default)| Enable running computations through oneDNN and always use oneMKL. |
 | GGML_SYCL_FA_ONEDNN | 1 (default) or 0 | Enable the oneDNN fused SDPA (flash-attention) path on supported GPUs. Set to 0 to always use the native SYCL flash-attention kernel. |
@@ -675,12 +675,6 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
    target_compile_options(${GGML_CPU_NAME} PRIVATE ${ARCH_FLAGS})
    target_compile_definitions(${GGML_CPU_NAME} PRIVATE ${ARCH_DEFINITIONS})

-    if (CMAKE_C_COMPILER_ID STREQUAL "GNU" AND NOT GGML_SYSTEM_ARCH STREQUAL "x86")
-        message(STATUS "Skipping PCH for ${GGML_CPU_NAME}: GCC PCH is only enabled for x86 (arch: ${GGML_SYSTEM_ARCH})")
-    else()
-        target_precompile_headers(${GGML_CPU_NAME} PRIVATE ggml-impl.h)
-    endif()
-
    if (EMSCRIPTEN)
        set_target_properties(${GGML_CPU_NAME} PROPERTIES COMPILE_FLAGS "-msimd128")
    endif()
@@ -417,6 +417,7 @@ void ggml_vec_dot_mxfp4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const vo
    sumf = vec_hsum_f32x4(v_acc);
    *s = sumf;
 #else
+    UNUSED(nb);
    UNUSED(x);
    UNUSED(y);
    UNUSED(ib);
@@ -70,6 +70,7 @@ void ggml_quantize_mat_q8_0_4x4(const float * GGML_RESTRICT x, void * GGML_RESTR
 #endif
 }

+#if defined(__VXE__) || defined(__VXE2__)
 static inline int16x8_t vxe_dot_acc(const int8x16_t v_x, const int8x16_t v_y, const int16x8_t v_acc) {
    return vec_meadd(v_x, v_y, vec_moadd(v_x, v_y, v_acc));
 }
@@ -84,6 +85,7 @@ static inline int32x4_t vxe_fold(const int16x8_t v_sumi) {
    const int16x8_t v_ones = vec_splats((int16_t)1);
    return vec_add(vec_mule(v_sumi, v_ones), vec_mulo(v_sumi, v_ones));
 }
+#endif

 void ggml_gemv_q4_0_4x4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int nr, int nc) {
    const int qk = QK8_0;
@@ -5,10 +5,10 @@
 //
 // cache line
 //
-
-#if defined(__cpp_lib_hardware_interference_size)
-#define CACHE_LINE_SIZE std::hardware_destructive_interference_size
-#else
+// TODO: rework CACHE_LINE_SIZE so std::hardware_destructive_interference_size
+// can be used consistently between C and C++ TUs; the previous macro form
+// diverged based on include order and undersized the work buffer.
+// ref: https://github.com/ggml-org/llama.cpp/pull/28882
 #if defined(__POWER9_VECTOR__)
 #define CACHE_LINE_SIZE 128
 #elif defined(__VXE__) || defined(__VXE2__)
@@ -16,17 +16,8 @@
 #else
 #define CACHE_LINE_SIZE 64
 #endif
-#endif

-// -Winterference-size was introduced in GCC 12
-#if defined(__cplusplus) && defined(__GNUC__) && !defined(__clang__) && __GNUC__ >= 12
-#pragma GCC diagnostic push
-#pragma GCC diagnostic ignored "-Winterference-size"
-#endif
 static const size_t CACHE_LINE_SIZE_F32 = CACHE_LINE_SIZE/sizeof(float);
-#if defined(__cplusplus) && defined(__GNUC__) && !defined(__clang__) && __GNUC__ >= 12
-#pragma GCC diagnostic pop
-#endif

 // Work buffer size for im2col operations in CONV2D
 #define GGML_IM2COL_WORK_SIZE (16 * 1024 * 1024)
@@ -329,6 +329,12 @@ static bool fp16_mma_hardware_available(const int cc) {
        (GGML_CUDA_CC_IS_MTHREADS(cc) && cc >= GGML_CUDA_CC_QY2);
 }

+// To be used for feature selection of external libraries, e.g. cuBLAS.
+static bool fast_bf16_hardware_available(const int cc) {
+        return (GGML_CUDA_CC_IS_AMD(cc) && (cc >= GGML_CUDA_CC_RDNA3 || GGML_CUDA_CC_IS_CDNA(cc)))
+            || (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE);
+}
+
 static bool bf16_mma_hardware_available(const int cc) {
    return (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_AMPERE) ||
        GGML_CUDA_CC_IS_CDNA(cc) || cc >= GGML_CUDA_CC_RDNA3 ||
@@ -1620,11 +1620,19 @@ static void ggml_cuda_mul_mat_cublas_impl(ggml_backend_cuda_context & ctx, const
 }

 static void ggml_cuda_mul_mat_cublas(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
+    const int cc = ggml_cuda_info().devices[ctx.device].cc;
    ggml_type compute_type = src0->type;
    if (ggml_is_quantized(compute_type)) {
-        compute_type = fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
-    } else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(ggml_cuda_info().devices[ctx.device].cc)) {
+        compute_type = fast_fp16_hardware_available(cc) ? GGML_TYPE_F16 : GGML_TYPE_F32;
+    } else if (compute_type == GGML_TYPE_F16 && !fast_fp16_hardware_available(cc)) {
        compute_type = GGML_TYPE_F32;
+    } else if (compute_type == GGML_TYPE_BF16 && !fast_bf16_hardware_available(cc)) {
+        if (GGML_CUDA_CC_IS_AMD(cc) && src1->ne[1] > 32) {
+            compute_type = GGML_TYPE_F32;
+        }
+        if (GGML_CUDA_CC_IS_NVIDIA(cc) && src1->ne[1] > (cc >= GGML_CUDA_CC_VOLTA ? 8 : 128)) {
+            compute_type = GGML_TYPE_F32;
+        }
    }
    if (dst->op_params[0] == GGML_PREC_F32) {
        compute_type = GGML_TYPE_F32;
@@ -110,15 +110,21 @@ if (GGML_SYCL_SUPPORT_LEVEL_ZERO_API)
    # Link against Level Zero loader for direct device memory allocation.
    # Avoids sycl::malloc_device triggering DMA-buf/TTM system RAM staging
    # in the xe kernel driver during multi-GPU inference.
-    find_path(LEVEL_ZERO_INCLUDE_DIR level_zero/ze_api.h HINTS ${ONEAPI_ROOT}/include ${LEVEL_ZERO_V1_SDK_PATH}/include)
+    find_path(LEVEL_ZERO_DEV_INCLUDE_DIR level_zero/ze_api.h HINTS ${ONEAPI_ROOT}/include ${LEVEL_ZERO_V1_SDK_PATH}/include)
    find_library(ZE_LOADER_LIB ze_loader HINTS ${ONEAPI_ROOT}/lib ${LEVEL_ZERO_V1_SDK_LIB_PATH} ENV LD_LIBRARY_PATH)
-    if(ZE_LOADER_LIB AND LEVEL_ZERO_INCLUDE_DIR)
+    if(ZE_LOADER_LIB AND LEVEL_ZERO_DEV_INCLUDE_DIR)
        target_link_libraries(ggml-sycl PRIVATE ${ZE_LOADER_LIB})
        target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_SUPPORT_LEVEL_ZERO_API)
        message(STATUS "Level Zero loader found: ${ZE_LOADER_LIB}")
-        message(STATUS "Level Zero headers found: ${LEVEL_ZERO_INCLUDE_DIR}")
+        message(STATUS "Level Zero development headers found: ${LEVEL_ZERO_DEV_INCLUDE_DIR}")
    else()
-        message(WARNING "Level Zero loader or headers not found, Level Zero support disabled")
+        message(WARNING "Level Zero loader or development headers not found, "
+                        "Level Zero API support disabled. "
+                        "Please install the Level Zero SDK/development package "
+                        "to support Level Zero API features. "
+                        "Level Zero API is not mandatory for SYCL backend, "
+                        "but it is required by the special features for better "
+                        "function & performance on Intel GPUs.")
    endif()
 endif()

@@ -44,6 +44,7 @@
 #include "ssm_conv.hpp"
 #include "softmax.hpp"
 #include "topk-moe.hpp"
+#include "topk-radix.hpp"
 #include "tsembd.hpp"
 #include "upscale.hpp"
 #include "wkv.hpp"
@@ -17,6 +17,7 @@
 #include <cstdio>

 extern int g_ggml_sycl_debug;
+extern int g_ggml_sycl_dev_debug;

 #if defined(__clang__) && __has_builtin(__builtin_expect)
 // Hint the optimizer to pipeline the more likely following instruction in branches
@@ -33,4 +34,10 @@ extern int g_ggml_sycl_debug;
            fprintf(stderr, __VA_ARGS__); \
    } while (0)

+#define GGML_SYCL_DEV_DEBUG(...)          \
+    do {                                  \
+        if (UNLIKELY(g_ggml_sycl_dev_debug))  \
+            fprintf(stderr, __VA_ARGS__); \
+    } while (0)
+
 #endif  // GGML_SYCL_BASE_HPP
@@ -401,29 +401,10 @@ struct ggml_backend_sycl_context {
    dnnl::stream stream_dnnl() {
        return stream_dnnl(device, 0);
    }
-    dnnl::memory get_scratchpad_mem(const dnnl::memory::desc & scratchpad_md,
-                                    const dnnl::engine & eng, const queue_ptr q) {
-        ggml_sycl_pool_alloc<uint8_t> * pool;
-        auto it = scratchpad_map.find(q);
-        if (it == scratchpad_map.end()) {
-            scratchpad_map[q] = std::make_unique<ggml_sycl_pool_alloc<uint8_t>>(this->pool());
-            pool = scratchpad_map[q].get();
-        } else {
-            pool = it->second.get();
-        }
-
-        size_t scratchpad_size = scratchpad_md.get_size();
-        if (scratchpad_size > pool->actual_size) {
-            pool->realloc(scratchpad_size);
-        }
-        void * mem_ptr = pool->get();
-        return dnnl::memory(scratchpad_md, eng, mem_ptr);
-    }
 #endif

    // pool
    std::unique_ptr<ggml_sycl_pool> pools[GGML_SYCL_MAX_DEVICES];
-    std::unordered_map<sycl::queue *, std::unique_ptr<ggml_sycl_pool_alloc<uint8_t>>> scratchpad_map;

    std::unique_ptr<ggml_sycl_fattn_kv_buffers> fattn_bufs[GGML_SYCL_MAX_DEVICES];

@@ -66,8 +66,10 @@ public:
        auto matmul_pd = dnnl::matmul::primitive_desc(eng, a_in_md, b_in_md, c_md, primitive_attr);
        auto c_mem = dnnl::memory(matmul_pd.dst_desc(), eng, c);

-        auto scratchpad_md = matmul_pd.scratchpad_desc();
-        auto scratchpad_mem = ctx.get_scratchpad_mem(scratchpad_md, eng, q);
+        const auto scratchpad_md = matmul_pd.scratchpad_desc();
+        ggml_sycl_pool_alloc<uint8_t> scratchpad(ctx.pool());
+        void * scratchpad_ptr = scratchpad_md.get_size() > 0 ? scratchpad.alloc(scratchpad_md.get_size()) : nullptr;
+        auto scratchpad_mem = dnnl::memory(scratchpad_md, eng, scratchpad_ptr);

        auto matmul_prim = dnnl::matmul(matmul_pd);

@@ -91,6 +91,7 @@

 static bool g_sycl_loaded = false;
 int g_ggml_sycl_debug = 0;
+int g_ggml_sycl_dev_debug = 0;
 int g_ggml_sycl_enable_optimize = 1;
 int g_ggml_sycl_enable_graph = 0;
 int g_ggml_sycl_enable_dnn = 1;
@@ -113,8 +114,8 @@ int g_ggml_sycl_enable_host_pinned_mem = 1;
 int g_ggml_sycl_host_pinned_mem_2g = 0;
 int g_ggml_sycl_get_mem_api = MEMORY_API_TYPE_LEVEL_ZERO;

-
 static ggml_sycl_device_info ggml_sycl_init() {
+    GGML_SYCL_DEBUG("[SYCL] call ggml_sycl_init\n");
    ggml_sycl_device_info info = {};

    // Do not hard crash when there exists no SYCL devices.
@@ -205,12 +206,9 @@ static ggml_sycl_device_info ggml_sycl_init() {
    }

 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
-    // Large buffers can be allocated before ggml_check_sycl() initializes other
-    // g_ggml_sycl_enable_* globals, so initialize this one as early as we can.
+    //update g_ggml_sycl_use_level_zero_api according to the device support
    g_ggml_sycl_use_level_zero_api =
-        info.ext_oneapi_level_zero && ggml_sycl_get_env("GGML_SYCL_USE_LEVEL_ZERO_API", 1);
-#else
-    g_ggml_sycl_use_level_zero_api = 0;
+        info.ext_oneapi_level_zero && g_ggml_sycl_use_level_zero_api;
 #endif

    return info;
@@ -314,23 +312,40 @@ static const char* dev2dev_int2str(int dev2dev) {
 * It's the first internal function to be called by them in SYCL backend.
 * This function is used to do initialize work for the SYCL backend and set the global variables.
 */
+#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+static ze_result_t init_zes() {
+    ze_result_t res = zesInit(0);
+    if (res != ZE_RESULT_SUCCESS) {
+        GGML_SYCL_DEBUG("Warning: [%s] zesInit failed with code %d. Sysman free-memory query be unavailable.\n",
+                            __func__, (int) res);
+    }
+    return res;
+}
+
+ze_result_t get_zes_init_res() {
+    static ze_result_t zes_init_res = init_zes();
+    GGML_SYCL_DEBUG("[SYCL] call %s: zesInit result: %d\n", __func__, (int) zes_init_res);
+    return zes_init_res;
+}
+#endif
+
 void initialize_sycl_begining() {
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
-    ze_result_t zes_init = zesInit(0);
-    if (zes_init != ZE_RESULT_SUCCESS) {
-        std::cerr << "Warning: zesInit failed [ggml_check_sycl] with code " << static_cast<int>(zes_init)
-                            << ". Sysman free-memory query may be unavailable.\n";
-    }
+    //must be called in initialization stage, before any other Level Zero API calls
+    GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__);
+    get_zes_init_res();
 #endif
 }

 static void ggml_check_sycl() try {
+    GGML_SYCL_DEBUG("[SYCL] ggml_check_sycl()\n");
    static bool initialized = false;

    if (!initialized) {
        initialize_sycl_begining();

        g_ggml_sycl_debug = ggml_sycl_get_env("GGML_SYCL_DEBUG", 0);
+        g_ggml_sycl_dev_debug = ggml_sycl_get_env("GGML_SYCL_DEV_DEBUG", 0);
        g_ggml_sycl_enable_optimize = ggml_sycl_get_env("GGML_SYCL_ENABLE_OPT", 1);
        g_ggml_sycl_enable_graph = ggml_sycl_get_env("GGML_SYCL_ENABLE_GRAPH", 0);
        g_ggml_sycl_enable_dnn = ggml_sycl_get_env("GGML_SYCL_ENABLE_DNN", 1);
@@ -344,9 +359,13 @@ static void ggml_check_sycl() try {
        g_ggml_sycl_enable_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_ESIMD", 1);
        g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);

+#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+        g_ggml_sycl_use_level_zero_api = ggml_sycl_get_env("GGML_SYCL_USE_LEVEL_ZERO_API", 1);
+#else
+        g_ggml_sycl_use_level_zero_api = 0;
+#endif
        g_ggml_sycl_dev2dev_memcpy = ggml_sycl_get_env("GGML_SYCL_DEV2DEV_MEMCPY", DEV2DEV_MEMCPY_SYCL);
        g_ggml_sycl_get_mem_api = ggml_sycl_get_env("GGML_SYCL_GET_MEM_API", MEMORY_API_TYPE_LEVEL_ZERO);
-
        if (g_ggml_sycl_use_level_zero_api == 0) {
            g_ggml_sycl_dev2dev_memcpy = DEV2DEV_MEMCPY_SYCL;
            g_ggml_sycl_get_mem_api = MEMORY_API_TYPE_SYCL;
@@ -405,6 +424,7 @@ static void ggml_check_sycl() try {

        GGML_LOG_INFO("Running with Environment Variables:\n");
        GGML_LOG_INFO("  GGML_SYCL_DEBUG: %d\n", g_ggml_sycl_debug);
+        GGML_LOG_INFO("  GGML_SYCL_DEV_DEBUG: %d\n", g_ggml_sycl_dev_debug);

 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
        GGML_LOG_INFO("  GGML_SYCL_DEV2DEV_MEMCPY: %d (%s)\n", g_ggml_sycl_dev2dev_memcpy, dev2dev_int2str(g_ggml_sycl_dev2dev_memcpy));
@@ -945,6 +965,7 @@ inline void * aligned_malloc_host(size_t alignment, size_t size) {
 static ggml_backend_buffer_t
 ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft,
                                           size_t size) try {
+    GGML_SYCL_DEBUG("[SYCL] call %s: size=%zu\n", __func__, size);
    ggml_check_sycl();

    ggml_backend_sycl_buffer_type_context * buft_ctx = (ggml_backend_sycl_buffer_type_context *)buft->context;
@@ -1464,10 +1485,11 @@ static ggml_backend_buffer_type_i ggml_backend_sycl_split_buffer_type_interface
 };

 ggml_backend_buffer_type_t ggml_backend_sycl_split_buffer_type(const float * tensor_split) {
+    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_split_buffer_type\n");
+
    static std::mutex mutex;
    std::lock_guard<std::mutex> lock(mutex);

-    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_split_buffer_type\n");
    ggml_check_sycl();
    // FIXME: this is not thread safe
    static std::map<std::array<float, GGML_SYCL_MAX_DEVICES>, struct ggml_backend_buffer_type> buft_map;
@@ -1520,6 +1542,7 @@ static const char * ggml_backend_sycl_host_buffer_type_name(ggml_backend_buffer_

 //host pinned memory
 static void * ggml_backend_sycl_host_malloc(size_t size) {
+    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_host_malloc\n");
    void * ptr = nullptr;
    try {
        ggml_check_sycl();
@@ -3098,10 +3121,14 @@ static void ggml_sycl_op_top_k(ggml_backend_sycl_context & ctx, ggml_tensor * ds
    const int64_t ncols = src0->ne[0];
    const int64_t nrows = ggml_nrows(src0);

-    GGML_ASSERT(k > 0 && k <= 32);
+    GGML_ASSERT(k > 0);
    GGML_ASSERT(k <= ncols);

-    top_k_f32_sycl(ctx, src0_dd, dst_dd, ncols, nrows, k, main_stream);
+    if (k <= SYCL_TOP_K_SCAN_MERGE_MAX_K) {
+        top_k_f32_sycl(ctx, src0_dd, dst_dd, ncols, nrows, k, main_stream);
+    } else {
+        ggml_sycl_top_k_radix(ctx, src0_dd, dst_dd, ncols, nrows, k, main_stream);
+    }
 }

 inline void ggml_sycl_op_argmax(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
@@ -5341,8 +5368,8 @@ catch (sycl::exception const &exc) {
 }

 static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct ggml_tensor * dst) try {
+    GGML_SYCL_DEBUG("[SYCL] ggml_sycl_compute_forward: dst=%s, op=%s\n", dst->name, ggml_op_name(dst->op));
    if (!g_sycl_loaded) return false;
-    initialize_sycl_begining();

    if (dst->src[0] != nullptr && ggml_backend_buffer_is_sycl_split(dst->src[0]->buffer)) {
        ggml_sycl_set_peer_access(dst->src[1]->ne[1], ctx.device);
@@ -5725,11 +5752,27 @@ catch (sycl::exception const &exc) {
  std::exit(1);
 }

+bool sycl_get_mem_info(int device, size_t * free, size_t * total) {
+    GGML_SYCL_DEBUG("[SYCL] [%s] g_ggml_sycl_get_mem_api=%d\n",
+        __func__, g_ggml_sycl_get_mem_api);
+
+    MemoryAPIType mem_api_type = MemoryAPIType::MEMORY_API_TYPE_SYCL;
+
+#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+    mem_api_type = get_zes_init_res() == ZE_RESULT_SUCCESS ?
+        (MemoryAPIType) g_ggml_sycl_get_mem_api : MemoryAPIType::MEMORY_API_TYPE_SYCL;
+#else
+    mem_api_type = MemoryAPIType::MEMORY_API_TYPE_SYCL;
+#endif
+    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device),
+        *free, *total, mem_api_type);
+    GGML_SYCL_DEBUG("[SYCL] [%s] total = %zu free = %zu\n", __func__, *total, *free);
+    return res;
+}
+
 void ggml_backend_sycl_get_device_memory(int device, size_t * free, size_t * total) try {
    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_get_device_memory\n");
-    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(device), *free, *total,
-                               (MemoryAPIType) g_ggml_sycl_get_mem_api);
-    if (!res) {
+    if (!sycl_get_mem_info(device, free, total)) {
        GGML_ABORT("[%s] failed to get device memory size", __func__);
    }
    ggml_sycl_memtrace_report_device("device memory query", device, *free, *total);
@@ -6177,12 +6220,12 @@ static const char * ggml_backend_sycl_device_get_description(ggml_backend_dev_t
 }

 static void ggml_backend_sycl_device_get_memory(ggml_backend_dev_t dev, size_t * free, size_t * total) {
+    GGML_SYCL_DEBUG("[SYCL] call %s\n", __func__);
    ggml_backend_sycl_device_context * ctx = (ggml_backend_sycl_device_context *) dev->context;
-    bool res = get_memory_size(dpct::dev_mgr::instance().get_device(ctx->device), *free, *total,
-                               (MemoryAPIType) g_ggml_sycl_get_mem_api);
-    if (!res) {
+    if (!sycl_get_mem_info(ctx->device, free, total)) {
        GGML_ABORT("[%s] failed to get device memory size", __func__);
    }
+    GGML_SYCL_DEBUG("[SYCL] call %s total %zu free %zu\n", __func__, *total, *free);
    ggml_sycl_memtrace_report_device("device memory query (dev)", ctx->device, *free, *total);
 }

@@ -6605,7 +6648,7 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
                op->type == GGML_TYPE_I32 &&
                src0->type == GGML_TYPE_F32 &&
                ggml_is_contiguous(src0) &&
-                k > 0 && k <= 32;
+                k > 0 && k <= src0->ne[0];
        }
        case GGML_OP_POOL_2D:
        case GGML_OP_POOL_1D:
@@ -7061,6 +7104,7 @@ static const ggml_backend_reg_i ggml_backend_sycl_reg_interface = {
 // backend registry

 ggml_backend_reg_t ggml_backend_sycl_reg() {
+    GGML_SYCL_DEBUG("[SYCL] call ggml_backend_sycl_reg\n");
    static ggml_backend_reg reg;
    static bool initialized = false;

@@ -7068,7 +7112,7 @@ ggml_backend_reg_t ggml_backend_sycl_reg() {
        static std::mutex mutex;
        std::lock_guard<std::mutex> lock(mutex);
        if (!initialized) {
-            initialize_sycl_begining();
+            ggml_check_sycl();
            ggml_backend_sycl_reg_context * ctx = new ggml_backend_sycl_reg_context;
            const int min_batch_size = getenv("GGML_OP_OFFLOAD_MIN_BATCH") ? atoi(getenv("GGML_OP_OFFLOAD_MIN_BATCH")) : 32;

@@ -6,13 +6,13 @@
 #include <level_zero/zes_api.h>
 #endif

-#include "base.hpp"
-#include "mem.hpp"
-
 #include <cstdint>
 #include <iostream>
 #include <vector>

+#include "base.hpp"
+#include "mem.hpp"
+
 const char * mem_api_int2str(int mem_api) {
    if (mem_api == MEMORY_API_TYPE_SYCL) {
        return "SYCL API";
@@ -24,7 +24,12 @@ const char * mem_api_int2str(int mem_api) {
 }

 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
+/*
+* Depend on to call zesInit(0) before any other Level Zero API calls, otherwise the Level Zero API calls may fail.
+*/
 bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
+    GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using Level Zero API.\n", __func__);
+
    free_bytes  = 0;
    total_bytes = 0;

@@ -37,41 +42,28 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot
 #endif

    try {
-        ze_result_t zes_init = zesInit(0);
-        if (zes_init != ZE_RESULT_SUCCESS) {
-            std::cerr << "Warning: zesInit failed with code " << static_cast<int>(zes_init)
-                      << ". Sysman free-memory query may be unavailable.\n";
-        }

        if (dev.get_platform().get_backend() != kL0Backend) {
-            GGML_SYCL_DEBUG("Device backend is not Level Zero; falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Device backend is not Level Zero.\n");
            return false;
        }

        ze_device_handle_t ze_dev = sycl::get_native<kL0Backend>(dev);
        if (ze_dev == nullptr) {
-            GGML_SYCL_DEBUG("Level Zero device handle is null; falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Level Zero device handle is null.\n");
            return false;
        }

        ze_result_t r = zesDeviceEnumMemoryModules(ze_dev, &module_count, nullptr);
        if (r != ZE_RESULT_SUCCESS || module_count == 0) {
-            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules. Falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules.\n");
            return false;
        }

        std::vector<zes_mem_handle_t> modules(module_count);
        r = zesDeviceEnumMemoryModules(ze_dev, &module_count, modules.data());
        if (r != ZE_RESULT_SUCCESS || module_count == 0) {
-            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules. Falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Failed to enumerate Level Zero memory modules.\n");
            return false;
        }

@@ -90,73 +82,70 @@ bool query_free_memory_by_ze(sycl::device dev, size_t & free_bytes, size_t & tot
        }

        if (total_bytes == 0) {
-            GGML_SYCL_DEBUG("Level Zero memory query returned zero total bytes. Falling back to SYCL memory query.\n");
-            total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-            free_bytes  = total_bytes;
+            GGML_SYCL_DEBUG("Level Zero memory query returned zero total bytes.\n");
            return false;
        }
-        return true;
+        return total_bytes >= free_bytes;
+
    } catch (const sycl::exception & e) {
        GGML_SYCL_DEBUG("Level Zero memory query failed: %s\n", e.what());
-        total_bytes = dev.get_info<sycl::info::device::global_mem_size>();
-        free_bytes  = total_bytes;
        return false;
    }
 }
 #endif

 bool get_memory_size_by_sycl_api(sycl::device dev, size_t & free_bytes, size_t & total_bytes) {
-    GGML_SYCL_DEBUG("[%s]Querying free memory using SYCL API.\n", __func__);
+    GGML_SYCL_DEBUG("[SYCL] call %s: Querying free memory using SYCL API.\n", __func__);
    total_bytes = dev.get_info<sycl::info::device::global_mem_size>();

 #if (defined(__SYCL_COMPILER_VERSION) && __SYCL_COMPILER_VERSION >= 20221105)
    if (dev.has(sycl::aspect::ext_intel_free_memory)) {
        try {
-            GGML_SYCL_DEBUG("Querying free memory using SYCL aspect::ext_intel_free_memory.");
+            GGML_SYCL_DEBUG("Querying free memory using SYCL aspect::ext_intel_free_memory.\n");
            free_bytes = dev.get_info<sycl::ext::intel::info::device::free_memory>();
            return true;
        } catch (const sycl::exception &) {
            GGML_SYCL_DEBUG(
-                "Failed to query free memory using SYCL aspect::ext_intel_free_memory. Using total memory as free "
-                "memory.");
-            free_bytes = total_bytes;
+                "Failed to query free memory using SYCL aspect::ext_intel_free_memory.\n");
            return false;
        }
    } else {
        GGML_SYCL_DEBUG(
-            "Device does not support SYCL aspect::ext_intel_free_memory. Using total memory as free memory.");
-        free_bytes = total_bytes;
+            "Device does not support SYCL aspect::ext_intel_free_memory.\n");
    }
 #else
-    GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105. Using total memory as free memory.");
-    free_bytes = total_bytes;
+    GGML_SYCL_DEBUG("SYCL Compiler version is older than 20221105.\n");
 #endif
-    return true;
+    return false;
 }

 bool get_memory_size(sycl::device dev, size_t & free_bytes, size_t & total_bytes, MemoryAPIType api_type) {
-    const auto name       = dev.get_info<sycl::info::device::name>();
-    const auto vendor     = dev.get_info<sycl::info::device::vendor>();
-    const auto global_mem = dev.get_info<sycl::info::device::global_mem_size>();

-    GGML_SYCL_DEBUG("[%s]GPU Name:          %s\n", __func__, name.c_str());
-    GGML_SYCL_DEBUG("[%s]GPU Vendor:        %s\n", __func__, vendor.c_str());
-    GGML_SYCL_DEBUG("[%s]GPU Global Memory: %zu bytes\n", __func__, static_cast<size_t>(global_mem));
+    GGML_SYCL_DEBUG("[%s]GPU Name:          %s\n", __func__,
+        dev.get_info<sycl::info::device::name>().c_str());
+    GGML_SYCL_DEBUG("[%s]GPU Vendor:        %s\n", __func__,
+        dev.get_info<sycl::info::device::vendor>().c_str());

    if (api_type == MEMORY_API_TYPE_LEVEL_ZERO) {
 #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
-        GGML_SYCL_DEBUG("[%s]Querying free memory using Level Zero API.\n", __func__);
-        if (!query_free_memory_by_ze(dev, free_bytes, total_bytes)) {
-            //fallback to SYCL API if Level Zero API fails
-            GGML_SYCL_DEBUG("[%s]Falling back to SYCL API for memory query.\n", __func__);
-            return get_memory_size_by_sycl_api(dev, free_bytes, total_bytes);
+        GGML_SYCL_DEBUG("[%s] Querying free memory using Level Zero API.\n", __func__);
+        if (query_free_memory_by_ze(dev, free_bytes, total_bytes)) {
+            return true;
        }
-        return true;
-#else
-        GGML_SYCL_DEBUG("[%s]Level Zero API support is not enabled. Please enable it to use this feature.\n", __func__);
-        return false;
+        //fallback to SYCL API if Level Zero API fails
+        GGML_SYCL_DEBUG("[%s] Falling back to SYCL API for memory query.\n", __func__);
 #endif
-    } else {  //MEMORY_API_TYPE_SYCL
-        return get_memory_size_by_sycl_api(dev, free_bytes, total_bytes);
    }
+
+    //MEMORY_API_TYPE_SYCL
+    if(get_memory_size_by_sycl_api(dev, free_bytes, total_bytes)){
+        return true;
+    }
+
+    //Todo, fallback to other methods to get free memory size, such as using OS-specific APIs (e.g., /proc/meminfo on Linux, GlobalMemoryStatusEx on Windows, etc.)
+    GGML_SYCL_DEBUG(
+        "[%s] Can't get free mem size by Level Zero and SYCL API. Using total memory as free memory.\n", __func__);
+    free_bytes = total_bytes;
+
+    return true;
 }
@@ -0,0 +1,531 @@
+#include "topk-radix.hpp"
+
+#include "common.hpp"
+
+#include <algorithm>
+
+// Large-k top-k by radix select on an order-preserving unsigned key.
+//
+// The k-th largest key of a row is found by four most-significant-first passes over its
+// 8-bit digits: histogram the digit over the candidate set, walk the buckets from the
+// top, and recurse into the bucket where the running count reaches what is still
+// needed. Everything strictly above that bucket is in the top-k. A final pass emits
+// every column whose key beats the pivot, then exactly as many pivot-equal columns as
+// are still missing, so duplicate keys yield exactly k distinct indices.
+//
+// SLM holds only the histogram, so unlike the scan-merge kernels the cost does not grow
+// with k. One work-group owns a row and runs every pass, so a top-k is one launch and
+// needs no pool scratch. The row is re-read once per pass rather than compacted, which
+// keeps the candidate set implicit: (key & mask) == prefix.
+//
+// The output is the set of winning indices in no particular order, which is what the
+// reference op provides (it swaps its first two outputs to say so) and what
+// test-backend-ops compares.
+
+static constexpr int SYCL_TOP_K_RADIX_BITS       = 8;
+static constexpr int SYCL_TOP_K_RADIX_BUCKETS    = 1 << SYCL_TOP_K_RADIX_BITS;
+// Private histogram copies, interleaved per bucket so neighbouring lanes hit
+// neighbouring banks. Lanes of one instruction spread over the copies, which is what
+// bounds the atomic serialisation on tie-heavy rows.
+static constexpr int SYCL_TOP_K_RADIX_HIST_COPIES = 8;
+static constexpr int SYCL_TOP_K_RADIX_HIST_SIZE   = SYCL_TOP_K_RADIX_BUCKETS * SYCL_TOP_K_RADIX_HIST_COPIES;
+// Past the histogram: pivot digit, pivot bucket count, remaining need, then the two
+// emit counters.
+static constexpr int SYCL_TOP_K_RADIX_SLM_WORDS   = SYCL_TOP_K_RADIX_HIST_SIZE + 5;
+
+// Larger float <=> larger key. The reference comparator is a plain float '>', under which
+// -0.0 and +0.0 tie, so -0.0 is folded onto +0.0 first. NaN has no defined order in the
+// reference (its comparator is not a strict weak order on NaN); here a positive NaN keys
+// above +inf and a negative NaN below -inf, which at least makes the result deterministic.
+static inline uint32_t top_k_radix_key(float f) {
+    uint32_t u = sycl::bit_cast<uint32_t>(f);
+    if (u == 0x80000000u) {
+        u = 0u;
+    }
+    return (u & 0x80000000u) ? ~u : (u | 0x80000000u);
+}
+
+static void top_k_radix_select_f32(
+    const float *   src,
+    int32_t *       dst_idx,
+    const int       ncols,
+    const int       k,
+    uint32_t *      slm,
+    const sycl::nd_item<1> & item_ct1
+) {
+    using local_atomic = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed,
+                                          sycl::memory_scope::work_group,
+                                          sycl::access::address_space::local_space>;
+
+    const int tid        = item_ct1.get_local_id(0);
+    const int block_size = item_ct1.get_local_range(0);
+
+    uint32_t * hist     = slm;
+    uint32_t * s_digit  = slm + SYCL_TOP_K_RADIX_HIST_SIZE;
+    uint32_t * s_bucket = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 1;
+    uint32_t * s_need   = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 2;
+    uint32_t * s_cnt_gt = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 3;
+    uint32_t * s_cnt_eq = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 4;
+
+    if (tid == 0) {
+        *s_cnt_gt = 0;
+        *s_cnt_eq = 0;
+    }
+
+    const int copy = tid & (SYCL_TOP_K_RADIX_HIST_COPIES - 1);
+
+    uint32_t prefix = 0;   // digits fixed so far, in place
+    uint32_t mask   = 0;   // which bits of prefix are fixed
+    uint32_t need   = (uint32_t) k;
+
+    for (int shift = 32 - SYCL_TOP_K_RADIX_BITS; shift >= 0; shift -= SYCL_TOP_K_RADIX_BITS) {
+        for (int i = tid; i < SYCL_TOP_K_RADIX_HIST_SIZE; i += block_size) {
+            hist[i] = 0;
+        }
+        item_ct1.barrier(sycl::access::fence_space::local_space);
+
+        for (int col = tid; col < ncols; col += block_size) {
+            const uint32_t key = top_k_radix_key(src[col]);
+            if ((key & mask) == prefix) {
+                const uint32_t bucket = (key >> shift) & (SYCL_TOP_K_RADIX_BUCKETS - 1);
+                local_atomic(hist[bucket * SYCL_TOP_K_RADIX_HIST_COPIES + copy]).fetch_add(1u);
+            }
+        }
+        item_ct1.barrier(sycl::access::fence_space::local_space);
+
+        // Lane t takes bucket 255 - t, so an inclusive scan over lanes counts from the top
+        // bucket downward. The pivot is the unique bucket whose cumulative count first
+        // reaches need; the previous cumulative count is what the higher buckets contribute.
+        uint32_t cnt = 0;
+        if (tid < SYCL_TOP_K_RADIX_BUCKETS) {
+            const uint32_t * h = hist + (SYCL_TOP_K_RADIX_BUCKETS - 1 - tid) * SYCL_TOP_K_RADIX_HIST_COPIES;
+            for (int c = 0; c < SYCL_TOP_K_RADIX_HIST_COPIES; c++) {
+                cnt += h[c];
+            }
+        }
+        const uint32_t incl = sycl::inclusive_scan_over_group(item_ct1.get_group(), cnt, sycl::plus<uint32_t>());
+
+        if (tid < SYCL_TOP_K_RADIX_BUCKETS && incl >= need && incl - cnt < need) {
+            *s_digit = (uint32_t) (SYCL_TOP_K_RADIX_BUCKETS - 1 - tid);
+            *s_bucket = cnt;
+            *s_need   = need - (incl - cnt);
+        }
+        item_ct1.barrier(sycl::access::fence_space::local_space);
+
+        const uint32_t digit      = *s_digit;
+        const uint32_t bucket_cnt = *s_bucket;
+        need   = *s_need;
+        prefix |= digit << shift;
+        mask   |= (uint32_t) (SYCL_TOP_K_RADIX_BUCKETS - 1) << shift;
+
+        // Every candidate in the pivot bucket is wanted: the remaining digits cannot
+        // change the answer, and the masked emit below is exact as it stands.
+        if (bucket_cnt == need) {
+            break;
+        }
+        // The next pass rewrites hist and s_*; the reads above must land first.
+        item_ct1.barrier(sycl::access::fence_space::local_space);
+    }
+
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    // Exactly k - need columns have (key & mask) > prefix; the first need of the pivot-equal
+    // columns fill the tail. Both counters live in SLM since the whole row is this group.
+    const uint32_t base_eq = (uint32_t) k - need;
+
+    for (int col = tid; col < ncols; col += block_size) {
+        const uint32_t kp = top_k_radix_key(src[col]) & mask;
+        if (kp > prefix) {
+            const uint32_t pos = local_atomic(*s_cnt_gt).fetch_add(1u);
+            dst_idx[pos] = col;
+        } else if (kp == prefix) {
+            const uint32_t pos = local_atomic(*s_cnt_eq).fetch_add(1u);
+            if (pos < need) {
+                dst_idx[base_eq + pos] = col;
+            }
+        }
+    }
+}
+
+static void top_k_radix_f32_sycl(
+    ggml_backend_sycl_context & ctx,
+    const float * src,
+    int32_t * dst_indices,
+    const int64_t ncols,
+    const int64_t nrows,
+    const int k,
+    dpct::queue_ptr main_stream
+) {
+    GGML_ASSERT(ncols <= INT32_MAX);
+
+    // One group per row; every pass is a strided sweep of the row, so lanes in flight is the
+    // only lever, and the device's own limit is the answer -- there is nothing here that
+    // wants a smaller group. Must still cover the 256 buckets for the scan step.
+    const int block_size = ggml_sycl_info().max_work_group_sizes[ctx.device];
+    GGML_ASSERT(block_size >= SYCL_TOP_K_RADIX_BUCKETS);
+
+    const sycl::range<1> block_dims(block_size);
+    const sycl::range<1> grid_dims(nrows);
+
+    main_stream->submit([&](sycl::handler &cgh) {
+        sycl::local_accessor<uint32_t, 1> slm(sycl::range<1>(SYCL_TOP_K_RADIX_SLM_WORDS), cgh);
+
+        cgh.parallel_for(
+            sycl::nd_range<1>(grid_dims * block_dims, block_dims),
+            [=](sycl::nd_item<1> item_ct1) {
+                const int row = item_ct1.get_group(0);
+
+                top_k_radix_select_f32(
+                    src + (int64_t) row * ncols, dst_indices + (int64_t) row * k,
+                    (int) ncols, k,
+                    slm.get_multi_ptr<sycl::access::decorated::no>().get(),
+                    item_ct1);
+            });
+    });
+}
+
+// One work-group owns a whole row above, which leaves the device idle whenever a graph
+// has fewer rows than it has cores -- the common case at batch size 1, where the
+// sparse-attention indexer and the backend sampler both top-k a single row. The kernels
+// below spread one row over several groups instead.
+//
+// A digit pass now needs the whole row's histogram before any group can pick the pivot,
+// so the per-pass state moves to global memory and the passes become separate launches:
+// a work-group barrier no longer spans the row. Each group still accumulates into SLM
+// and contributes 256 global atomics at the end, so global traffic is per-group, not
+// per-element. The last group to finish a pass (the one whose fetch_add returns G - 1)
+// does the scan for the row and clears the histogram for the next pass, which keeps the
+// launch count at one per digit rather than two.
+//
+// Running all four digits unconditionally costs nothing in correctness: once a bucket
+// holds exactly the elements still needed, later digits only extend the prefix, and the
+// count of columns above that longer prefix grows by exactly as much as `need` shrinks.
+// The emit below therefore stays exact whatever pass the answer settled on.
+
+static constexpr int SYCL_TOP_K_RADIX_ROW_DONE   = SYCL_TOP_K_RADIX_BUCKETS + 0;
+static constexpr int SYCL_TOP_K_RADIX_ROW_PREFIX = SYCL_TOP_K_RADIX_BUCKETS + 1;
+static constexpr int SYCL_TOP_K_RADIX_ROW_MASK   = SYCL_TOP_K_RADIX_BUCKETS + 2;
+static constexpr int SYCL_TOP_K_RADIX_ROW_NEED   = SYCL_TOP_K_RADIX_BUCKETS + 3;
+static constexpr int SYCL_TOP_K_RADIX_ROW_CNT_GT = SYCL_TOP_K_RADIX_BUCKETS + 4;
+static constexpr int SYCL_TOP_K_RADIX_ROW_CNT_EQ = SYCL_TOP_K_RADIX_BUCKETS + 5;
+static constexpr int SYCL_TOP_K_RADIX_ROW_WORDS  = SYCL_TOP_K_RADIX_BUCKETS + 6;
+
+// How wide the split goes is a property of the device, not of the model: enough groups to
+// cover the cores, and no more. Past that the extra groups add histogram traffic without
+// adding bandwidth (measured on this device: 20 and 40 groups tie, 60 and 160 lose).
+//
+// nsm is max_compute_units / 16, i.e. it counts an Xe core as 16 EUs. That is a core's
+// width on Xe-HPG, but an Xe2 core is 8 XVEs wide, so on Battlemage the field reads half
+// the cores actually present (10 for a 20-core B60). The measured curve is flat from one
+// group per core to two and only falls off at three, so a factor of two covers the device
+// on Xe2 and lands in the flat region on Xe-HPG. It is the one number here that a correct
+// core count would remove; it was tuned on Xe2 and has not been measured on Xe-HPG.
+static constexpr int SYCL_TOP_K_RADIX_GROUPS_PER_NSM = 2;
+// Splitting trades one kernel for five. Below the width at which the single-group kernel
+// runs longer than those four extra launches, it wins on its own; measured break-even on
+// this device sits just under 64K columns.
+static constexpr int SYCL_TOP_K_RADIX_MIN_SPLIT_COLS = 65536;
+// A partition thinner than this cannot keep a group's sweep busy.
+static constexpr int SYCL_TOP_K_RADIX_MIN_PART_COLS  = 4096;
+
+static int top_k_radix_split_groups(const int device, const int64_t ncols, const int64_t nrows) {
+    const int64_t target = (int64_t) SYCL_TOP_K_RADIX_GROUPS_PER_NSM * ggml_sycl_info().devices[device].nsm;
+
+    // One group per row already, so a graph with rows enough to cover the device gains
+    // nothing from splitting and would only pay the extra launches.
+    if (ncols < SYCL_TOP_K_RADIX_MIN_SPLIT_COLS || nrows >= target) {
+        return 1;
+    }
+
+    const int64_t by_rows = target / nrows;   // floor: never overshoot a row that is nearly covered
+    const int64_t by_cols = ncols / SYCL_TOP_K_RADIX_MIN_PART_COLS;
+
+    return (int) std::max<int64_t>(1, std::min(by_rows, by_cols));
+}
+
+using top_k_radix_gatomic = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed,
+                                             sycl::memory_scope::device,
+                                             sycl::access::address_space::global_space>;
+
+static void top_k_radix_split_pass_f32(
+    const float *   src,
+    uint32_t *      state,
+    const int       ncols,
+    const int       k,
+    const int       shift,
+    const bool      first,
+    const int       part,
+    const int       nparts,
+    uint32_t *      slm,
+    const sycl::nd_item<1> & item_ct1
+) {
+    using local_atomic = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed,
+                                          sycl::memory_scope::work_group,
+                                          sycl::access::address_space::local_space>;
+
+    const int tid        = item_ct1.get_local_id(0);
+    const int block_size = item_ct1.get_local_range(0);
+
+    uint32_t * hist   = slm;
+    uint32_t * s_last = slm + SYCL_TOP_K_RADIX_HIST_SIZE;
+    uint32_t * s_row  = slm + SYCL_TOP_K_RADIX_HIST_SIZE + 1;   // prefix, mask, need
+
+    // The previous launch is the barrier that publishes these, so a plain load is enough.
+    // One lane reads them and the group takes them from SLM: a device-scope atomic load
+    // is uncached here, and having every work-item issue three of them off the same
+    // address costs more than the whole sweep below.
+    if (tid == 0) {
+        s_row[0] = first ? 0u : state[SYCL_TOP_K_RADIX_ROW_PREFIX];
+        s_row[1] = first ? 0u : state[SYCL_TOP_K_RADIX_ROW_MASK];
+        s_row[2] = first ? (uint32_t) k : state[SYCL_TOP_K_RADIX_ROW_NEED];
+    }
+
+    for (int i = tid; i < SYCL_TOP_K_RADIX_HIST_SIZE; i += block_size) {
+        hist[i] = 0;
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    const uint32_t prefix = s_row[0];
+    const uint32_t mask   = s_row[1];
+    const uint32_t need   = s_row[2];
+
+    const int copy  = tid & (SYCL_TOP_K_RADIX_HIST_COPIES - 1);
+    const int chunk = (ncols + nparts - 1) / nparts;
+    const int col0  = part * chunk;
+    const int col1  = std::min(ncols, col0 + chunk);
+
+    for (int col = col0 + tid; col < col1; col += block_size) {
+        const uint32_t key = top_k_radix_key(src[col]);
+        if ((key & mask) == prefix) {
+            const uint32_t bucket = (key >> shift) & (SYCL_TOP_K_RADIX_BUCKETS - 1);
+            local_atomic(hist[bucket * SYCL_TOP_K_RADIX_HIST_COPIES + copy]).fetch_add(1u);
+        }
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    // One global atomic per bucket per group, not per element.
+    for (int b = tid; b < SYCL_TOP_K_RADIX_BUCKETS; b += block_size) {
+        uint32_t sum = 0;
+        for (int c = 0; c < SYCL_TOP_K_RADIX_HIST_COPIES; c++) {
+            sum += hist[b * SYCL_TOP_K_RADIX_HIST_COPIES + c];
+        }
+        if (sum) {
+            top_k_radix_gatomic(state[b]).fetch_add(sum);
+        }
+    }
+
+    // Publish this group's bins, then claim the scan if this group is the row's last.
+    // The group-wide barrier flushes the atomics above; only the claiming lane needs the
+    // release, so the device-scope fence is paid once per group rather than per work-item.
+    item_ct1.barrier(sycl::access::fence_space::global_and_local);
+    if (tid == 0) {
+        sycl::atomic_fence(sycl::memory_order::release, sycl::memory_scope::device);
+        sycl::atomic_ref<uint32_t, sycl::memory_order::acq_rel, sycl::memory_scope::device,
+                         sycl::access::address_space::global_space> done(state[SYCL_TOP_K_RADIX_ROW_DONE]);
+        *s_last = (done.fetch_add(1u) == (uint32_t) (nparts - 1)) ? 1u : 0u;
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    if (*s_last == 0u) {
+        return;
+    }
+    sycl::atomic_fence(sycl::memory_order::acquire, sycl::memory_scope::device);
+
+    // Lane t takes bucket 255 - t, so an inclusive scan counts down from the top bucket.
+    uint32_t cnt = 0;
+    if (tid < SYCL_TOP_K_RADIX_BUCKETS) {
+        cnt = top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_BUCKETS - 1 - tid]).load();
+    }
+    const uint32_t incl = sycl::inclusive_scan_over_group(item_ct1.get_group(), cnt, sycl::plus<uint32_t>());
+
+    if (tid < SYCL_TOP_K_RADIX_BUCKETS && incl >= need && incl - cnt < need) {
+        const uint32_t digit = (uint32_t) (SYCL_TOP_K_RADIX_BUCKETS - 1 - tid);
+        top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_PREFIX]).store(prefix | (digit << shift));
+        top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_MASK]).store(
+            mask | ((uint32_t) (SYCL_TOP_K_RADIX_BUCKETS - 1) << shift));
+        top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_NEED]).store(need - (incl - cnt));
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    // Clear for the next pass; the next launch is the barrier that orders this.
+    for (int b = tid; b < SYCL_TOP_K_RADIX_BUCKETS; b += block_size) {
+        top_k_radix_gatomic(state[b]).store(0u);
+    }
+    if (tid == 0) {
+        top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_DONE]).store(0u);
+    }
+}
+
+static void top_k_radix_split_emit_f32(
+    const float *   src,
+    int32_t *       dst_idx,
+    uint32_t *      state,
+    const int       ncols,
+    const int       k,
+    const int       part,
+    const int       nparts,
+    uint32_t *      slm,
+    const sycl::nd_item<1> & item_ct1
+) {
+    using local_atomic = sycl::atomic_ref<uint32_t, sycl::memory_order::relaxed,
+                                          sycl::memory_scope::work_group,
+                                          sycl::access::address_space::local_space>;
+
+    const int tid        = item_ct1.get_local_id(0);
+    const int block_size = item_ct1.get_local_range(0);
+
+    uint32_t * s_gt      = slm;
+    uint32_t * s_eq      = slm + 1;
+    uint32_t * s_base_gt = slm + 2;
+    uint32_t * s_base_eq = slm + 3;
+
+    uint32_t * s_row = slm + 4;   // prefix, mask, need
+
+    if (tid == 0) {
+        *s_gt = 0;
+        *s_eq = 0;
+        s_row[0] = state[SYCL_TOP_K_RADIX_ROW_PREFIX];
+        s_row[1] = state[SYCL_TOP_K_RADIX_ROW_MASK];
+        s_row[2] = state[SYCL_TOP_K_RADIX_ROW_NEED];
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    const uint32_t prefix = s_row[0];
+    const uint32_t mask   = s_row[1];
+    const uint32_t need   = s_row[2];
+
+    // Exactly k - need columns beat the pivot; the first need pivot-equal ones fill the tail.
+    const uint32_t base_eq = (uint32_t) k - need;
+
+    const int chunk = (ncols + nparts - 1) / nparts;
+    const int col0  = part * chunk;
+    const int col1  = std::min(ncols, col0 + chunk);
+
+    // Counting first and reserving one range per group keeps the row's two counters out of
+    // the inner loop: a per-element global atomic on a single address serialises the whole
+    // emit, and at k in the thousands that alone outweighs every read the kernel does.
+    for (int col = col0 + tid; col < col1; col += block_size) {
+        const uint32_t kp = top_k_radix_key(src[col]) & mask;
+        if (kp > prefix) {
+            local_atomic(*s_gt).fetch_add(1u);
+        } else if (kp == prefix) {
+            local_atomic(*s_eq).fetch_add(1u);
+        }
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    if (tid == 0) {
+        const uint32_t n_gt = *s_gt;
+        const uint32_t n_eq = *s_eq;
+        *s_base_gt = n_gt ? top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_CNT_GT]).fetch_add(n_gt) : 0u;
+        *s_base_eq = n_eq ? top_k_radix_gatomic(state[SYCL_TOP_K_RADIX_ROW_CNT_EQ]).fetch_add(n_eq) : 0u;
+        *s_gt = 0;
+        *s_eq = 0;
+    }
+    item_ct1.barrier(sycl::access::fence_space::local_space);
+
+    const uint32_t base_gt_g = *s_base_gt;
+    const uint32_t base_eq_g = *s_base_eq;
+
+    for (int col = col0 + tid; col < col1; col += block_size) {
+        const uint32_t kp = top_k_radix_key(src[col]) & mask;
+        if (kp > prefix) {
+            dst_idx[base_gt_g + local_atomic(*s_gt).fetch_add(1u)] = col;
+        } else if (kp == prefix) {
+            const uint32_t pos = base_eq_g + local_atomic(*s_eq).fetch_add(1u);
+            if (pos < need) {
+                dst_idx[base_eq + pos] = col;
+            }
+        }
+    }
+}
+
+static void top_k_radix_split_f32_sycl(
+    ggml_backend_sycl_context & ctx,
+    const float * src,
+    int32_t * dst_indices,
+    const int64_t ncols,
+    const int64_t nrows,
+    const int k,
+    const int nparts,
+    dpct::queue_ptr main_stream
+) {
+    GGML_ASSERT(ncols <= INT32_MAX);
+    GGML_ASSERT(nparts > 1);
+
+    const int block_size = ggml_sycl_info().max_work_group_sizes[ctx.device];
+    GGML_ASSERT(block_size >= SYCL_TOP_K_RADIX_BUCKETS);
+
+    const size_t state_words = (size_t) nrows * SYCL_TOP_K_RADIX_ROW_WORDS;
+    ggml_sycl_pool_alloc<uint32_t> state_alloc(ctx.pool(), state_words);
+    uint32_t * state = state_alloc.get();
+
+    // Zero histogram, done counter and both emit counters. prefix/mask/need are seeded by
+    // the first pass, which ignores the stored values.
+    // The queue is in-order, so the passes below are already ordered after this fill.
+    SYCL_CHECK(CHECK_TRY_ERROR(main_stream->memset(state, 0, state_words * sizeof(uint32_t))));
+
+    const sycl::range<1> block_dims(block_size);
+    const sycl::range<1> grid_dims(nrows * nparts);
+
+    bool first = true;
+    for (int shift = 32 - SYCL_TOP_K_RADIX_BITS; shift >= 0; shift -= SYCL_TOP_K_RADIX_BITS) {
+        const bool is_first = first;
+        first = false;
+        main_stream->submit([&](sycl::handler &cgh) {
+            sycl::local_accessor<uint32_t, 1> slm(sycl::range<1>(SYCL_TOP_K_RADIX_HIST_SIZE + 4), cgh);
+
+            cgh.parallel_for(
+                sycl::nd_range<1>(grid_dims * block_dims, block_dims),
+                [=](sycl::nd_item<1> item_ct1) {
+                    const int g    = item_ct1.get_group(0);
+                    const int row  = g / nparts;
+                    const int part = g % nparts;
+
+                    top_k_radix_split_pass_f32(
+                        src + (int64_t) row * ncols,
+                        state + (int64_t) row * SYCL_TOP_K_RADIX_ROW_WORDS,
+                        (int) ncols, k, shift, is_first, part, nparts,
+                        slm.get_multi_ptr<sycl::access::decorated::no>().get(),
+                        item_ct1);
+                });
+        });
+    }
+
+    main_stream->submit([&](sycl::handler &cgh) {
+        sycl::local_accessor<uint32_t, 1> slm(sycl::range<1>(8), cgh);
+
+        cgh.parallel_for(
+            sycl::nd_range<1>(grid_dims * block_dims, block_dims),
+            [=](sycl::nd_item<1> item_ct1) {
+                const int g    = item_ct1.get_group(0);
+                const int row  = g / nparts;
+                const int part = g % nparts;
+
+                top_k_radix_split_emit_f32(
+                    src + (int64_t) row * ncols,
+                    dst_indices + (int64_t) row * k,
+                    state + (int64_t) row * SYCL_TOP_K_RADIX_ROW_WORDS,
+                    (int) ncols, k, part, nparts,
+                    slm.get_multi_ptr<sycl::access::decorated::no>().get(),
+                    item_ct1);
+            });
+    });
+}
+
+void ggml_sycl_top_k_radix(
+    ggml_backend_sycl_context & ctx,
+    const float *   src,
+    int32_t *       dst_indices,
+    const int64_t   ncols,
+    const int64_t   nrows,
+    const int       k,
+    dpct::queue_ptr main_stream
+) {
+    const int nparts = top_k_radix_split_groups(ctx.device, ncols, nrows);
+    if (nparts > 1) {
+        top_k_radix_split_f32_sycl(ctx, src, dst_indices, ncols, nrows, k, nparts, main_stream);
+    } else {
+        top_k_radix_f32_sycl(ctx, src, dst_indices, ncols, nrows, k, main_stream);
+    }
+}
@@ -0,0 +1,24 @@
+#pragma once
+
+#include "common.hpp"
+
+// The legacy implementation uses SLM to implement sorting and top_k selection.
+// SLM is limited to 128KB on Xe, which limits how much can be sorted to k<32.
+// After a k=8, the radix selection becomes beneficial for most cases, because
+// scan-merge has (block + 1) * k pairs of (value, index). Given normal sorting of nlog(n),
+// radix-select becomes beneficial quite early. This sets it to 8 - however, the other parameters
+// (columns and rows) may also be a driving factor.
+// We select the legacy implementation for k below this constant because the overhead of radix select
+// exceeds the benefit for very small problems
+constexpr int SYCL_TOP_K_SCAN_MERGE_MAX_K = 8;
+
+// Top-k of every row of src, k indices per row into dst_indices, in no particular order.
+// Picks between the one-group-per-row and the split-row kernel from the shape and the device.
+void ggml_sycl_top_k_radix(
+    ggml_backend_sycl_context & ctx,
+    const float *   src,
+    int32_t *       dst_indices,
+    const int64_t   ncols,
+    const int64_t   nrows,
+    const int       k,
+    dpct::queue_ptr main_stream);
@@ -541,6 +541,7 @@ class MODEL_ARCH(IntEnum):
    ARWKV7           = auto()
    MAMBA            = auto()
    MAMBA2           = auto()
+    MAPLE            = auto()
    JAMBA            = auto()
    XVERSE           = auto()
    COMMAND_R        = auto()
@@ -1295,6 +1296,7 @@ MODEL_ARCH_NAMES: dict[MODEL_ARCH, str] = {
    MODEL_ARCH.ARWKV7:           "arwkv7",
    MODEL_ARCH.MAMBA:            "mamba",
    MODEL_ARCH.MAMBA2:           "mamba2",
+    MODEL_ARCH.MAPLE:            "maple",
    MODEL_ARCH.JAMBA:            "jamba",
    MODEL_ARCH.XVERSE:           "xverse",
    MODEL_ARCH.COMMAND_R:        "command-r",
@@ -3487,6 +3489,23 @@ MODEL_TENSORS: dict[MODEL_ARCH, list[MODEL_TENSOR]] = {
        MODEL_TENSOR.SSM_NORM,
        MODEL_TENSOR.SSM_OUT,
    ],
+    MODEL_ARCH.MAPLE: [
+        MODEL_TENSOR.TOKEN_EMBD,
+        MODEL_TENSOR.OUTPUT_NORM,
+        MODEL_TENSOR.OUTPUT,
+        MODEL_TENSOR.ATTN_NORM,
+        MODEL_TENSOR.ATTN_Q,
+        MODEL_TENSOR.ATTN_Q_NORM,
+        MODEL_TENSOR.ATTN_K,
+        MODEL_TENSOR.ATTN_K_NORM,
+        MODEL_TENSOR.ATTN_V,
+        MODEL_TENSOR.ATTN_OUT,
+        MODEL_TENSOR.FFN_NORM,
+        MODEL_TENSOR.FFN_GATE_INP,
+        MODEL_TENSOR.FFN_GATE_EXP,
+        MODEL_TENSOR.FFN_DOWN_EXP,
+        MODEL_TENSOR.FFN_UP_EXP,
+    ],
    MODEL_ARCH.JAMBA: [
        MODEL_TENSOR.TOKEN_EMBD,
        MODEL_TENSOR.OUTPUT_NORM,
@@ -62,6 +62,7 @@ static const std::map<llm_arch, const char *> LLM_ARCH_NAMES = {
    { LLM_ARCH_STARCODER2,       "starcoder2"       },
    { LLM_ARCH_MAMBA,            "mamba"            },
    { LLM_ARCH_MAMBA2,           "mamba2"           },
+    { LLM_ARCH_MAPLE,            "maple"            },
    { LLM_ARCH_JAMBA,            "jamba"            },
    { LLM_ARCH_FALCON_H1,        "falcon-h1"        },
    { LLM_ARCH_XVERSE,           "xverse"           },
@@ -67,6 +67,7 @@ enum llm_arch {
    LLM_ARCH_STARCODER2,
    LLM_ARCH_MAMBA,
    LLM_ARCH_MAMBA2,
+    LLM_ARCH_MAPLE,
    LLM_ARCH_JAMBA,
    LLM_ARCH_FALCON_H1,
    LLM_ARCH_XVERSE,
@@ -871,17 +871,18 @@ static void llama_grammar_advance_stack(
    std::set<llama_grammar_stack, decltype(stack_cmp)> seen(stack_cmp);

    while (!todo.empty()) {
-        llama_grammar_stack curr_stack = std::move(todo.back());
+        llama_grammar_stack curr_stack_candidate = std::move(todo.back());
        todo.pop_back();

-        if (seen.find( curr_stack) != seen.end()) {
+        auto [curr_stack_it, inserted] = seen.insert(std::move(curr_stack_candidate));
+        if (!inserted) {
            continue;
        }
-        seen.insert(curr_stack);
+        const llama_grammar_stack & curr_stack = *curr_stack_it;

        if (curr_stack.empty()) {
            if (std::find(new_stacks.begin(), new_stacks.end(), curr_stack) == new_stacks.end()) {
-                new_stacks.emplace_back(std::move(curr_stack));
+                new_stacks.emplace_back(curr_stack);
            }
            continue;
        }
@@ -924,7 +925,7 @@ static void llama_grammar_advance_stack(
        case LLAMA_GRETYPE_TOKEN_NOT:
            if (std::find(new_stacks.begin(), new_stacks.end(), curr_stack) == new_stacks.end()) {
                // only add the stack if it's not a duplicate of one we already have
-                new_stacks.emplace_back(std::move(curr_stack));
+                new_stacks.emplace_back(curr_stack);
            }
            break;
        default:
@@ -2225,7 +2225,7 @@ ggml_tensor * llm_graph_context::build_moe_ffn(
                    const float limit = hparams.swiglu_clamp_exp[il];
                    constexpr float eps = 1e-6f;
                    if (limit > eps) {
-                        if (arch == LLM_ARCH_DEEPSEEK4 || (arch == LLM_ARCH_DFLASH && hparams.dsv4_hc_mult > 0) || arch == LLM_ARCH_HY_V4) {
+                        if (arch == LLM_ARCH_MAPLE || arch == LLM_ARCH_DEEPSEEK4 || (arch == LLM_ARCH_DFLASH && hparams.dsv4_hc_mult > 0) || arch == LLM_ARCH_HY_V4) {
                            cur = ggml_swiglu_clamp(ctx0, cur, up, limit);
                        } else {
                            up = ggml_clamp(ctx0, up, -limit, limit);
@@ -33,6 +33,7 @@ bool llama_model_saver_supports_arch(llm_arch arch) {
        case LLM_ARCH_LAGUNA:
        case LLM_ARCH_GRANITE_SWA:
        case LLM_ARCH_DOTS3NOTE: // TODO: need to handle SWA pattern and MLA+SWA config
+        case LLM_ARCH_MAPLE:
            return false;
        default:
            return true;
@@ -162,6 +162,8 @@ static llama_model * llama_model_mapping(llm_arch arch, const llama_model_params
            return new llama_model_mamba(params);
        case LLM_ARCH_MAMBA2:
            return new llama_model_mamba2(params);
+        case LLM_ARCH_MAPLE:
+            return new llama_model_maple(params);
        case LLM_ARCH_JAMBA:
            return new llama_model_jamba(params);
        case LLM_ARCH_XVERSE:
@@ -3019,6 +3021,7 @@ llama_rope_type llama_model_rope_type(const llama_model * model) {
        case LLM_ARCH_SPARK2_5:
        case LLM_ARCH_TALKIE:
        case LLM_ARCH_MELLUM:
+        case LLM_ARCH_MAPLE:
            return LLAMA_ROPE_TYPE_NEOX;

        case LLM_ARCH_DFLASH:
@@ -4,7 +4,7 @@ void llama_model_gemma4_assistant::load_arch_hparams(llama_model_loader & ml) {
    hparams.n_embd_inp_impl = hparams.n_embd_out();

    hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
-    ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer_all);
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);

    uint32_t n_kv_shared_layers = 0;
    ml.get_key(LLM_KV_ATTENTION_SHARED_KV_LAYERS, n_kv_shared_layers, false);
@@ -2,7 +2,7 @@

 void llama_model_gemma4::load_arch_hparams(llama_model_loader & ml) {
    hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
-    ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer());
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);

    uint32_t n_kv_shared_layers = 0;
    ml.get_key(LLM_KV_ATTENTION_SHARED_KV_LAYERS, n_kv_shared_layers, false);
@@ -0,0 +1,150 @@
+#include "models.h"
+
+void llama_model_maple::load_arch_hparams(llama_model_loader & ml) {
+    hparams.swa_type = LLAMA_SWA_TYPE_STANDARD;
+
+    ml.get_key(LLM_KV_ATTENTION_LAYERNORM_RMS_EPS, hparams.f_norm_rms_eps);
+    ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW,    hparams.n_swa);
+    ml.get_key_or_arr(LLM_KV_EXPERT_FEED_FORWARD_LENGTH, hparams.n_ff_exp_arr, hparams.n_layer_all);
+
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);
+
+    hparams.rope_freq_base_train_swa  = hparams.rope_freq_base_train;
+    hparams.rope_freq_scale_train_swa = hparams.rope_freq_scale_train;
+    ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA, hparams.rope_freq_base_train_swa, false);
+
+    ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_EXP, hparams.swiglu_clamp_exp, hparams.n_layer_all);
+
+    switch (hparams.n_layer()) {
+        case 24: type = LLM_TYPE_20B; break;
+        default: type = LLM_TYPE_UNKNOWN;
+    }
+}
+
+void llama_model_maple::load_arch_tensors(llama_model_loader &) {
+    LLAMA_LOAD_LOCALS;
+
+    const int64_t n_ff_exp = hparams.n_ff_exp();
+    const int64_t head_dim = hparams.n_embd_head_k();
+
+    tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
+
+    output_norm = create_tensor(tn(LLM_TENSOR_OUTPUT_NORM, "weight"), {n_embd}, 0);
+    output      = create_tensor(tn(LLM_TENSOR_OUTPUT,      "weight"), {n_embd, n_vocab}, 0);
+
+    if (n_expert == 0) {
+        throw std::runtime_error("n_expert must be > 0 for Maple");
+    }
+    if (n_expert_used == 0) {
+        throw std::runtime_error("n_expert_used must be > 0 for Maple");
+    }
+
+    for (int i = 0; i < n_layer; ++i) {
+        auto & layer = layers[i];
+
+        layer.attn_norm = create_tensor(tn(LLM_TENSOR_ATTN_NORM, "weight", i), {n_embd}, 0);
+
+        create_tensor_qkv(layer, i, n_embd, n_head * head_dim, n_head_kv * head_dim, n_head_kv * head_dim, 0);
+        layer.wo = create_tensor(tn(LLM_TENSOR_ATTN_OUT, "weight", i), {n_head * head_dim, n_embd}, 0);
+
+        layer.attn_q_norm = create_tensor(tn(LLM_TENSOR_ATTN_Q_NORM, "weight", i), {head_dim}, 0);
+        layer.attn_k_norm = create_tensor(tn(LLM_TENSOR_ATTN_K_NORM, "weight", i), {head_dim}, 0);
+        layer.ffn_norm    = create_tensor(tn(LLM_TENSOR_FFN_NORM,    "weight", i), {n_embd}, 0);
+
+        layer.ffn_gate_inp  = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP,  "weight", i), {n_embd, n_expert}, 0);
+        layer.ffn_gate_exps = create_tensor(tn(LLM_TENSOR_FFN_GATE_EXPS, "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
+        layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {n_ff_exp, n_embd, n_expert}, 0);
+        layer.ffn_up_exps   = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS,   "weight", i), {n_embd, n_ff_exp, n_expert}, 0);
+    }
+}
+
+std::unique_ptr<llm_graph_context> llama_model_maple::build_arch_graph(const llm_graph_params & params) const {
+    return std::make_unique<graph>(*this, params);
+}
+
+llama_model_maple::graph::graph(const llama_model & model, const llm_graph_params & params) :
+    llm_graph_context(params) {
+    const int64_t n_embd_head = hparams.n_embd_head_k();
+
+    GGML_ASSERT(n_embd_head == hparams.n_embd_head_v());
+
+    ggml_tensor * inpL = build_inp_embd(model.tok_embd);
+    ggml_tensor * inp_pos = build_inp_pos();
+    auto * inp_attn = build_attn_inp_kv_iswa();
+    ggml_tensor * inp_out_ids = build_inp_out_ids();
+
+    for (int il = 0; il < n_layer; ++il) {
+        ggml_tensor * inpSA = inpL;
+
+        ggml_tensor * cur = build_norm(inpL, model.layers[il].attn_norm, nullptr, LLM_NORM_RMS, il);
+        cb(cur, "attn_norm", il);
+
+        {
+            auto [Qcur, Kcur, Vcur] = build_qkv(model.layers[il], cur, n_embd_head, n_head, n_head_kv, il);
+
+            Qcur = build_norm(Qcur, model.layers[il].attn_q_norm, nullptr, LLM_NORM_RMS, il);
+            Kcur = build_norm(Kcur, model.layers[il].attn_k_norm, nullptr, LLM_NORM_RMS, il);
+            cb(Qcur, "Qcur_normed", il);
+            cb(Kcur, "Kcur_normed", il);
+
+            if (hparams.is_swa(il)) {
+                const int64_t n_rot_l = hparams.n_rot(il);
+                const float freq_base_l = model.get_rope_freq_base(cparams, il);
+                const float freq_scale_l = model.get_rope_freq_scale(cparams, il);
+
+                Qcur = ggml_rope_ext(ctx0, Qcur, inp_pos, nullptr, n_rot_l, rope_type, n_ctx_orig, freq_base_l,
+                                     freq_scale_l, ext_factor, attn_factor, beta_fast, beta_slow);
+                Kcur = ggml_rope_ext(ctx0, Kcur, inp_pos, nullptr, n_rot_l, rope_type, n_ctx_orig, freq_base_l,
+                                     freq_scale_l, ext_factor, attn_factor, beta_fast, beta_slow);
+            }
+            cb(Qcur, "Qcur", il);
+            cb(Kcur, "Kcur", il);
+            cb(Vcur, "Vcur", il);
+
+            cur = build_attn(inp_attn,
+                    model.layers[il].wo, nullptr, model.layers[il].wo_s,
+                    Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, 1.0f / sqrtf(float(n_embd_head)), il);
+            cb(cur, "attn_out", il);
+        }
+
+        if (il == n_layer - 1 && inp_out_ids) {
+            cur   = ggml_get_rows(ctx0, cur, inp_out_ids);
+            inpSA = ggml_get_rows(ctx0, inpSA, inp_out_ids);
+        }
+
+        ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpSA);
+        cb(ffn_inp, "ffn_inp", il);
+
+        cur = build_norm(ffn_inp, model.layers[il].ffn_norm, nullptr, LLM_NORM_RMS, il);
+        cb(cur, "ffn_norm", il);
+
+        cur = build_moe_ffn(cur,
+                model.layers[il].ffn_gate_inp,
+                model.layers[il].ffn_up_exps,
+                model.layers[il].ffn_gate_exps,
+                model.layers[il].ffn_down_exps,
+                nullptr,
+                n_expert, n_expert_used,
+                LLM_FFN_SILU, true,
+                1.0f,
+                LLAMA_EXPERT_GATING_FUNC_TYPE_SOFTMAX,
+                il);
+        cb(cur, "ffn_moe_out", il);
+
+        cur = ggml_add(ctx0, cur, ffn_inp);
+        cur = build_cvec(cur, il);
+        cb(cur, "l_out", il);
+
+        inpL = cur;
+    }
+
+    ggml_tensor * cur = build_norm(inpL, model.output_norm, nullptr, LLM_NORM_RMS, -1);
+    cb(cur, "result_norm", -1);
+    res->t_embd = cur;
+
+    cur = build_lora_mm(model.output, cur, model.output_s);
+    cb(cur, "result_output", -1);
+    res->t_logits = cur;
+
+    ggml_build_forward_expand(gf, cur);
+}
@@ -9,7 +9,7 @@ void llama_model_mimo2::load_arch_hparams(llama_model_loader & ml) {
    ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW,   hparams.n_swa);
    ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA,         hparams.rope_freq_base_train_swa, false);

-    ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer());
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);

    float value_scale = 0.0f;
    if (ml.get_key(LLM_KV_ATTENTION_VALUE_SCALE, value_scale, false) && value_scale != 1.0f) {
@@ -945,6 +945,19 @@ struct llama_model_mamba2 : public llama_model_base {
 };


+struct llama_model_maple : public llama_model_base {
+    llama_model_maple(const struct llama_model_params & params) : llama_model_base(params) {}
+    void load_arch_hparams(llama_model_loader & ml) override;
+    void load_arch_tensors(llama_model_loader & ml) override;
+
+    struct graph : public llm_graph_context {
+        graph(const llama_model & model, const llm_graph_params & params);
+    };
+
+    std::unique_ptr<llm_graph_context> build_arch_graph(const llm_graph_params & params) const override;
+};
+
+
 struct llama_model_jamba : public llama_model_base {
    llama_model_jamba(const struct llama_model_params & params) : llama_model_base(params) {}
    void load_arch_hparams(llama_model_loader & ml) override;
@@ -145,8 +145,14 @@ void llama_model_nemotron_h::load_arch_tensors(llama_model_loader & ml) {
        const int64_t n_head_i       = hparams.n_head(i);
        const int64_t n_embd_k_gqa_i = hparams.n_embd_k_gqa(i);
        const int64_t n_embd_v_gqa_i = hparams.n_embd_v_gqa(i);
-        const int64_t n_ff_exp       = hparams.n_ff_exp(i) ? (int64_t)hparams.n_ff_exp(i) : n_ff / (int64_t)hparams.n_expert_used(i);
-        const int64_t n_ff_shexp     = hparams.n_ff_shexp;
+        const int64_t n_expert_used_i = hparams.n_expert_used(i);
+        const int64_t n_ff_exp_i      = hparams.n_ff_exp(i);
+        if (n_ff_exp_i == 0 && n_expert_used_i == 0) {
+            throw std::runtime_error(format("%s: layer %d declares neither expert_feed_forward_length nor expert_used_count, "
+                                            "cannot determine the expert FFN size", __func__, i));
+        }
+        const int64_t n_ff_exp   = n_ff_exp_i ? n_ff_exp_i : n_ff / n_expert_used_i;
+        const int64_t n_ff_shexp = hparams.n_ff_shexp;

        // NextN input-fusion tensors
        layer.nextn.enorm            = create_tensor(tn(LLM_TENSOR_NEXTN_ENORM,            "weight", i), {n_embd}, mtp_flags);
@@ -23,7 +23,7 @@ void llama_model_step35::load_arch_hparams(llama_model_loader & ml) {
    ml.get_key(LLM_KV_ATTENTION_SLIDING_WINDOW,  hparams.n_swa);
    ml.get_key(LLM_KV_ROPE_FREQ_BASE_SWA,        hparams.rope_freq_base_train_swa, false);

-    ml.get_key_or_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl, hparams.n_layer_all);
+    ml.get_arr(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, hparams.is_swa_impl);

    ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_EXP,   hparams.swiglu_clamp_exp,   hparams.n_layer_all, false);
    ml.get_key_or_arr(LLM_KV_SWIGLU_CLAMP_SHEXP, hparams.swiglu_clamp_shexp, hparams.n_layer_all, false);
@@ -11298,6 +11298,30 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
        }
    }

+    // qwen4exp sparse-attention indexer: nrows = n_tokens/n_stream, so tg gives nrows==1.
+    // Sweep nrows to expose how much of the device a single row leaves idle.
+    for (auto cols : {8192, 32768, 131072}) {
+        for (auto nrows : {1, 2, 4, 8, 16, 32}) {
+            test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {cols, nrows, 1, 1}, 2048));
+        }
+    }
+    // backend sampler: one row of the vocab (llama-sampler.cpp top_k)
+    for (auto k : {20, 40}) {
+        test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {151936, 1, 1, 1}, k));
+    }
+
+    // short rows, many of them: MoE routing and group selection. The opposite corner from
+    // the indexer, and the one where a work-group per row is the wasteful choice.
+    for (auto cols : {2, 16, 128, 1024}) {
+        for (auto nrows : {1024, 8192}) {
+            for (auto k : {1, 2, 8, 16, 32}) {
+                if (k <= cols) {
+                    test_cases.emplace_back(new test_top_k(GGML_TYPE_F32, {cols, nrows, 1, 1}, k));
+                }
+            }
+        }
+    }
+
    for (auto nrows : {1, 4, 8, 16}) {
        for (auto cols : {128, 1024, 4096, 8192, 16384, 32768, 65536, 131072, 200000, 2000000}) {
            test_cases.emplace_back(new test_cumsum(GGML_TYPE_F32, {cols, nrows, 1, 1}));
@@ -239,7 +239,8 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
        // SWA pattern: every 5th layer is full attention (matches E2B layer_types)
        ms.add_kv(LLM_KV_ATTENTION_SLIDING_WINDOW_PATTERN, uint32_t(5));
    } else if (arch == LLM_ARCH_COHERE2MOE || arch == LLM_ARCH_MIMO2 || arch == LLM_ARCH_STEP35 || arch == LLM_ARCH_SPARK2_5 ||
-            arch == LLM_ARCH_MUSE_GLIMMER || arch == LLM_ARCH_GRANITE_SWA || arch == LLM_ARCH_DOTS3NOTE) {
+            arch == LLM_ARCH_MUSE_GLIMMER || arch == LLM_ARCH_GRANITE_SWA || arch == LLM_ARCH_DOTS3NOTE ||
+            arch == LLM_ARCH_MAPLE) {
        std::vector<uint32_t> pattern;
        pattern.reserve(n_layer);
        for (uint32_t il = 0; il < n_layer; il++) {
@@ -323,6 +324,11 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
        ms.add_kv(LLM_KV_EXPERT_WEIGHTS_SCALE,                  1.0f);
        ms.add_kv(LLM_KV_EXPERT_WEIGHTS_NORM,                   true);
    }
+
+    if (arch == LLM_ARCH_MAPLE) {
+        ms.add_kv(LLM_KV_SWIGLU_CLAMP_EXP, 7.0f);
+    }
+
    ms.add_kv(LLM_KV_TOKENIZER_MODEL,         "no_vocab");
    // ms.add_kv(LLM_KV_DENSE_2_FEAT_OUT,     n_embd);
    // ms.add_kv(LLM_KV_DENSE_3_FEAT_IN,      n_embd);
@@ -505,6 +511,7 @@ static bool moe_mandatory(const llm_arch arch) {
        case LLM_ARCH_MISTRAL4:
        case LLM_ARCH_MELLUM:
        case LLM_ARCH_LAGUNA:
+        case LLM_ARCH_MAPLE:
            return true;
        default:
            return false;
@@ -563,7 +570,8 @@ static bool arch_supported(const llm_arch arch) {
    }
    // FIXME: these hit scheduler/view-backed-output issues with WebGPU on CI.
 #ifdef GGML_USE_WEBGPU
-    if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_DOTS3NOTE || arch == LLM_ARCH_QWEN4EXP) {
+    if (arch == LLM_ARCH_DEEPSEEK32 || arch == LLM_ARCH_GLM_DSA || arch == LLM_ARCH_DOTS3NOTE || arch == LLM_ARCH_QWEN4EXP ||
+            arch == LLM_ARCH_HY_V4) {
        return false;
    }
 #endif // GGML_USE_WEBGPU
@@ -221,7 +221,7 @@ static const remote_model_spec model_specs[] = {
    { "ggml-org/Step-3.5-Flash-GGUF",                 "Q4_K"   },
    { "ggml-org/Qwen3-Coder-Next-GGUF",               "Q8_0"   },
    { "ggml-org/Qwen3-14B-GGUF",                      "Q8_0"   },
-    { "ggml-org/NVIDIA-Nemotron-Nano-3-30B-A3B-GGUF", "Q8_0"   },
+    { "ggml-org/NVIDIA-Nemotron-3-Nano-30B-A3B-GGUF", "Q8_0"   },
    { "ggml-org/gpt-oss-120b-GGUF",                   "mxfp4"  },
    { "ggml-org/gemma-3-4b-it-GGUF",                  "Q8_0"   },
    { "bartowski/Meta-Llama-3.1-70B-Instruct-GGUF",   "Q4_K_M" },
Author	SHA1	Message	Date
Sigbjørn Skjæret	97e4ca7358	models : fix incorrect uses of get_key_or_arr (#28868 )	2026-09-14 14:05:37 +03:00
Sigbjørn Skjæret	1aca1f9fcd	models : fix mimo2 swa pattern load (#28865 )	2026-09-14 14:05:17 +03:00
Aaron Teo	be2c6d7d1f	tests(s390x): add non-vxe build to tests (#28776 ) * tests: add non-vxe build to tests Signed-off-by: Aaron Teo <[email protected]> ggml-cpu: add unused macro to fix ci Signed-off-by: Aaron Teo <[email protected]> Revert "ggml-cpu: temporarily add #28775 patch until its merged" This reverts commit d4645257b6b7e65c47b1b46baec3eb46a3f40968. Signed-off-by: Aaron Teo <[email protected]> * ggml-cpu: revert back to upstream/master Signed-off-by: Aaron Teo <[email protected]> --------- Signed-off-by: Aaron Teo <[email protected]>	2026-09-14 14:04:58 +03:00
Alex	3d10bcd197	llama: add Maple 20B-A1B ternary MoE architecture (CPU) (#27000 ) * gguf-py: add Maple tensor constants Add MODEL_ARCH.MAPLE, its "maple" name, and the tensor list for the Maple 20B-A1B ternary MoE architecture: token embeddings, output, attention with Q/K RMS norms, and per-expert FFN tensors. * convert: add Maple HF->GGUF converter Register MapleForCausalLM in the HF architecture map and add the converter for the Maple 20B-A1B ternary MoE model: 24 layers, 256 experts with 8 active, sliding-window attention (SWA-512) interleaved with global attention at a 3:1 ratio, partial rotary factor 0.5, and per-expert weight stacking into merged 3D tensors. * llama: add Maple architecture (20B-A1B ternary MoE) Add the Maple 20B-A1B ternary MoE architecture: 24 layers, 256 experts with 8 active, sliding-window attention (SWA-512) interleaved with global attention at a 3:1 ratio, and ternary TQ1_0/TQ2_0 quantization support. - register LLM_ARCH_MAPLE between MAMBA2 and JAMBA - implement llama_model_maple: Q/K RMS norms after projection (GEMMA4 style), rope applied only on SWA layers (nope_on_global_attention), ISWA KV cache, and MoE FFN with swiglu gate clamp at +7 (DEEPSEEK4 style) - mark MAPLE as unsupported by the model saver (roundtrip skipped) * tests: mark Maple as MoE-mandatory Maple is always-MoE: the model throws when n_expert == 0, so the test harness must only run the MoE config for LLM_ARCH_MAPLE. * maple: apply review feedback (n_ff_exp_arr, get_arr, rope params) - load_arch_hparams: use n_ff_exp_arr + n_ff_exp() accessor (upstream changed these from a scalar member during the rebase) - sliding_window_pattern: get_arr, the pattern is mandatory for this arch - partial_rotary_factor: read only from rope_parameters (base.py mirrors the top-level key automatically) - document why TOKEN_EMBD/OUTPUT are forced to F16 (they are the two dense tensors in Maple, and the reference GGUFs ship them as F16) - add @ModelBase.example("deepgrove/maple-preview") * tests: add Maple to the SWA pattern array list get_arr for maple.attention.sliding_window_pattern requires an array, but the harness only emitted a per-layer array for the arches in its list, so test-llama-archs -a maple failed to load the model. Assisted-by: DeepSeek Harness * maple: move swiglu_clamp_exp to the converter The loader prefilled 7.0 and read the key optionally. The converter now writes it and the loader reads it as required, because llama-graph.cpp skips the clamp when the limit is 0 and an optional read would silently run unclamped. The test harness provides the key for the same reason. Also drops tensor_force_quant: base.py already forces FFN_GATE_INP to F32 and TOKEN_EMBD/OUTPUT to F16 for ternary file types. Assisted-by: DeepSeek Harness * convert: fix the LazyBase func signature in the Maple converter ty flagged the stack() closure: it takes no argument, while LazyBase is annotated with func: Callable[[Any], Any]. Pass the tensor list through args instead of closing over it, the same way kimi_k3 does, so the callable shape matches. Assisted-by: DeepSeek Harness	2026-09-14 14:04:05 +03:00
cwriterandcwriter	21f6b0d22c	sycl: rfc: Use radix select for top_k (#28670 ) * sycl: GPU-resident TOP_K for large k, parallelised over the device The SYCL backend refused GGML_OP_TOP_K above k = 32 and let it fall back to the CPU, a backend round-trip per call. The limit was not conservatism: the scan-merge kernels keep (split_block + 1) * k candidate (value, index) pairs in SLM, so at k = 128 a work-group already needs 132 KB and cannot launch. qwen4exp's sparse-attention indexer asks for k = 2048 in 12 layers on every token, so this fired at every context length. Add a radix select for large k. The k-th largest is found by four most-significant-first passes over an order-preserving unsigned key: histogram the digit over the candidate set, walk the buckets from the top, and recurse into the one where the running count reaches what is still needed. SLM holds the histogram rather than candidates, so the footprint is independent of k. A final pass emits every column beating the pivot plus exactly as many pivot-equal columns as are still missing, so duplicate keys still yield exactly k distinct indices. Output order is not required and is not paid for: ggml-cpu/ops.cpp swaps its first two outputs to say so. The key folds -0.0 onto +0.0 so its equivalence classes match the reference comparator, under which the two tie. NaN has no defined order in the reference (its comparator is not a strict weak order there); here +NaN keys above +inf and -NaN below -inf, which at least makes the result deterministic. One work-group per row leaves the device idle whenever a graph has fewer rows than it has cores, which at batch size 1 means one work-group full stop: qwen4exp tops-k a tensor of shape [n_kv, n_tokens/n_stream, n_stream], so token generation gives nrows == 1, and the backend sampler reshapes logits to a single row as well. Measured, ne=[200000,1] and ne=[200000,16] cost 358.0 us and 363.4 us -- sixteen rows for 1.5% more wall-clock. So also spread a row over several groups when there are too few rows to cover the device. Per-pass state moves to global memory and each digit pass becomes its own launch, since a work-group barrier can no longer span the row. Groups accumulate in SLM and contribute 256 global atomics each, keeping global traffic per-group rather than per-element, and the last group of a row -- the one whose fetch_add returns G-1 -- performs that pass's scan, holding the launch count at one per digit plus one emit. The group count comes from the device and is floor-divided by nrows, so a row count that already covers the device is left whole and pays nothing. Below 64K columns the single-group kernel finishes inside the cost of the extra launches and stays in charge. Reading the row's prefix/mask/need through a device-scope atomic_ref costs more than the sweep it guards: those loads are uncached, so passes 2-4 ran at 49 us against 12 us for pass 1. One lane reads them into SLM and the group takes them from there -- 208 us -> 44.6 us at ne=[131072,1], k=2048. The block size now takes the device's max_work_group_size instead of a cap of 512. The cap was never a floor, so a device reporting 512 is unaffected; one allowing 1024 was being given half its width. Finally, put the scan-merge gate where the two paths actually cross. That kernel's cost climbs with k while the radix select's does not; measured over widths from 2 to 200K columns and row counts from 1 to 8192, radix is ahead everywhere from k = 8 up and behind at k <= 2, where scan-merge's smaller fixed cost wins. The short-row corner (ncols=2, nrows=65536, as in bailingmoe2 group selection) is exactly where radix loses at low k, and the gate keeps it on scan-merge. Op-level against the CPU-fallback path this replaces, and against the single-group radix select for the split: 4.98x at ne=[131072,1] k=2048, 6.65x at ne=[151936,1] k=40, 13.35x at k=20, 118x at ne=[65000,16] k=32. No measured shape regressed. End to end on 3x Arc Pro B60 with Qwen3.8-Flash-Next UD-IQ4_XS, llama-bench tg64, the parallelisation is worth 5.91 -> 6.05 t/s at d=131072 and a wash at shallower depths. Perplexity over wikitext-2 is unchanged within noise at both 512 and 81920 context. test-backend-ops: 525/525 TOP_K (previously every k > 32 case was refused), 880/880 MUL_MAT_ID. Perf coverage added for k > 32 at large widths and for the short-row corner, neither of which was exercised before. * move topk-select to topk-radix.{cpp\|hpp} --------- Co-authored-by: cwriter <cwriter@localhost>	2026-09-14 14:02:44 +03:00
Georgi Gerganov	2f539596c6	ggml-cpu : disable PCH and fix CACHE_LINE_SIZE ambiguity to fix heap corruption (#28882 ) Disable the ggml-cpu precompiled header and remove the std::hardware_destructive_interference_size branch from CACHE_LINE_SIZE. The PCH force-includes ggml-impl.h before ops.h, which pulls in <new> via <array>/<vector> and defines __cpp_lib_hardware_interference_size. This makes the C++ kernels use CACHE_LINE_SIZE = 256 (hardware destructive interference size) while the C work-buffer sizing code in ggml-cpu.c always uses the fallback 64. The mismatch undersizes the rope work buffer by (CACHE_LINE_SIZE/4 - 16) * n_threads * 4 bytes, causing a heap-buffer-overflow that corrupts the heap and later crashes in ggml_compute_forward_rope_flt. Disabling the ggml-cpu PCH restores the natural include order so ops.h is processed before <new>, keeping CACHE_LINE_SIZE consistent. Removing the std::hardware_destructive_interference_size branch makes the value deterministic and include-order independent. ref: https://github.com/ggml-org/llama.cpp/issues/28858 Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp	2026-09-14 13:03:41 +03:00
Georgi Gerganov	89fe242405	ci : trigger self-hosted CI on changes to ci/run.sh (#28859 ) The workflow's push/pull_request path filters did not include the ci/run.sh script that all of its jobs execute, so changes to it never re-triggered the self-hosted CI. Assisted-by: pi:llama.cpp/Qwen3.8-27B	2026-09-14 11:51:06 +03:00
Georgi Gerganov	15d8f2d592	ci : remove gg_sum summary logic (#28857 ) Assisted-by: pi:llama.cpp/DeepSeek-V4-Flash-Vision-Exp	2026-09-14 11:50:43 +03:00
Łukasz Ślusarczyk	661643e430	sycl : fix oneDNN scratchpad breaking the pool free order (#28704 )	2026-09-14 02:24:06 -04:00
Daniel Bevenius	093a2f86c3	common : move llama_n_rs_seq to before llama_decode (#28749 ) This commit moves the llama_n_rs_seq function call to before the llama_decode call and returns directly if the check is true, removing the setting of res and the goto statement. The motivation for this change is to avoid the llama_decode call if it is not needed.	2026-09-14 05:24:05 +02:00
thelittlefiremanandJohannes Gäßler	ad6c66839a	ggml-cuda: fallback to F32 on device without BF16 hardware acceleration (#28846 ) * ggml-cuda: fallback to F32 on device without BF16 hardware acceleration: (Nvidia >= AMPERE, AMD >= RDNA3 or = CDNA) * apply logic to NVIDIA as well --------- Co-authored-by: Johannes Gäßler <[email protected]>	2026-09-14 00:05:10 +02:00
Clint Herron	7a16a6ce32	grammar : coalesce find + insert into a single insert and adjust move/copy mechanics (#26885 ) 1) Combine two consecutive lookups (find + insert) into a single insert-attempt/lookup routine so that we don't per form two O(log(n)) lookup operations in a row anymore -- we only need to do it once and then see if the insert succeeded. 2) Instead of copying every potential stack (expensive) and then moving it (cheap) to new_stacks when it's a final output state, we switch the order so that we move every potential stack (cheap), and then only copy it (expensive) to new stacks when it's a final output state. There are a LOT of intermediate states that get generated, and unless they become final output states, then all of these expensive intermediate copies are wasted. Before: lookup -> lookup/insert + copy -> optional move to output New: lookup/insert + move -> optional copy to output	2026-09-13 16:56:46 -05:00
fairydreamingandStanisław Szymczyk	5f436dddb4	tests : exclude HY_V4 from WebGPU test-llama-archs tests (#28855 ) Co-authored-by: Stanisław Szymczyk <[email protected]>	2026-09-13 19:24:11 +02:00
Yaniss Amazouz	e49d2c2760	models : guard the expert FFN size fallback in nemotron-h against a zero divisor (#28779 ) The NextN/MTP tail loop derives the expert FFN size as n_ff/n_expert_used when expert_feed_forward_length gives nothing for the layer. Both values come from per-layer arrays that legitimately hold 0 on layers that are not MoE, so a checkpoint whose predict layers hold 0 in both divides by zero and dies with SIGFPE at load time, with no error message. Report the malformed metadata instead.	2026-09-13 19:20:50 +02:00
Bernard Ladenthin	6978052985	ggml-cpu(s390x): guard VXE-only repack helpers (#28775 )	2026-09-14 01:18:53 +08:00
Michael Taylor	243a3082d4	tests : fix typo in test-quant-type-selection for nemotron 3 nano (#28835 ) Corrects a typo in `tests/test-quant-type-selection` for the Nvidia Nemotron 3 Nano 30B A3B model, which was referred to as nvidia-nemotron-nano-3-30b-a3b. The error made the test skip that test case, rather than failing the test. [no release]	2026-09-13 18:50:46 +02:00
Neo Zhang Titaniumtown Neo Zhang JianyuNeo Zhang <NA>	b6b003d2cb	sycl : Fix get mem error (#28227 ) * fix for unsupport zes API * optimize the code * adjust the log level * rm unused head files * Update docs/backend/SYCL.md Co-authored-by: Titaniumtown <[email protected]> * fix the error to detect level zero SDK/dev package, stop build after detect the error * update the message * fix the build error when missed to install level zero dev package * rm GGML_SYCL_DEV_DEBUG, mv read env vars in all entry functions --------- Co-authored-by: Neo Zhang Jianyu <[email protected]> Co-authored-by: Titaniumtown <[email protected]> Co-authored-by: Neo Zhang <NA>	2026-09-13 18:31:34 +03:00
Georgi Gerganov	c95f8e47b8	ci : run editorconfig and code-style checks on ubuntu-slim (#28854 ) Move the EditorConfig Checker and Code Style Checker workflows from the `[self-hosted, fast]` runners to `ubuntu-slim`, which is an established runner label in the repo. Assisted-by: pi:llama.cpp/Qwen3.8-27B	2026-09-13 18:16:45 +03:00
Georgi Gerganov	bc52a12b38	pi : prefer PI_MODEL_NAME env var for model disclosure (#28853 ) Assisted-by: pi:llama.cpp/Qwen3.8-27B	2026-09-13 18:13:32 +03:00