sycl: make --fit respect --fit-target better (#27629 )

improve the --fit algorithm to take into account the actual peak required VRAM for a given context size on a SYCL backend. This includes both properly accounting for how much VRAM is required when the allocated context is fully used (which makes the reported context drop below what it did before, but stop it OOMing) as well as preventing some overly-conservative calculations which meant too much VRAM was being reserved. Tested on a Arc b70 with unsloth's qwen3.8 (Q4_K_XL), able to get 262144 context, fully usable, with q8_0 KV and MTP and 4k ubatch size using --fit-target 1
vulkan: combine duplicated fastdiv functions, rename the one optimizing small divs (#27526 )
2026-08-29 11:37:42 +02:00 · 2026-08-29 05:00:09 -04:00 · 2026-08-29 10:59:48 +03:00 · 2026-08-29 15:12:23 +08:00 · 2026-08-29 10:09:24 +03:00 · 2026-08-28 14:01:59 -07:00
40 changed files with 931 additions and 217 deletions
@@ -2735,9 +2735,9 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
        "- auto: on, but only for tensors larger than 4 GiB\n"
        "- off: always keep them resident",
        [](common_params & params, const std::string & value) {
-            /**/ if (value == "on")   { params.tensor_read_lazy = LLAMA_TENSOR_READ_LAZY_ON;   }
-            else if (value == "auto") { params.tensor_read_lazy = LLAMA_TENSOR_READ_LAZY_AUTO; }
-            else if (value == "off")  { params.tensor_read_lazy = LLAMA_TENSOR_READ_LAZY_OFF;  }
+            /**/ if (value == "on")   { params.lazy_mode = LLAMA_LAZY_MODE_ON;   }
+            else if (value == "auto") { params.lazy_mode = LLAMA_LAZY_MODE_AUTO; }
+            else if (value == "off")  { params.lazy_mode = LLAMA_LAZY_MODE_OFF;  }
            else { throw std::invalid_argument("invalid value"); }
        }
    ).set_env("LLAMA_ARG_TENSOR_READ_LAZY"));
@@ -1688,7 +1688,7 @@ struct llama_model_params common_model_params_to_llama(common_params & params) {
    mparams.main_gpu        = params.main_gpu;
    mparams.split_mode      = params.split_mode;
    mparams.load_mode       = params.load_mode;
-    mparams.tensor_read_lazy = params.tensor_read_lazy;
+    mparams.lazy_mode = params.lazy_mode;
    mparams.tensor_split    = params.tensor_split;
    mparams.check_tensors   = params.check_tensors;
    mparams.use_extra_bufts = !params.no_extra_bufts;
@@ -483,7 +483,7 @@ struct common_params {
    enum llama_split_mode split_mode = LLAMA_SPLIT_MODE_LAYER; // how to split the model across GPUs
    enum llama_load_mode  load_mode  = LLAMA_LOAD_MODE_AUTO; // how to load the model

-    enum llama_tensor_read_lazy tensor_read_lazy = LLAMA_TENSOR_READ_LAZY_AUTO; // on-demand reading of tensors marked by the arch
+    enum llama_lazy_mode lazy_mode = LLAMA_LAZY_MODE_AUTO; // on-demand reading of tensors marked by the arch

    common_cpu_params cpuparams;
    common_cpu_params cpuparams_batch;
@@ -24,7 +24,18 @@ must be included in the .cat file digitally signed with a trusted certificate.
 This document covers details on how to generate personal certificate files (.pfx) and how to configure the system
 to allow for test signatures (aka test-signing).

-## Install the latest Adreno OpenCL SDK
+## Install Windows SDKs
+
+The recommended method is `setup-sdk.py`:
+
+```
+> python scripts\snapdragon\setup-sdk.py --list-sdk-releases
+> python scripts\snapdragon\setup-sdk.py --hexagon --opencl
+```
+
+It installs the selected SDKs under `C:\Qualcomm` and sets their corresponding environment variables for the current user. Start a new terminal after it completes; native Windows builds check all SDK paths before CMake runs.
+
+Select the SDKs to install with `--hexagon` and `--opencl`; use both to prepare a dual-backend build. To select a different available version, pass it to the SDK option, for example `--hexagon 6.4.0.2`. SDK versions install side by side, so you can switch versions without deleting an existing installation. Use `--force` to reinstall the selected SDKs. Use a new CMake build directory after each switch because CMake caches the SDK paths.

 Either use the trimmed down version (optimized for CI) from

@@ -2936,12 +2936,13 @@ struct ggml_cplan ggml_graph_plan(
                        const int64_t ne10 = node->src[1]->ne[0]; // W
                        const int64_t ne11 = node->src[1]->ne[1]; // H
                        const int64_t ne12 = node->src[1]->ne[2]; // Channels In
+                        const int64_t ne13 = node->src[1]->ne[3]; // Batch

                        GGML_ASSERT(node->src[0]->type == GGML_TYPE_F16 || node->src[0]->type == GGML_TYPE_F32);
                        GGML_ASSERT(node->src[1]->type == GGML_TYPE_F32);

                        cur += ggml_type_size(node->src[0]->type) * ne00 * ne01 * ne02 * ne03;
-                        cur += ggml_type_size(node->src[0]->type) * ne10 * ne11 * ne12;
+                        cur += ggml_type_size(node->src[0]->type) * ne10 * ne11 * ne12 * ne13;

                    } break;
                case GGML_OP_TOP_K:
@@ -7267,18 +7267,21 @@ static void ggml_compute_forward_conv_transpose_2d_impl(
            }
        }

-        // permute source data (src1) from (Sw x Sh x Cin) to (Cin x Sw x Sh)
+        // permute source data (src1) from (Sw x Sh x Cin) to (Cin x Sw x Sh), for all batches
        {
            kernel_t * const wdata = (kernel_t *) params->wdata + nk;
-            for (int i12 = 0; i12 < ne12; i12++) {
-                for (int i11 = 0; i11 < ne11; i11++) {
-                    const float * const src = (float *)((char *) src1->data + i12*nb12 + i11*nb11);
-                    kernel_t * dst_data = wdata + i11*ne10*ne12;
-                    for (int i10 = 0; i10 < ne10; i10++) {
-                        if constexpr (std::is_same_v<kernel_t, ggml_fp16_t>) {
-                            dst_data[i10*ne12 + i12] = GGML_CPU_FP32_TO_FP16(src[i10]);
-                        } else {
-                            dst_data[i10*ne12 + i12] = src[i10];
+            for (int i13 = 0; i13 < ne13; i13++) {
+                kernel_t * const wdata_b = wdata + i13*ne10*ne11*ne12;
+                for (int i12 = 0; i12 < ne12; i12++) {
+                    for (int i11 = 0; i11 < ne11; i11++) {
+                        const float * const src = (float *)((char *) src1->data + i13*nb13 + i12*nb12 + i11*nb11);
+                        kernel_t * dst_data = wdata_b + i11*ne10*ne12;
+                        for (int i10 = 0; i10 < ne10; i10++) {
+                            if constexpr (std::is_same_v<kernel_t, ggml_fp16_t>) {
+                                dst_data[i10*ne12 + i12] = GGML_CPU_FP32_TO_FP16(src[i10]);
+                            } else {
+                                dst_data[i10*ne12 + i12] = src[i10];
+                            }
                        }
                    }
                }
@@ -7305,24 +7308,27 @@ static void ggml_compute_forward_conv_transpose_2d_impl(
    kernel_t * const wdata_src = wdata + nk;

    for (int i2 = ip0; i2 < ip1; i2++) { // Cout
-        float * dst_data = (float *)((char *) dst->data + i2*nb2);
        kernel_t * wdata_kernel = wdata + i2*ne01*ne00*ne03;
-        for (int i11 = 0; i11 < ne11; i11++) {
-            for (int i10 = 0; i10 < ne10; i10++) {
-                const int i1n = i11*ne10*ne12 + i10*ne12;
-                for (int i01 = 0; i01 < ne01; i01++) {
-                    for (int i00 = 0; i00 < ne00; i00++) {
-                        float v = 0;
-                        if constexpr (std::is_same_v<kernel_t, ggml_fp16_t>) {
-                            ggml_vec_dot_f16(ne03, &v, 0,
-                                    wdata_src + i1n, 0,
-                                    wdata_kernel + i01*ne00*ne03 + i00*ne03, 0, 1);
-                        } else {
-                            ggml_vec_dot_f32(ne03, &v, 0,
-                                    wdata_src + i1n, 0,
-                                    wdata_kernel + i01*ne00*ne03 + i00*ne03, 0, 1);
+        for (int i3 = 0; i3 < ne3; i3++) { // batch
+            float * dst_data = (float *)((char *) dst->data + i3*nb3 + i2*nb2);
+            kernel_t * wdata_src_b = wdata_src + i3*ne10*ne11*ne12;
+            for (int i11 = 0; i11 < ne11; i11++) {
+                for (int i10 = 0; i10 < ne10; i10++) {
+                    const int i1n = i11*ne10*ne12 + i10*ne12;
+                    for (int i01 = 0; i01 < ne01; i01++) {
+                        for (int i00 = 0; i00 < ne00; i00++) {
+                            float v = 0;
+                            if constexpr (std::is_same_v<kernel_t, ggml_fp16_t>) {
+                                ggml_vec_dot_f16(ne03, &v, 0,
+                                        wdata_src_b + i1n, 0,
+                                        wdata_kernel + i01*ne00*ne03 + i00*ne03, 0, 1);
+                            } else {
+                                ggml_vec_dot_f32(ne03, &v, 0,
+                                        wdata_src_b + i1n, 0,
+                                        wdata_kernel + i01*ne00*ne03 + i00*ne03, 0, 1);
+                            }
+                            dst_data[(i11*stride + i01)*ne0 + i10*stride + i00] += v;
                        }
-                        dst_data[(i11*stride + i01)*ne0 + i10*stride + i00] += v;
                    }
                }
            }
@@ -660,6 +660,7 @@ typedef struct {
    uint64_t nb0;
    uint64_t nb1;
    uint64_t nb2;
+    uint64_t nb3;
 } ggml_metal_kargs_conv_transpose_2d;

 typedef struct {
@@ -4645,6 +4645,7 @@ int ggml_metal_op_conv_transpose_2d(ggml_metal_op_t ctx, int idx) {
    const int32_t OW = op->ne[0];
    const int32_t OH = op->ne[1];
    const int32_t OC = op->ne[2];
+    const int32_t N  = op->src[1]->ne[3];

    ggml_metal_kargs_conv_transpose_2d args = {
        /*.IC  =*/ IC,
@@ -4657,6 +4658,7 @@ int ggml_metal_op_conv_transpose_2d(ggml_metal_op_t ctx, int idx) {
        /*.nb0 =*/ nb0,
        /*.nb1 =*/ nb1,
        /*.nb2 =*/ nb2,
+        /*.nb3 =*/ nb3,
    };

    auto pipeline = ggml_metal_library_get_pipeline_conv_transpose_2d(lib, op);
@@ -4671,7 +4673,7 @@ int ggml_metal_op_conv_transpose_2d(ggml_metal_op_t ctx, int idx) {
    const size_t smem = GGML_PAD(KW * KH * sizeof(float), 16);
    ggml_metal_encoder_set_threadgroup_memory_size(enc, smem, 0);

-    ggml_metal_encoder_dispatch_threadgroups(enc, OW, OH, OC, KW, KH, 1);
+    ggml_metal_encoder_dispatch_threadgroups(enc, OW, OH, OC * N, KW, KH, 1);

    return 1;
 }
@@ -279,6 +279,194 @@ constexpr fa_vec_entry_t fa_vec_tuned_table[] = {
    { { GGML_METAL_DEVICE_M1_PRO, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
    { { GGML_METAL_DEVICE_M1_PRO, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },

+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 64, 64, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 128, 128, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 128, 128, 3, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 128, 128, 3, 4 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 192, 128, 1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 192, 128, 2, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 192, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 320, 256, 2, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 320, 256, 3, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_F16, 320, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 32, 32, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 32, 32, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 32, 32, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 32, 32, 3, 2 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 32, 32, 3, 3 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 64, 64, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 64, 64, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 64, 64, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 64, 64, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 96, 96, 1, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 96, 96, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 96, 96, 2, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 96, 96, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 96, 96, 3, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 96, 96, 3, 4 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 128, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 128, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 192, 192, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 192, 192, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 192, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 192, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 256, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 256, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 320, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 320, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 320, 256, 1, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 512, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 512, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 576, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_0, 576, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 32, 32, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 32, 32, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 32, 32, 1, 4 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 32, 32, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 32, 32, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 32, 32, 3, 3 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 64, 64, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 64, 64, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 96, 96, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 96, 96, 2, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 96, 96, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 96, 96, 3, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 128, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 128, 128, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 128, 128, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 128, 128, 1, 4 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 128, 128, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 128, 128, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 192, 192, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 192, 192, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 192, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 192, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 256, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 256, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 320, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 320, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 320, 256, 1, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 512, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 512, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 576, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q4_1, 576, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 32, 32, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 32, 32, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 32, 32, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 32, 32, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 64, 64, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 64, 64, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 64, 64, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 64, 64, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 96, 96, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 128, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 128, 128, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 128, 128, 1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 128, 128, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 128, 128, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 128, 128, 2, 4 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 128, 128, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 192, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 192, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 192, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 192, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 192, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 128, 1, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 128, 2, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 192, 128, 3, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 256, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 256, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 256, 256, 1, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 320, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 320, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 320, 256, 1, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 512, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 512, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 576, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_0, 576, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 32, 32, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 32, 32, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 32, 32, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 32, 32, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 64, 64, -1, 1 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 64, 64, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 64, 64, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 64, 64, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 96, 96, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 96, 96, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 128, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 128, 128, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 128, 128, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 128, 128, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 128, 128, 2, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 128, 128, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 192, 192, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 192, 192, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 192, 192, 1, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 192, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 192, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 192, 128, 1, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 192, 128, 2, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 192, 128, 3, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 256, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 256, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 256, 256, 1, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 256, 256, 2, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 320, 256, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 320, 256, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 320, 256, 1, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 320, 256, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 320, 256, 2, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 320, 256, 2, 4 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 320, 256, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 320, 256, 3, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 512, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 512, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 576, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q5_1, 576, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 32, 32, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 32, 32, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 32, 32, 2, 4 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 32, 32, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 32, 32, 3, 3 }, { 4, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 64, 64, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 64, 64, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 64, 64, 1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 64, 64, 1, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 64, 64, 2, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 64, 64, 3, 2 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 96, 96, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 96, 96, 2, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 96, 96, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 96, 96, 3, 3 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 96, 96, 3, 4 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 128, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 128, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 192, 192, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 192, 192, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 192, 128, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 192, 128, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 256, 256, -1, 0 }, { 1, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 256, 256, -1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 320, 256, 1, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 320, 256, 1, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 320, 256, 2, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 320, 256, 2, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 320, 256, 3, 1 }, { 2, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 320, 256, 3, 3 }, { 4, 2 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 512, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 512, 512, -1, 1 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 576, 512, -1, 0 }, { 1, 4 } },
+    { { GGML_METAL_DEVICE_M1_MAX, GGML_TYPE_Q8_0, 576, 512, -1, 1 }, { 1, 4 } },
+
    { { GGML_METAL_DEVICE_M2_ULTRA, GGML_TYPE_F16, 64, 64, -1, 0 }, { 1, 4 } },
    { { GGML_METAL_DEVICE_M2_ULTRA, GGML_TYPE_F16, 64, 64, -1, 1 }, { 1, 4 } },
    { { GGML_METAL_DEVICE_M2_ULTRA, GGML_TYPE_F16, 128, 128, 2, 1 }, { 1, 4 } },
@@ -366,7 +366,8 @@ kernel void kernel_conv_transpose_2d(

    const int64_t out_x = tgpig[0];
    const int64_t out_y = tgpig[1];
-    const int64_t out_c = tgpig[2];
+    const int64_t batch = tgpig[2] / args.OC;
+    const int64_t out_c = tgpig[2] % args.OC;

    const int64_t kw = tpitg[0];
    const int64_t kh = tpitg[1];
@@ -390,7 +391,7 @@ kernel void kernel_conv_transpose_2d(

        if (in_x >= args.IW) continue;

-        const int64_t input_idx = (args.IW * args.IH) * in_c + (args.IW) * in_y + in_x;
+        const int64_t input_idx = (args.IW * args.IH) * (args.IC * batch + in_c) + (args.IW) * in_y + in_x;
        const int64_t kernel_idx = (args.KH * args.KW * args.OC) * in_c + (args.KH * args.KW) * out_c + (args.KW) * kh + kw;

        v += (float)src0[kernel_idx] * src1[input_idx];
@@ -408,7 +409,7 @@ kernel void kernel_conv_transpose_2d(
            total += shared_sum[i];
        }

-        device float * dst_ptr = (device float *) (dst + out_x*args.nb0 + out_y * args.nb1 + out_c*args.nb2);
+        device float * dst_ptr = (device float *) (dst + batch*args.nb3 + out_c*args.nb2 + out_y * args.nb1 + out_x*args.nb0);
        dst_ptr[0] = total;
    }
 }
@@ -6,6 +6,7 @@
 #include "convert.hpp"
 #include "vecdotq.hpp"
 #include "fattn-buffers.hpp"
+#include "fattn.hpp"

 #include "ggml.h"

@@ -926,6 +927,7 @@ void launch_fattn(

    ggml_sycl_fattn_alloc        K_f16(fbuf.K);
    ggml_sycl_fattn_alloc        V_f16(fbuf.V);
+    const ggml_sycl_fattn_extra  extra = ggml_sycl_fattn_get_extra(dst);
    ggml_sycl_pool_alloc<int>    KV_max(pool);
    ggml_sycl_pool_alloc<float>  dst_tmp(pool);
    ggml_sycl_pool_alloc<sycl::float2> dst_tmp_meta(pool);
@@ -944,10 +946,11 @@ void launch_fattn(
        const size_t bs = ggml_blck_size(K->type);
        const size_t ts = ggml_type_size(K->type);

-        K_f16.alloc(ggml_nelements(K));
+        sycl::half * K_f16_ptr = extra.K_buffer_ptr ? (sycl::half *) extra.K_buffer_ptr
+                                                    : K_f16.alloc(ggml_nelements(K));
        if (ggml_is_contiguously_allocated(K)) {
            to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(K->type, dst);
-            to_fp16(K_data, K_f16.ptr, ggml_nelements(K), main_stream);
+            to_fp16(K_data, K_f16_ptr, ggml_nelements(K), main_stream);

            nb11 = nb11 * bs * sizeof(sycl::half) / ts;
            nb12 = nb12 * bs * sizeof(sycl::half) / ts;
@@ -958,13 +961,13 @@ void launch_fattn(
            const int64_t s01 = nb11 / ts;
            const int64_t s02 = nb12 / ts;
            const int64_t s03 = nb13 / ts;
-            to_fp16(K_data, K_f16.ptr, K->ne[0], K->ne[1], K->ne[2], K->ne[3], s01, s02, s03, main_stream);
+            to_fp16(K_data, K_f16_ptr, K->ne[0], K->ne[1], K->ne[2], K->ne[3], s01, s02, s03, main_stream);

            nb11 = K->ne[0] * sizeof(sycl::half);
            nb12 = K->ne[1] * nb11;
            nb13 = K->ne[2] * nb12;
        }
-        K_data = (char *) K_f16.ptr;
+        K_data = (char *) K_f16_ptr;
    }

    if (need_f16_V && V->type != GGML_TYPE_F16) {
@@ -977,11 +980,12 @@ void launch_fattn(
            const size_t bs = ggml_blck_size(V->type);
            const size_t ts = ggml_type_size(V->type);

-            V_f16.alloc(ggml_nelements(V));
+            sycl::half * V_f16_ptr = extra.V_buffer_ptr ? (sycl::half *) extra.V_buffer_ptr
+                                                        : V_f16.alloc(ggml_nelements(V));
            if (ggml_is_contiguously_allocated(V)) {
                to_fp16_sycl_t to_fp16 = ggml_get_to_fp16_sycl(V->type, dst);
-                to_fp16(V_data, V_f16.ptr, ggml_nelements(V), main_stream);
-                V_data = (char *) V_f16.ptr;
+                to_fp16(V_data, V_f16_ptr, ggml_nelements(V), main_stream);
+                V_data = (char *) V_f16_ptr;

                nb21 = nb21 * bs * sizeof(sycl::half) / ts;
                nb22 = nb22 * bs * sizeof(sycl::half) / ts;
@@ -992,13 +996,13 @@ void launch_fattn(
                const int64_t s01 = nb21 / ts;
                const int64_t s02 = nb22 / ts;
                const int64_t s03 = nb23 / ts;
-                to_fp16(V_data, V_f16.ptr, V->ne[0], V->ne[1], V->ne[2], V->ne[3], s01, s02, s03, main_stream);
+                to_fp16(V_data, V_f16_ptr, V->ne[0], V->ne[1], V->ne[2], V->ne[3], s01, s02, s03, main_stream);

                nb21 = V->ne[0] * sizeof(sycl::half);
                nb22 = V->ne[1] * nb21;
                nb23 = V->ne[2] * nb22;
            }
-            V_data = (char *) V_f16.ptr;
+            V_data = (char *) V_f16_ptr;
        }
    }

@@ -14,9 +14,21 @@
 // set minimum query length to treat as prefill (32)
 #define GGML_SYCL_FA_ONEDNN_MIN_Q 32

-bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) {
+bool ggml_sycl_fattn_onednn_binds_kv(const ggml_tensor * K, const ggml_tensor * V) {
+    if (K->type != GGML_TYPE_F16 || V->type != GGML_TYPE_F16) {
+        return false;
+    }
+    auto bindable = [](const ggml_tensor * t) {
+        return t->nb[0] == sizeof(sycl::half) && t->nb[1] % sizeof(sycl::half) == 0 &&
+               t->nb[2] % sizeof(sycl::half) == 0 && t->nb[3] % sizeof(sycl::half) == 0;
+    };
+    return bindable(K) && bindable(V);
+}
+
+bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst, bool use_shape_limit) {
 #if !GGML_SYCL_DNNL
    GGML_UNUSED(dst);
+    GGML_UNUSED(use_shape_limit);
    return false;
 #else
    if (!g_ggml_sycl_fa_onednn) {
@@ -44,7 +56,7 @@ bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) {
        if (!k_ok || !v_ok) {
            return false;
        }
-        if (Q->ne[1] < 32 || K->ne[1] < 1024) {
+        if (use_shape_limit && (Q->ne[1] < 32 || K->ne[1] < 1024)) {
            return false;
        }
        for (const ggml_tensor * t : {K, V}) {
@@ -94,7 +106,7 @@ bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst) {
        return false;
    }
    // Prefill only.
-    if (Q->ne[1] < GGML_SYCL_FA_ONEDNN_MIN_Q) {
+    if (use_shape_limit && Q->ne[1] < GGML_SYCL_FA_ONEDNN_MIN_Q) {
        return false;
    }
    return true;
@@ -240,9 +252,16 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
    dnnl::engine    eng    = ctx.engine_dnnl(stream);
    dnnl::stream    strm   = ctx.stream_dnnl(stream);

+    const ggml_sycl_fattn_extra extra = ggml_sycl_fattn_get_extra(dst);
+
    // Q: always f32 -- copy to dense f16.
-    ggml_sycl_pool_alloc<sycl::half> Qf(ctx.pool(), (size_t) H * q * d);
-    cont_to_f16_sycl<float>((const char *) Q->data, Qf.get(), d, q, H, mb, Q->nb[1], Q->nb[2], Q->nb[3], stream);
+    std::optional<ggml_sycl_pool_alloc<sycl::half>> Qf_pool;
+    sycl::half * Qf_ptr = (sycl::half *) extra.Q_buffer_ptr;
+    if (!Qf_ptr) {
+        Qf_pool.emplace(ctx.pool(), (size_t) H * q * d);
+        Qf_ptr = Qf_pool->get();
+    }
+    cont_to_f16_sycl<float>((const char *) Q->data, Qf_ptr, d, q, H, mb, Q->nb[1], Q->nb[2], Q->nb[3], stream);

    // K/V: bind the f16 cache in place. llama.cpp permutes it to [token][head][dim], so its head
    // plane is strided rather than dense, which is what an explicit stride vector expresses.
@@ -253,11 +272,12 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
    std::array<int64_t, 5> v_str = k_str;
    std::optional<ggml_sycl_pool_alloc<sycl::half>> Kf_pool;
    std::optional<ggml_sycl_pool_alloc<sycl::half>> Vf_pool;
+    // Helper: hand out reserved space, or fall back to the pool.
+    auto stage_k = [&](size_t n) { if (extra.K_buffer_ptr) { return (sycl::half *) extra.K_buffer_ptr; }
+                                  Kf_pool.emplace(ctx.pool(), n); return Kf_pool->get(); };
+    auto stage_v = [&](size_t n) { if (extra.V_buffer_ptr) { return (sycl::half *) extra.V_buffer_ptr; }
+                                  Vf_pool.emplace(ctx.pool(), n); return Vf_pool->get(); };

-    auto bindable = [](const ggml_tensor * t) {
-        return t->nb[0] == sizeof(sycl::half) && t->nb[1] % sizeof(sycl::half) == 0 &&
-               t->nb[2] % sizeof(sycl::half) == 0 && t->nb[3] % sizeof(sycl::half) == 0;
-    };
    auto elem_strides = [](const ggml_tensor * t) {
        const int64_t s1 = (int64_t) (t->nb[1] / t->nb[0]);
        const int64_t s2 = (int64_t) (t->nb[2] / t->nb[0]);
@@ -266,22 +286,19 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
        return std::array<int64_t, 5>{ s3, s2, s2, s1, 1 };
    };

-    if (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16 && bindable(K) && bindable(V)) {
+    if (ggml_sycl_fattn_onednn_binds_kv(K, V)) {
        K_ptr = (sycl::half *) K->data;
        V_ptr = (sycl::half *) V->data;
        k_str = elem_strides(K);
        v_str = elem_strides(V);
    } else if (K->type == GGML_TYPE_F16 && V->type == GGML_TYPE_F16) {
-        Kf_pool.emplace(ctx.pool(), (size_t) Hkv * seq * d);
-        Vf_pool.emplace(ctx.pool(), (size_t) Hkv * seq * d);
-        cont_to_f16_sycl<sycl::half>((const char *) K->data, Kf_pool->get(), d, seq, Hkv, mb, K->nb[1], K->nb[2], K->nb[3], stream);
-        cont_to_f16_sycl<sycl::half>((const char *) V->data, Vf_pool->get(), d, seq, Hkv, mb, V->nb[1], V->nb[2], V->nb[3], stream);
-        K_ptr = Kf_pool->get();
-        V_ptr = Vf_pool->get();
+        K_ptr = stage_k((size_t) Hkv * seq * d);
+        V_ptr = stage_v((size_t) Hkv * seq * d);
+        cont_to_f16_sycl<sycl::half>((const char *) K->data, K_ptr, d, seq, Hkv, mb, K->nb[1], K->nb[2], K->nb[3], stream);
+        cont_to_f16_sycl<sycl::half>((const char *) V->data, V_ptr, d, seq, Hkv, mb, V->nb[1], V->nb[2], V->nb[3], stream);
    } else if (ggml_is_quantized(K->type)) {
        // Quantized K/V: dequant to dense F16 using pool, same lifetime as F16 path.
-        Kf_pool.emplace(ctx.pool(), ggml_nelements(K));
-        K_ptr = Kf_pool->get();
+        K_ptr = stage_k((size_t) ggml_nelements(K));
        {
            const char * K_data = (const char *)K->data;
            const bool k_non_dense = ((int64_t)K->ne[1] * K->nb[1] != K->nb[2]) && K->ne[2] > 1;
@@ -315,8 +332,7 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
        // data pointer), their logical values differ because the quantized
        // elements at different positions/offsets represent different K/V
        // data. Master's F16 path also never aliases K and V.
-        Vf_pool.emplace(ctx.pool(), ggml_nelements(V));
-        V_ptr = Vf_pool->get();
+        V_ptr = stage_v((size_t) ggml_nelements(V));
        {
            const char * V_data = (const char *)V->data;
            const bool v_non_dense = ((int64_t)V->ne[1] * V->nb[1] != V->nb[2]) && V->ne[2] > 1;
@@ -347,12 +363,10 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
        }
    } else {
        // F32: strided copy to dense F16 via cont_to_f16_sycl<float>.
-        Kf_pool.emplace(ctx.pool(), ggml_nelements(K));
-        K_ptr = Kf_pool->get();
+        K_ptr = stage_k((size_t) ggml_nelements(K));
        cont_to_f16_sycl<float>((const char *) K->data, K_ptr, K->ne[0], K->ne[1], K->ne[2], K->ne[3],
                                K->nb[1], K->nb[2], K->nb[3], stream);
-        Vf_pool.emplace(ctx.pool(), ggml_nelements(V));
-        V_ptr = Vf_pool->get();
+        V_ptr = stage_v((size_t) ggml_nelements(V));
        cont_to_f16_sycl<float>((const char *) V->data, V_ptr, V->ne[0], V->ne[1], V->ne[2], V->ne[3],
                                V->nb[1], V->nb[2], V->nb[3], stream);
    }
@@ -366,11 +380,21 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
    // instead -- the value is captured into the command, so no host memory has to outlive the
    // call, and the enqueue stays async.
    const sycl::half scale_h = (sycl::half) (1.0f / kq_scale);
-    ggml_sycl_pool_alloc<sycl::half> scbuf(ctx.pool(), 1);
-    sycl::half * const scale_dev = scbuf.get();
+    std::optional<ggml_sycl_pool_alloc<sycl::half>> scbuf;
+    sycl::half * scale_dev = (sycl::half *) extra.scale_buffer_ptr;
+    if (!scale_dev) {
+        scbuf.emplace(ctx.pool(), 1);
+        scale_dev = scbuf->get();
+    }
    stream->single_task([=]() { *scale_dev = scale_h; });

-    ggml_sycl_pool_alloc<sycl::half> outf(ctx.pool(), (size_t) H * q * d);   // f16 contiguous SDPA out [mb,H,q,d]
+    // f16 contiguous SDPA out [mb,H,q,d]
+    std::optional<ggml_sycl_pool_alloc<sycl::half>> outf_pool;
+    sycl::half * outf_ptr = (sycl::half *) extra.out_buffer_ptr;
+    if (!outf_ptr) {
+        outf_pool.emplace(ctx.pool(), (size_t) H * q * d);
+        outf_ptr = outf_pool->get();
+    }

    // compile once per (device, shape, KV strides), reuse across layers/calls. Stride 2 always
    // repeats stride 1 and stride 4 is always 1, so the key covers every entry that can differ.
@@ -392,7 +416,7 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
    }

    auto id2ptr = [&](size_t r) -> void * {
-        if (r == E.id_q)     return Qf.get();
+        if (r == E.id_q)     return Qf_ptr;
        if (r == E.id_k)     return K_ptr;
        if (r == E.id_v)     return V_ptr;
        if (r == E.id_scale) return scale_dev;
@@ -404,10 +428,10 @@ void ggml_sycl_flash_attn_ext_onednn(ggml_backend_sycl_context & ctx, ggml_tenso
    for (auto & lt : E.ins) {
        ti.emplace_back(lt, eng, id2ptr(lt.get_id()));
    }
-    tensor to(E.out, eng, outf.get());
+    tensor to(E.out, eng, outf_ptr);
    E.cp.execute(strm, ti, {to});

-    permute_sdpa_out_sycl(outf.get(), (float *) dst->data, mb, H, q, d, stream);
+    permute_sdpa_out_sycl(outf_ptr, (float *) dst->data, mb, H, q, d, stream);
    // Single device needs no sync: the dnnl stream wraps this same in-order queue, so the SDPA
    // serializes with the staging kernels before it and the permute/pool reuse after it. The
    // garbage output formerly blamed on the missing sync here was the scale use-after-return
@@ -5,7 +5,11 @@

 // Static-only check: fused-XMX oneDNN Graph SDPA path==flash-attn op
 // (f16 KV, no softcap/ALiBi, single stream, tuned head_dim, prefill-sized q.)
-bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst);
+bool ggml_sycl_flash_attn_ext_onednn_supported(const ggml_tensor * dst, bool use_shape_limit = true);
+
+// True when the oneDNN path binds an F16 KV cache in place instead of staging a dense copy of
+// it. Depends only on the types and strides of K and V, so the answer holds for every call.
+bool ggml_sycl_fattn_onednn_binds_kv(const ggml_tensor * K, const ggml_tensor * V);

 // Run flash attention through oneDNN's fused xmx SDPA
 // execute the cached SDPA partition, write the f32 dst. Falls back to the TILE kernel on any failure.
@@ -378,3 +378,76 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst
 bool ggml_sycl_flash_attn_ext_supported(int device, const ggml_tensor * dst) {
    return ggml_sycl_get_best_fattn_kernel(device, dst) != BEST_FATTN_KERNEL_NONE;
 }
+
+static uintptr_t ggml_sycl_fattn_reserve_halves(ggml_sycl_fattn_extra & extra, size_t n_halves) {
+    if (n_halves == 0) {
+        return 0;
+    }
+    extra.end = GGML_PAD(extra.end, SYCL_BUFFER_ALIGNMENT);
+    const uintptr_t block = extra.end;
+    extra.end += n_halves * sizeof(sycl::half);
+    return block;
+}
+
+ggml_sycl_fattn_extra ggml_sycl_fattn_get_extra(const ggml_tensor * dst) {
+    ggml_sycl_fattn_extra extra;
+
+    extra.end = (uintptr_t) dst->data + ggml_nbytes(dst);
+
+    if (dst->op != GGML_OP_FLASH_ATTN_EXT) {
+        return extra;
+    }
+
+    const ggml_tensor * Q = dst->src[0];
+    const ggml_tensor * K = dst->src[1];
+    const ggml_tensor * V = dst->src[2];
+    if (!Q || !K || !V) {
+        return extra;
+    }
+
+    const int64_t d = K->ne[0];
+    const int64_t H = Q->ne[2];
+    const int64_t q = Q->ne[1];
+
+    // calculate the worst-case memory consumption across all kernels
+    const bool onednn_supported = ggml_sycl_flash_attn_ext_onednn_supported(dst, /* use_shape_limit */ false);
+
+    const bool tile_needs_K = K->type != GGML_TYPE_F16;
+    const bool tile_needs_V = V->type != GGML_TYPE_F16;
+
+    const bool V_is_K_view = V->view_src &&
+        (V->view_src == K || (V->view_src == K->view_src && V->view_offs == K->view_offs));
+
+    size_t need_K = 0, need_V = 0, need_Q = 0, need_out = 0, need_scale = 0;
+    if (onednn_supported) {
+        need_Q     = (size_t) H * q * d;
+        need_out   = (size_t) H * q * d;
+        need_scale = 1;
+        // an f16 cache is bound in place, so it needs no staging copy
+        if (!ggml_sycl_fattn_onednn_binds_kv(K, V)) {
+            need_K = (size_t) ggml_nelements(K);
+            need_V = (size_t) ggml_nelements(V);
+        }
+    }
+    if (tile_needs_K) {
+        need_K = std::max(need_K, (size_t) ggml_nelements(K));
+    }
+    if (tile_needs_V) {
+        need_V = std::max(need_V, (size_t) ggml_nelements(V));
+    }
+
+    extra.Q_buffer_ptr = ggml_sycl_fattn_reserve_halves(extra, need_Q);
+    extra.K_buffer_ptr = ggml_sycl_fattn_reserve_halves(extra, need_K);
+    extra.V_buffer_ptr = (V_is_K_view && !onednn_supported && need_V)
+                       ? extra.K_buffer_ptr
+                       : ggml_sycl_fattn_reserve_halves(extra, need_V);
+    extra.scale_buffer_ptr = ggml_sycl_fattn_reserve_halves(extra, need_scale);
+    extra.out_buffer_ptr   = ggml_sycl_fattn_reserve_halves(extra, need_out);
+
+    return extra;
+}
+
+size_t ggml_sycl_flash_attn_ext_get_alloc_size(const ggml_tensor * dst) {
+    const ggml_sycl_fattn_extra extra = ggml_sycl_fattn_get_extra(dst);
+    return (size_t) (extra.end - (uintptr_t) dst->data);
+}
@@ -19,6 +19,24 @@ void ggml_sycl_flash_attn_ext(ggml_backend_sycl_context & ctx, ggml_tensor * dst

 bool ggml_sycl_flash_attn_ext_supported(int device, const ggml_tensor * dst);

+// Scratch that flash attention needs beyond the output tensor
+struct ggml_sycl_fattn_extra {
+    uintptr_t K_buffer_ptr     = 0;   // F16 copy of the K cache
+    uintptr_t V_buffer_ptr     = 0;   // F16 copy of the V cache
+    uintptr_t Q_buffer_ptr     = 0;   // dense F16 copy of Q, oneDNN only
+    uintptr_t scale_buffer_ptr = 0;   // the softmax scale as an F16 scalar, oneDNN only
+    uintptr_t out_buffer_ptr   = 0;   // F16 SDPA output before conversion to F32, oneDNN only
+    uintptr_t end              = 0;   // one past the last reserved byte; sizes the allocation
+};
+
+// ggml_sycl_fattn_get_extra() is the single source of truth for the layout: it both sizes
+// the reservation and hands out the pointers, so the two cannot disagree.
+// Each field is the address of one reserved block, or 0 if that block was not reserved,
+// in which case the caller allocates from the scratch pool instead.
+ggml_sycl_fattn_extra ggml_sycl_fattn_get_extra(const ggml_tensor * dst);
+
+size_t ggml_sycl_flash_attn_ext_get_alloc_size(const ggml_tensor * dst);
+
 void ggml_sycl_flash_attn_ext_mkl(ggml_backend_sycl_context & ctx, ggml_tensor * dst);

 #endif // GGML_SYCL_FATTN_HPP
@@ -955,7 +955,10 @@ static size_t ggml_backend_sycl_buffer_type_get_max_size(ggml_backend_buffer_typ
 }

 static size_t ggml_backend_sycl_buffer_type_get_alloc_size(ggml_backend_buffer_type_t buft, const ggml_tensor * tensor) {
-    size_t size = ggml_nbytes(tensor);
+    // Reserve the additional scratch so it's visible to the graph allocator
+    size_t size = tensor->op == GGML_OP_FLASH_ATTN_EXT
+        ? ggml_sycl_flash_attn_ext_get_alloc_size(tensor)
+        : ggml_nbytes(tensor);
    int64_t ne0 = tensor->ne[0];

    if (ggml_is_quantized(tensor->type)) {
@@ -1347,7 +1347,6 @@ struct vk_mat_mat_id_push_constants {
    uint32_t stride_a; uint32_t stride_b; uint32_t stride_d;
    uint32_t batch_stride_a; uint32_t batch_stride_b; uint32_t batch_stride_d;
    uint32_t nei0; uint32_t nei1; uint32_t nbi1; uint32_t ne11;
-    uint32_t padded_N;
    uint32_t n_experts;
    uint32_t hoist_row_ids;
 };
@@ -2403,9 +2402,8 @@ struct ggml_backend_vk_context {
    // Cache most recent tensor that was converted into prealloc_y, and what pipeline it used to convert.
    vk_pipeline_struct * prealloc_y_last_pipeline_used {};
    const ggml_tensor * prealloc_y_last_tensor_used {};
-    // True when prealloc_y holds the padded fp16 layout used by the coopmat2 B decode-vector callback.
-    // If false, then it's contiguous.
-    bool prealloc_y_last_decode_vector_staging {};
+    // True when the K dimension in prealloc_y is padded.
+    bool prealloc_y_last_k_padded {};

    // Track which nodes have been used since the last sync, and whether they were written to
    std::vector<const ggml_tensor *> unsynced_nodes_written;
@@ -8984,13 +8982,13 @@ static void ggml_vk_matmul_id(
        uint32_t m, uint32_t n, uint32_t k, uint32_t stride_a, uint32_t stride_b, uint32_t stride_d,
        uint32_t batch_stride_a, uint32_t batch_stride_b, uint32_t batch_stride_d,
        uint32_t n_as, uint32_t nei0, uint32_t nei1, uint32_t nbi1, uint32_t ne11,
-        uint32_t padded_n, bool hoist_row_ids) {
+        bool hoist_row_ids) {
    VK_LOG_DEBUG("ggml_vk_matmul_id(a: (" << a.buffer->buffer << ", " << a.offset << ", " << a.size << "), b: (" << b.buffer->buffer << ", " << b.offset << ", " << b.size << "), d: (" << d.buffer->buffer << ", " << d.offset << ", " << d.size << "), ids: (" << ids.buffer->buffer << ", " << ids.offset << ", " << ids.size << "), expert_count: (" << expert_count_buf.buffer->buffer << ", " << expert_count_buf.offset << ", " << expert_count_buf.size << "), " <<
        "m: " << m << ", n: " << n << ", k: " << k << ", stride_a: " << stride_a << ", stride_b: " << stride_b << ", stride_d: " << stride_d << ", " <<
        "batch_stride_a: " << batch_stride_a << ", batch_stride_b: " << batch_stride_b << ", batch_stride_d: " << batch_stride_d << ", " <<
        "n_as: " << n_as << ", nei0: " << nei0 << ", nei1: " << nei1 << ", nbi1: " << nbi1 << ", ne11: " << ne11 << ")");
    const vk_mat_mat_id_push_constants pc = { m, n, k, stride_a, stride_b, stride_d, batch_stride_a, batch_stride_b, batch_stride_d,
-                                              nei0, nei1, nbi1, ne11, padded_n, n_as, uint32_t(hoist_row_ids) };
+                                              nei0, nei1, nbi1, ne11, n_as, uint32_t(hoist_row_ids) };
    ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { a, b, d, ids, expert_count_buf }, pc, { m, nei1, n_as });
 }

@@ -9455,27 +9453,27 @@ static void ggml_vk_mul_mat_q_f16(ggml_backend_vk_context * ctx, vk_context& sub
    if (y_non_contig) {
        if (ctx->prealloc_y_last_pipeline_used != to_fp16_vk_1.get() ||
            ctx->prealloc_y_last_tensor_used != src1 ||
-            ctx->prealloc_y_last_decode_vector_staging) {
+            ctx->prealloc_y_last_k_padded) {
            if (ctx->prealloc_y_need_sync) {
                ggml_vk_sync_buffers(ctx, subctx);
            }
            ggml_vk_cpy_to_contiguous(ctx, subctx, to_fp16_vk_1, src1, ggml_vk_subbuffer(ctx, d_Qy, qy_buf_offset), ggml_vk_subbuffer(ctx, d_Y, 0));
            ctx->prealloc_y_last_pipeline_used = to_fp16_vk_1.get();
            ctx->prealloc_y_last_tensor_used = src1;
-            ctx->prealloc_y_last_decode_vector_staging = false;
+            ctx->prealloc_y_last_k_padded = false;
        }
    }
    if (quantize_y) {
        if (ctx->prealloc_y_last_pipeline_used != to_q8_1.get() ||
            ctx->prealloc_y_last_tensor_used != src1 ||
-            ctx->prealloc_y_last_decode_vector_staging) {
+            ctx->prealloc_y_last_k_padded) {
            if (ctx->prealloc_y_need_sync) {
                ggml_vk_sync_buffers(ctx, subctx);
            }
            ggml_vk_quantize_q8_1(ctx, subctx, ggml_vk_subbuffer(ctx, d_Qy, qy_buf_offset), ggml_vk_subbuffer(ctx, d_Y, 0), y_ne);
            ctx->prealloc_y_last_pipeline_used = to_q8_1.get();
            ctx->prealloc_y_last_tensor_used = src1;
-            ctx->prealloc_y_last_decode_vector_staging = false;
+            ctx->prealloc_y_last_k_padded = false;
        }
    }

@@ -9734,27 +9732,27 @@ static void ggml_vk_mul_mat_vec_q_f16(ggml_backend_vk_context * ctx, vk_context&
        GGML_ASSERT(y_sz == ggml_type_size(src1->type) * y_ne);
        if (ctx->prealloc_y_last_pipeline_used != to_fp16_vk_1.get() ||
            ctx->prealloc_y_last_tensor_used != src1 ||
-            ctx->prealloc_y_last_decode_vector_staging) {
+            ctx->prealloc_y_last_k_padded) {
            if (ctx->prealloc_y_need_sync) {
                ggml_vk_sync_buffers(ctx, subctx);
            }
            ggml_vk_cpy_to_contiguous(ctx, subctx, to_fp16_vk_1, src1, d_Qy, d_Y);
            ctx->prealloc_y_last_pipeline_used = to_fp16_vk_1.get();
            ctx->prealloc_y_last_tensor_used = src1;
-            ctx->prealloc_y_last_decode_vector_staging = false;
+            ctx->prealloc_y_last_k_padded = false;
        }
    }
    if (quantize_y) {
        if (ctx->prealloc_y_last_pipeline_used != to_q8_1.get() ||
            ctx->prealloc_y_last_tensor_used != src1 ||
-            ctx->prealloc_y_last_decode_vector_staging) {
+            ctx->prealloc_y_last_k_padded) {
            if (ctx->prealloc_y_need_sync) {
                ggml_vk_sync_buffers(ctx, subctx);
            }
            ggml_vk_quantize_q8_1(ctx, subctx, d_Qy, d_Y, y_ne);
            ctx->prealloc_y_last_pipeline_used = to_q8_1.get();
            ctx->prealloc_y_last_tensor_used = src1;
-            ctx->prealloc_y_last_decode_vector_staging = false;
+            ctx->prealloc_y_last_k_padded = false;
        }
    }

@@ -10234,8 +10232,6 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
                              (src0->type == GGML_TYPE_BF16 && src1->type != GGML_TYPE_BF16) ||
                              !ggml_vk_dim01_contiguous(src1);

-    const uint32_t y_staged_row_stride = y_decode_vector_staging ? (uint32_t)ggml_vk_align_size(ne10, 4) : (uint32_t)ne10;
-
    const bool y_f32_kernel = src1->type == GGML_TYPE_F32 && !y_non_contig;

    bool quantize_y = ctx->device->integer_dot_product && src1->type == GGML_TYPE_F32 && ggml_is_contiguous(src1) && !y_non_contig && (ne11 * ne10) % 4 == 0;
@@ -10250,19 +10246,25 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
    }

    const bool qx_needs_dequant = mmp == nullptr || x_non_contig;
-    const bool qy_needs_dequant = !quantize_y && ((src1->type != f16_type && !y_f32_kernel) || y_non_contig);
+    bool qy_needs_dequant = !quantize_y && ((src1->type != f16_type && !y_f32_kernel) || y_non_contig);

    if (qx_needs_dequant) {
        // Fall back to dequant + f16 mulmat
        mmp = ggml_vk_get_mul_mat_mat_id_pipeline(ctx, f16_type, y_f32_kernel ? GGML_TYPE_F32 : f16_type, (ggml_prec)dst->op_params[0]);
    }

-    // Not implemented
-    GGML_ASSERT(y_non_contig || !qy_needs_dequant);  // NOLINT
-
    const ggml_type effective_src1_type = quantize_y ? GGML_TYPE_Q8_1 : (y_f32_kernel ? GGML_TYPE_F32 : src1->type);

    const uint32_t kpad = quantize_y ? 0 : ggml_vk_align_size(ne10, ggml_vk_guess_matmul_id_pipeline_align(ctx, mmp, ne01, nei1, qx_needs_dequant ? f16_type : src0->type, effective_src1_type));
+    // Coopmat2 MUL_MAT_ID BK specialization constants in ggml_vk_load_shaders are at most 64.
+    const uint32_t y_staged_row_stride = ctx->device->coopmat2 && !quantize_y ? ggml_vk_align_size(ne10, 64) : ne10;
+    const bool y_needs_k_padding = ne10 != y_staged_row_stride;
+    const bool y_needs_reformat = y_non_contig || y_needs_k_padding;
+    qy_needs_dequant = qy_needs_dequant || y_needs_k_padding;
+
+    // Not implemented
+    GGML_ASSERT(y_needs_reformat || !qy_needs_dequant);  // NOLINT
+
    const bool aligned = !quantize_y && ne10 == kpad && ne01 > 8 && nei1 > 8;

    vk_pipeline pipeline = ggml_vk_guess_matmul_id_pipeline(ctx, mmp, ne01, nei1, aligned, qx_needs_dequant ? f16_type : src0->type, effective_src1_type);
@@ -10270,10 +10272,8 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
    if (ggml_nbytes(src0) > ctx->device->properties.limits.maxStorageBufferRange) {
        pipeline = ggml_vk_get_64b_indexing_pipeline(ctx, pipeline);
    }
-    // Reserve extra storage in the N dimension for the Y matrix, so we can avoid bounds-checking
-    uint32_t padded_n = qy_needs_dequant ? ROUNDUP_POW2(ne11, pipeline->wg_denoms[1]) :ne11;
    const uint64_t x_ne = ggml_nelements(src0);
-    const uint64_t y_ne = (uint64_t)y_staged_row_stride * padded_n * ne12 * ne13;
+    const uint64_t y_ne = (uint64_t)y_staged_row_stride * ne11 * ne12 * ne13;
    const uint64_t d_ne = ggml_nelements(dst);

    const uint64_t qx_sz = ggml_type_size(src0->type) * x_ne / ggml_blck_size(src0->type);
@@ -10292,7 +10292,7 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
        y_staged_dst.type = f16_type;
        y_staged_dst.nb[0] = ggml_type_size(f16_type);
        y_staged_dst.nb[1] = y_staged_dst.nb[0] * y_staged_row_stride;
-        y_staged_dst.nb[2] = y_staged_dst.nb[1] * padded_n;
+        y_staged_dst.nb[2] = y_staged_dst.nb[1] * ne11;
        y_staged_dst.nb[3] = y_staged_dst.nb[2] * y_staged_dst.ne[2];
        return y_staged_dst;
    };
@@ -10302,10 +10302,10 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
    } else {
        to_fp16_vk_0 = ggml_vk_get_to_fp16(ctx, src0->type);
    }
-    if (y_non_contig) {
+    if (y_needs_reformat) {
        ggml_tensor y_staged_dst;
        const ggml_tensor * y_staged_dst_ptr = nullptr;
-        if (y_decode_vector_staging) {
+        if (y_needs_k_padding) {
            y_staged_dst = make_y_staged_dst();
            y_staged_dst_ptr = &y_staged_dst;
        }
@@ -10432,14 +10432,18 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
        ggml_vk_dispatch_pipeline(ctx, subctx, to_fp16_vk_0,
            { vk_subbuffer{ d_Qx, qx_buf_offset, qx_sz }, vk_subbuffer{ d_X, 0, x_sz } }, pc, { (uint32_t)x_ne, 1, 1});
    }
-    if (y_non_contig) {
+    if (y_needs_reformat) {
        if (ctx->prealloc_y_last_pipeline_used != to_fp16_vk_1.get() ||
            ctx->prealloc_y_last_tensor_used != src1 ||
-            ctx->prealloc_y_last_decode_vector_staging != y_decode_vector_staging) {
+            ctx->prealloc_y_last_k_padded != y_needs_k_padding) {
            if (ctx->prealloc_y_need_sync) {
                ggml_vk_sync_buffers(ctx, subctx);
            }
-            if (y_decode_vector_staging) {
+            if (y_needs_k_padding) {
+                GGML_ASSERT(y_sz % 4 == 0);
+                // Zero B padding because clamping only A can produce 0 * Inf or NaN.
+                subctx->s->buffer->buf.fillBuffer(d_Y->buffer, 0, y_sz, 0);
+                ggml_vk_sync_buffers(ctx, subctx);
                const ggml_tensor y_staged_dst = make_y_staged_dst();
                const uint32_t y_staged_dst_type_size = ggml_type_size(y_staged_dst.type);
                ggml_vk_cpy_to_strided(
@@ -10454,27 +10458,27 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
            }
            ctx->prealloc_y_last_pipeline_used = to_fp16_vk_1.get();
            ctx->prealloc_y_last_tensor_used = src1;
-            ctx->prealloc_y_last_decode_vector_staging = y_decode_vector_staging;
+            ctx->prealloc_y_last_k_padded = y_needs_k_padding;
        }
    }
    if (quantize_y) {
        if (ctx->prealloc_y_last_pipeline_used != to_q8_1.get() ||
            ctx->prealloc_y_last_tensor_used != src1 ||
-            ctx->prealloc_y_last_decode_vector_staging) {
+            ctx->prealloc_y_last_k_padded) {
            if (ctx->prealloc_y_need_sync) {
                ggml_vk_sync_buffers(ctx, subctx);
            }
            ggml_vk_quantize_q8_1(ctx, subctx, ggml_vk_subbuffer(ctx, d_Qy, qy_buf_offset), ggml_vk_subbuffer(ctx, d_Y, 0), y_ne);
            ctx->prealloc_y_last_pipeline_used = to_q8_1.get();
            ctx->prealloc_y_last_tensor_used = src1;
-            ctx->prealloc_y_last_decode_vector_staging = false;
+            ctx->prealloc_y_last_k_padded = false;
        }
    }
    ggml_vk_sync_buffers(ctx, subctx);

    uint32_t stride_batch_x = ne00*ne01;
-    uint32_t stride_b_y = y_decode_vector_staging ? y_staged_row_stride : ne10;
-    uint32_t stride_batch_y = y_decode_vector_staging ? y_staged_row_stride * padded_n : ne10*ne11;
+    uint32_t stride_b_y = y_needs_k_padding ? y_staged_row_stride : ne10;
+    uint32_t stride_batch_y = y_needs_k_padding ? y_staged_row_stride * ne11 : ne10*ne11;

    if (!ggml_vk_dim01_contiguous(src0) && !qx_needs_dequant) {
        stride_batch_x = src0->nb[0] / ggml_type_size(src0->type);
@@ -10491,13 +10495,13 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
        { d_D, d_buf_offset, d_sz }, { d_ids, ids_buf_offset, ids_sz }, expert_count_buf,
        ne01, ne21, ne10, ne10, stride_b_y, ne01,
        stride_batch_x, stride_batch_y, ne20*ne21,
-        n_as, nei0, nei1, nbi1 / ggml_type_size(ids->type), ne11, padded_n, hoist_row_ids
+        n_as, nei0, nei1, nbi1 / ggml_type_size(ids->type), ne11, hoist_row_ids
    );  // NOLINT

    if (x_non_contig || qx_needs_dequant) {
        ctx->prealloc_x_need_sync = true;
    }
-    if (y_non_contig || quantize_y) {
+    if (y_needs_reformat || quantize_y) {
        ctx->prealloc_y_need_sync = true;
    }
    ctx->prealloc_split_k_need_sync = true;
@@ -10648,27 +10652,27 @@ static void ggml_vk_mul_mat_vec_id_q_f16(ggml_backend_vk_context * ctx, vk_conte
        GGML_ASSERT(y_sz == ggml_type_size(src1->type) * y_ne);
        if (ctx->prealloc_y_last_pipeline_used != to_fp16_vk_1.get() ||
            ctx->prealloc_y_last_tensor_used != src1 ||
-            ctx->prealloc_y_last_decode_vector_staging) {
+            ctx->prealloc_y_last_k_padded) {
            if (ctx->prealloc_y_need_sync) {
                ggml_vk_sync_buffers(ctx, subctx);
            }
            ggml_vk_cpy_to_contiguous(ctx, subctx, to_fp16_vk_1, src1, d_Qy, d_Y);
            ctx->prealloc_y_last_pipeline_used = to_fp16_vk_1.get();
            ctx->prealloc_y_last_tensor_used = src1;
-            ctx->prealloc_y_last_decode_vector_staging = false;
+            ctx->prealloc_y_last_k_padded = false;
        }
    }
    if (quantize_y) {
        if (ctx->prealloc_y_last_pipeline_used != to_q8_1.get() ||
            ctx->prealloc_y_last_tensor_used != src1 ||
-            ctx->prealloc_y_last_decode_vector_staging) {
+            ctx->prealloc_y_last_k_padded) {
            if (ctx->prealloc_y_need_sync) {
                ggml_vk_sync_buffers(ctx, subctx);
            }
            ggml_vk_quantize_q8_1(ctx, subctx, d_Qy, d_Y, y_ne);
            ctx->prealloc_y_last_pipeline_used = to_q8_1.get();
            ctx->prealloc_y_last_tensor_used = src1;
-            ctx->prealloc_y_last_decode_vector_staging = false;
+            ctx->prealloc_y_last_k_padded = false;
        }
    }

@@ -15520,7 +15524,7 @@ static void ggml_vk_preallocate_buffers(ggml_backend_vk_context * ctx, vk_contex
        ctx->prealloc_y = ggml_vk_create_buffer_device(ctx->device, ctx->prealloc_size_y);
        ctx->prealloc_y_last_pipeline_used = nullptr;
        ctx->prealloc_y_last_tensor_used = nullptr;
-        ctx->prealloc_y_last_decode_vector_staging = false;
+        ctx->prealloc_y_last_k_padded = false;
    }
    if (ctx->prealloc_split_k == nullptr || (ctx->prealloc_size_split_k > 0 && ctx->prealloc_split_k->size < ctx->prealloc_size_split_k)) {
        VK_LOG_MEMORY("ggml_vk_preallocate_buffers(split_k_size: " << ctx->prealloc_size_split_k << ")");
@@ -16145,7 +16149,7 @@ static void ggml_vk_graph_cleanup(ggml_backend_vk_context * ctx) {
    VK_LOG_DEBUG("ggml_vk_graph_cleanup()");
    ctx->prealloc_y_last_pipeline_used = {};
    ctx->prealloc_y_last_tensor_used = nullptr;
-    ctx->prealloc_y_last_decode_vector_staging = false;
+    ctx->prealloc_y_last_k_padded = false;

    ctx->unsynced_nodes_written.clear();
    ctx->unsynced_nodes_read.clear();
@@ -16197,7 +16201,7 @@ static void ggml_vk_cleanup(ggml_backend_vk_context * ctx) {

    ctx->prealloc_y_last_pipeline_used = nullptr;
    ctx->prealloc_y_last_tensor_used = nullptr;
-    ctx->prealloc_y_last_decode_vector_staging = false;
+    ctx->prealloc_y_last_k_padded = false;

    ctx->prealloc_size_x = 0;
    ctx->prealloc_size_y = 0;
@@ -17395,7 +17399,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg

    ctx->prealloc_y_last_pipeline_used = nullptr;
    ctx->prealloc_y_last_tensor_used = nullptr;
-    ctx->prealloc_y_last_decode_vector_staging = false;
+    ctx->prealloc_y_last_k_padded = false;

    if (ctx->prealloc_size_add_rms_partials) {
        ggml_vk_preallocate_buffers(ctx, nullptr);
@@ -17800,20 +17804,32 @@ static void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph *
        return;
    }

-    auto const &is_empty = [](ggml_tensor * node) -> bool {
+    auto const &is_empty = [](const ggml_tensor * node) -> bool {
        return node->op == GGML_OP_NONE || node->op == GGML_OP_RESHAPE || node->op == GGML_OP_TRANSPOSE || node->op == GGML_OP_VIEW || node->op == GGML_OP_PERMUTE;
    };

-    auto const &is_src_of = [](const ggml_tensor *dst, const ggml_tensor *src) -> bool {
+    auto const &is_src_of = [&is_empty](const ggml_tensor *dst, const ggml_tensor *src) -> bool {
+        auto const &base = [](const ggml_tensor * tensor) {
+            return tensor->view_src ? tensor->view_src : tensor;
+        };
        for (uint32_t s = 0; s < GGML_MAX_SRC; ++s) {
            if (dst->src[s] == src) {
                return true;
            }
+            if (is_empty(dst) || is_empty(src)) {
+                continue;
+            }
+            // A source view of dst may read storage written through a different view by src.
+            if (dst->src[s] && base(dst->src[s]) == base(src)) {
+                return true;
+            }
+            // Moving dst forward may overwrite storage still read through a view by src.
+            if (src->src[s] && base(dst) == base(src->src[s])) {
+                return true;
+            }
        }
        // implicit dependency if they view the same tensor
-        const ggml_tensor *dst2 = dst->view_src ? dst->view_src : dst;
-        const ggml_tensor *src2 = src->view_src ? src->view_src : src;
-        if (dst2 == src2) {
+        if (base(dst) == base(src)) {
            return true;
        }
        return false;
@@ -19,6 +19,7 @@
 #endif

 #include "types.glsl"
+#include "utils.glsl"

 // shape notation: [dim(N), ..., dim(0)] -- stride(dim(j)) >= stride(dim(i)) if i > j
 layout(binding = 0) readonly buffer A {
@@ -193,14 +194,6 @@ uint32_t       Br    = tid / BS_NPQ;
 uint32_t       Bc    = tid % BS_NPQ;
 const uint32_t BrpWg = WG_SIZE / BS_NPQ;

-// see init_fastdiv_values in ggml-vulkan.cpp
-uint fastdiv(uint n, uint mp, uint L) {
-    uint msbs, lsbs;
-    // msbs = mulhi(n, mp)
-    umulExtended(n, mp, msbs, lsbs);
-    return (msbs + n) >> L;
-}
-
 #ifdef COOPMAT2
 #define ACC_TYPE float16_t

@@ -15,6 +15,7 @@
 #endif

 #include "types.glsl"
+#include "utils.glsl"

 // shape notation: [dim(N), ..., dim(0)] -- stride(dim(j)) >= stride(dim(i)) if i > j
 layout(binding = 0) readonly buffer A {
@@ -178,14 +179,6 @@ uint32_t       Br    = tid / BS_NPQ;
 uint32_t       Bc    = tid % BS_NPQ;
 const uint32_t BrpWg = WG_SIZE / BS_NPQ;

-// see init_fastdiv_values in ggml-vulkan.cpp
-uint fastdiv(uint n, uint mp, uint L) {
-    uint msbs, lsbs;
-    // msbs = mulhi(n, mp)
-    umulExtended(n, mp, msbs, lsbs);
-    return (msbs + n) >> L;
-}
-
 void split_crs(uint32_t crs_idx, out uint32_t ic, out uint32_t kd, out uint32_t kh, out uint32_t kw) {
    const uint32_t KHKW = KH * KW;
    const uint32_t KDKHKW = KD * KHKW;
@@ -8,6 +8,7 @@
 #endif

 #include "types.glsl"
+#include "utils.glsl"

 layout (push_constant) uniform parameter
 {
@@ -33,14 +34,6 @@ shared uint vals[BLOCK_SIZE];
 shared uint offsets[BLOCK_SIZE];
 shared uint cursors[BLOCK_SIZE];

-// see init_fastdiv_values in ggml-vulkan.cpp
-uint fastdiv(uint n, uint mp, uint L) {
-    uint msbs, lsbs;
-    // msbs = mulhi(n, mp)
-    umulExtended(n, mp, msbs, lsbs);
-    return (msbs + n) >> L;
-}
-
 // data_d layout when p.hoist_row_ids is set:
 //   [0,              n_experts)   per-expert row count
 //   [n_experts,    2*n_experts)   per-expert start offset into the row id region
@@ -1,6 +1,8 @@
 #extension GL_EXT_shader_16bit_storage : require
 #extension GL_EXT_control_flow_attributes : require

+#include "utils.glsl"
+
 layout (push_constant) uniform parameter
 {
    uint ne;
@@ -32,18 +34,6 @@ uint get_idx() {
 uint get_aoffset() { return p.misalign_offsets >> 16; }
 uint get_doffset() { return p.misalign_offsets & 0xFFFF; }

-// see init_fastdiv_values in ggml-vulkan.cpp
-uint fastdiv(uint n, uint mp, uint L) {
-    uint msbs, lsbs;
-    // msbs = mulhi(n, mp)
-    umulExtended(n, mp, msbs, lsbs);
-    return (msbs + n) >> L;
-}
-
-uint fastdiv_L(uint packed, uint slot) {
-    return (packed >> (slot * 8)) & 0x3Fu;
-}
-
 uint src0_idx(uint idx) {
    const uint i03 = fastdiv(idx, p.ne0_012mp, fastdiv_L(p.ne0_Ls, 0));
    const uint i03_offset = i03 * p.ne02*p.ne01*p.ne00;
@@ -1,5 +1,7 @@
 #extension GL_EXT_shader_16bit_storage : require

+#include "utils.glsl"
+

 layout(local_size_x = 512, local_size_y = 1, local_size_z = 1) in;

@@ -39,9 +41,3 @@ uint get_aoffset() { return p.misalign_offsets >> 16; }
 uint get_boffset() { return (p.misalign_offsets >> 8) & 0xFF; }
 uint get_doffset() { return p.misalign_offsets & 0xFF; }

-// see init_fastdiv_values in ggml-vulkan.cpp
-uint fastdiv(uint n, uint mp, uint L) {
-    uint msbs, lsbs;
-    umulExtended(n, mp, msbs, lsbs);
-    return (msbs + n) >> L;
-}
@@ -88,7 +88,6 @@ layout (push_constant) uniform parameter
    uint nei1;
    uint nbi1;
    uint ne11;
-    uint padded_N;
    uint n_experts;
    uint hoist_row_ids;
 #else
@@ -56,6 +56,8 @@ layout (push_constant) uniform parameter
    uint nei1;
    uint nbi1;
    uint ne11;
+    uint n_experts;
+    uint hoist_row_ids;
 #else
    uint base_work_group_z;
    uint num_batches;
@@ -64,12 +66,8 @@ layout (push_constant) uniform parameter
    uint ne12;
    uint broadcast2;
    uint broadcast3;
-#endif
    // N dimension for the B matrix can be >= p.N
    uint padded_N;
-#ifdef MUL_MAT_ID
-    uint n_experts;
-    uint hoist_row_ids;
 #endif
 } p;

@@ -332,7 +330,9 @@ void main() {
    tensorLayoutNV<2> tensorLayoutA = createTensorLayoutNV(2);
    tensorLayoutNV<2, gl_CooperativeMatrixClampModeConstantNV> tensorLayoutAClamp = createTensorLayoutNV(2, gl_CooperativeMatrixClampModeConstantNV);
    tensorLayoutNV<2> tensorLayoutB = createTensorLayoutNV(2);
+#ifndef MUL_MAT_ID
    tensorLayoutNV<2, gl_CooperativeMatrixClampModeConstantNV> tensorLayoutBClamp = createTensorLayoutNV(2, gl_CooperativeMatrixClampModeConstantNV);
+#endif
    tensorLayoutNV<2, gl_CooperativeMatrixClampModeConstantNV> tensorLayoutD = createTensorLayoutNV(2, gl_CooperativeMatrixClampModeConstantNV);

 #if QUANT_K > 1
@@ -345,12 +345,19 @@ void main() {

    // Use end_k rather than p.K as the dimension because that's what
    // we need to bound check against when using split_k.
-    // Bounds check B against padded_N, but bounds check D against N.
    tensorLayoutA = setTensorLayoutDimensionNV(tensorLayoutA, p.M, end_k);
+#ifdef MUL_MAT_ID
+    // MUL_MAT_ID pads each B row to stride_b so partial K tiles read zeros without clamping.
+    tensorLayoutB = setTensorLayoutDimensionNV(tensorLayoutB, BN, p.stride_b);
+#else
+    // Bounds check B against padded_N, but bounds check D against N.
    tensorLayoutB = setTensorLayoutDimensionNV(tensorLayoutB, p.padded_N, end_k);
+#endif
    tensorLayoutD = setTensorLayoutDimensionNV(tensorLayoutD, p.N, p.M);
    tensorLayoutAClamp = setTensorLayoutDimensionNV(tensorLayoutAClamp, p.M, end_k);
+#ifndef MUL_MAT_ID
    tensorLayoutBClamp = setTensorLayoutDimensionNV(tensorLayoutBClamp, p.padded_N, end_k);
+#endif

    tensorLayoutD = setTensorLayoutStrideNV(tensorLayoutD, p.stride_d, 1);

@@ -527,7 +534,9 @@ void main() {

        tensorLayoutB = setTensorLayoutStrideNV(tensorLayoutB, stride_b, 1);

+#ifndef MUL_MAT_ID
        tensorLayoutBClamp = setTensorLayoutStrideNV(tensorLayoutBClamp, stride_b, 1);
+#endif

        uint k_iters = (end_k - start_k + BK - 1) / BK;

@@ -56,7 +56,6 @@ layout (push_constant) uniform parameter
    uint nei1;
    uint nbi1;
    uint ne11;
-    uint padded_N;
    uint n_experts;
    uint hoist_row_ids;
 #else
@@ -1,4 +1,6 @@

+#include "utils.glsl"
+
 // vk_op_sum_rows_push_constants
 layout (push_constant) uniform parameter
 {
@@ -15,11 +17,3 @@ layout (push_constant) uniform parameter
 uint get_aoffset() { return p.misalign_offsets >> 16; }
 uint get_doffset() { return p.misalign_offsets & 0xFFFF; }

-// see init_fastdiv_values in ggml-vulkan.cpp
-uint fastdiv(uint n, uint mp, uint L) {
-    uint msbs, lsbs;
-    // msbs = mulhi(n, mp)
-    umulExtended(n, mp, msbs, lsbs);
-    return (msbs + n) >> L;
-}
-
@@ -9,14 +9,26 @@ uint fastmod(uint a, uint b) {
    return a % b;
 }

-uint fastdiv(uint a, uint b) {
+// see init_fastdiv_values in ggml-vulkan.cpp
+uint fastdiv(uint n, uint mp, uint L) {
+    uint msbs, lsbs;
+    // msbs = mulhi(n, mp)
+    umulExtended(n, mp, msbs, lsbs);
+    return (msbs + n) >> L;
+}
+
+uint fastdiv_L(uint packed, uint slot) {
+    return (packed >> (slot * 8)) & 0x3Fu;
+}
+
+uint fastdiv_small(uint a, uint b) {
    return (a < b) ? 0 : (a / b);
 }

 void get_indices(uint idx, out uint i00, out uint i01, out uint i02, out uint i03, uint ne00, uint ne01, uint ne02, uint ne03) {
-    i03 = fastdiv(idx, (ne02*ne01*ne00));
+    i03 = fastdiv_small(idx, (ne02*ne01*ne00));
    const uint i03_offset = i03 * ne02*ne01*ne00;
-    i02 = fastdiv((idx - i03_offset), (ne01*ne00));
+    i02 = fastdiv_small((idx - i03_offset), (ne01*ne00));
    const uint i02_offset = i02*ne01*ne00;
    i01 = (idx - i03_offset - i02_offset) / ne00;
    i00 = idx - i03_offset - i02_offset - i01*ne00;
@@ -214,10 +214,10 @@ extern "C" {
    LLAMA_API const char * llama_load_mode_name(enum llama_load_mode load_mode);
    LLAMA_API enum llama_load_mode llama_load_mode_from_str(const char * str);

-    enum llama_tensor_read_lazy {
-        LLAMA_TENSOR_READ_LAZY_OFF  = 0, // always read the whole tensor up front
-        LLAMA_TENSOR_READ_LAZY_AUTO = 1, // lazy only for marked tensors larger than 4 GiB (requires mmap)
-        LLAMA_TENSOR_READ_LAZY_ON   = 2, // read the rows of tensors marked by the arch on demand (requires mmap)
+    enum llama_lazy_mode {
+        LLAMA_LAZY_MODE_OFF  = 0, // always read the whole tensor up front
+        LLAMA_LAZY_MODE_AUTO = 1, // lazy only for marked tensors larger than 4 GiB (requires mmap)
+        LLAMA_LAZY_MODE_ON   = 2, // read the rows of tensors marked by the arch on demand (requires mmap)
    };

    enum llama_context_type {
@@ -321,7 +321,7 @@ extern "C" {
        enum llama_split_mode split_mode; // how to split the model across multiple GPUs
        enum llama_load_mode  load_mode;  // how to load the model

-        enum llama_tensor_read_lazy tensor_read_lazy; // on-demand reading of tensors marked by the arch
+        enum llama_lazy_mode lazy_mode; // on-demand reading of tensors marked by the arch

        // the GPU that is used for the entire model when split_mode is LLAMA_SPLIT_MODE_NONE
        int32_t main_gpu;
@@ -11,6 +11,8 @@ import platform
 import shutil
 import logging

+from sdk import validate_windows_sdks
+
 logger = logging.getLogger("build")


@@ -65,6 +67,13 @@ def main():
        logger.error(f"Error: Invalid target format '{args.target}'. Must be android[:serial]/adb[:serial], linux:[user@]host/lnx:[user@]host/ubuntu:[user@]host, or windows/wos.")
        sys.exit(1)

+    if target_type == "windows":
+        logger.info("Windows target selected. Forcing native compilation...")
+        args.no_docker = True
+        if platform.system() != "Windows":
+            logger.warning("Warning: Windows compilation is intended to run on Windows arm64 hosts.")
+        validate_windows_sdks()
+
    # Determine preset and check if it's debug
    preset = args.preset
    if preset:
@@ -120,12 +129,6 @@ def main():

    jobs = args.jobs if args.jobs else os.cpu_count() or 4

-    if target_type == "windows":
-        logger.info("Windows target selected. Forcing native compilation...")
-        args.no_docker = True
-        if platform.system() != "Windows":
-            logger.warning("Warning: Windows compilation is intended to run on Windows arm64 hosts.")
-
    if args.no_docker:
        # Native/local host build
        logger.info("Running native/local CMake build...")
@@ -258,3 +261,6 @@ if __name__ == "__main__":
    except KeyboardInterrupt:
        logger.info("\nInterrupted by user.")
        sys.exit(130)
+    except RuntimeError as err:
+        logger.error("Error: %s", err)
+        sys.exit(1)
@@ -0,0 +1,62 @@
+import os
+from pathlib import Path
+
+
+SDK_CONFIGS = (
+    {
+        "name": "Hexagon SDK",
+        "repo": "snapdragon-toolchain/hexagon-sdk",
+        "default_version": "6.6.0.0",
+        "parent_dir": "Hexagon_SDK",
+        "archive_prefix": "hexagon-sdk-v",
+        "markers": ("hexagon_sdk.json",),
+    },
+    {
+        "name": "OpenCL SDK",
+        "repo": "snapdragon-toolchain/opencl-sdk",
+        "default_version": "2.3.2",
+        "parent_dir": "OpenCL_SDK",
+        "archive_prefix": "adreno-opencl-sdk-v",
+        "markers": ("include/CL", "lib/OpenCL.lib"),
+    },
+)
+
+
+def is_valid_sdk(config, target_dir):
+    return target_dir.is_dir() and all((target_dir / marker).exists() for marker in config["markers"])
+
+
+def get_hexagon_tools_dir(hexagon_dir):
+    tools_parent = hexagon_dir / "tools" / "HEXAGON_Tools"
+    if not tools_parent.is_dir():
+        raise RuntimeError(f"Expected Hexagon tools directory in {tools_parent}")
+    tools_dirs = [path for path in tools_parent.iterdir() if path.is_dir()]
+    if len(tools_dirs) != 1:
+        raise RuntimeError(f"Expected one Hexagon tools directory in {tools_parent}")
+    return tools_dirs[0]
+
+
+def validate_windows_sdks():
+    hexagon_config, opencl_config = SDK_CONFIGS
+    hexagon_dir = os.environ.get("HEXAGON_SDK_ROOT")
+    tools_dir = os.environ.get("HEXAGON_TOOLS_ROOT")
+    opencl_dir = os.environ.get("OPENCL_SDK_ROOT")
+    missing = []
+
+    expected_tools_dir = None
+    if not hexagon_dir or not is_valid_sdk(hexagon_config, Path(hexagon_dir)):
+        missing.append("HEXAGON_SDK_ROOT")
+    else:
+        try:
+            expected_tools_dir = get_hexagon_tools_dir(Path(hexagon_dir))
+        except RuntimeError:
+            pass
+    if not tools_dir or not expected_tools_dir or Path(tools_dir) != expected_tools_dir:
+        missing.append("HEXAGON_TOOLS_ROOT")
+    if not opencl_dir or not is_valid_sdk(opencl_config, Path(opencl_dir)):
+        missing.append("OPENCL_SDK_ROOT")
+    if missing:
+        raise RuntimeError(
+            f"Missing or invalid Windows SDK paths: {', '.join(missing)}. "
+            "Run scripts/snapdragon/setup-sdk.py first."
+        )
@@ -0,0 +1,233 @@
+#!/usr/bin/env python3
+#
+# Install Windows on Snapdragon SDKs for llama.cpp.
+#
+
+import sys
+import os
+import argparse
+import shutil
+import logging
+import json
+import hashlib
+import tarfile
+import tempfile
+from pathlib import Path
+from urllib.error import HTTPError, URLError
+from urllib.request import Request, urlopen
+
+from sdk import SDK_CONFIGS, get_hexagon_tools_dir, is_valid_sdk
+
+
+logger = logging.getLogger("setup_sdk")
+
+DEFAULT_SDK_BASE_DIR = r"C:\Qualcomm"
+
+
+def get_sdk_releases(config):
+    request = Request(
+        f"https://api.github.com/repos/{config['repo']}/releases?per_page=100",
+        headers={"Accept": "application/vnd.github+json", "User-Agent": "llama.cpp"},
+    )
+    try:
+        with urlopen(request, timeout=30) as response:
+            releases = json.load(response)
+    except (HTTPError, URLError, TimeoutError) as err:
+        raise RuntimeError(f"Cannot query {config['name']} releases: {err}") from err
+
+    result = []
+    for release in releases:
+        if release["draft"] or release["prerelease"]:
+            continue
+        version = release["tag_name"].removeprefix("v")
+        archive_name = f"{config['archive_prefix']}{version}-arm64-wos.tar.xz"
+        for asset in release["assets"]:
+            if asset["name"] != archive_name:
+                continue
+            result.append({
+                "version": version,
+                "name": asset["name"],
+                "url": asset["browser_download_url"],
+                "sha256": (asset.get("digest") or "").removeprefix("sha256:"),
+            })
+    return result
+
+
+def list_sdk_releases():
+    for config in SDK_CONFIGS:
+        logger.info("%s:", config["name"])
+        releases = get_sdk_releases(config)
+        if not releases:
+            logger.info("  no Windows on Snapdragon releases found")
+            continue
+        for release in releases:
+            logger.info("  %s: %s", release["version"], release["name"])
+
+
+def get_sdk_release(config, version):
+    version = version or config["default_version"]
+    version = version.removeprefix("v")
+    for release in get_sdk_releases(config):
+        if release["version"] == version:
+            if not release["sha256"]:
+                raise RuntimeError(f"{config['name']} {version} does not provide a SHA-256 digest")
+            return release
+    raise RuntimeError(
+        f"No Windows on Snapdragon release for {config['name']} {version}. "
+        "Run scripts/snapdragon/setup-sdk.py --list-sdk-releases to see available versions."
+    )
+
+
+def sha256sum(path):
+    digest = hashlib.sha256()
+    with open(path, "rb") as file:
+        for chunk in iter(lambda: file.read(1024 * 1024), b""):
+            digest.update(chunk)
+    return digest.hexdigest()
+
+
+def download_sdk(release, archive):
+    while True:
+        if archive.exists() and sha256sum(archive) == release["sha256"]:
+            logger.info("Using existing archive %s", archive)
+            return
+
+        offset = archive.stat().st_size if archive.exists() else 0
+        headers = {"User-Agent": "llama.cpp"}
+        if offset:
+            headers["Range"] = f"bytes={offset}-"
+            logger.info("Resuming download of %s at %d MiB", release["name"], offset // (1024 * 1024))
+        else:
+            logger.info("Downloading %s", release["name"])
+
+        try:
+            with urlopen(Request(release["url"], headers=headers), timeout=30) as response:
+                mode = "ab" if offset and response.status == 206 else "wb"
+                with open(archive, mode) as file:
+                    shutil.copyfileobj(response, file)
+        except HTTPError as err:
+            if err.code != 416:
+                raise RuntimeError(f"Cannot download {release['name']}: {err}") from err
+            archive.unlink(missing_ok=True)
+            continue
+        except (URLError, TimeoutError) as err:
+            raise RuntimeError(f"Cannot download {release['name']}: {err}") from err
+
+        if sha256sum(archive) == release["sha256"]:
+            return
+        raise RuntimeError(f"SHA-256 mismatch for {archive}. Re-run the command to resume the download.")
+
+
+def extract_sdk(config, archive, target_dir):
+    if not hasattr(tarfile, "data_filter"):
+        raise RuntimeError("SDK extraction requires Python 3.10.12 or later")
+
+    with tempfile.TemporaryDirectory(prefix=f".{target_dir.name}.tmp-", dir=target_dir.parent) as staging_path:
+        staging_dir = Path(staging_path)
+        with tarfile.open(archive, "r:xz") as tar:
+            tar.extractall(staging_dir, filter=tarfile.data_filter)
+
+        candidates = [staging_dir] + [path for path in staging_dir.iterdir() if path.is_dir()]
+        extracted_dirs = [path for path in candidates if is_valid_sdk(config, path)]
+        if len(extracted_dirs) != 1:
+            raise RuntimeError(f"{config['name']} archive does not contain the expected files")
+        extracted_dir = extracted_dirs[0]
+
+        backup_dir = None
+        if target_dir.exists():
+            backup_dir = target_dir.parent / f".{target_dir.name}.backup"
+            if backup_dir.exists():
+                raise RuntimeError(f"Cannot replace {target_dir}: backup directory {backup_dir} already exists")
+            target_dir.replace(backup_dir)
+        try:
+            extracted_dir.replace(target_dir)
+        except Exception:
+            if backup_dir:
+                backup_dir.replace(target_dir)
+            raise
+        if backup_dir:
+            shutil.rmtree(backup_dir)
+
+
+def install_sdk(config, version, base_dir, force):
+    version = (version or config["default_version"]).removeprefix("v")
+    target_dir = base_dir / config["parent_dir"] / version
+    if is_valid_sdk(config, target_dir) and not force:
+        logger.info("Using existing %s at %s", config["name"], target_dir)
+        return target_dir
+
+    release = get_sdk_release(config, version)
+    target_dir.parent.mkdir(parents=True, exist_ok=True)
+    archive = target_dir.parent / release["name"]
+    download_sdk(release, archive)
+    logger.info("Extracting %s to %s", config["name"], target_dir)
+    extract_sdk(config, archive, target_dir)
+    archive.unlink(missing_ok=True)
+    return target_dir
+
+
+def set_user_environment(values):
+    if os.name != "nt":
+        raise RuntimeError("SDK setup must run on Windows")
+
+    import winreg
+
+    with winreg.CreateKey(winreg.HKEY_CURRENT_USER, "Environment") as key:
+        for name, value in values.items():
+            winreg.SetValueEx(key, name, 0, winreg.REG_SZ, str(value))
+            os.environ[name] = str(value)
+
+    import ctypes
+
+    result = ctypes.c_ulong()
+    ctypes.windll.user32.SendMessageTimeoutW(0xffff, 0x001a, 0, "Environment", 0x0002, 5000, ctypes.byref(result))
+
+
+def setup_sdks(args):
+    base_dir = Path(args.sdk_base_dir).expanduser().resolve()
+    hexagon_config, opencl_config = SDK_CONFIGS
+    environment = {}
+
+    if args.hexagon is not None:
+        hexagon_dir = install_sdk(hexagon_config, args.hexagon, base_dir, args.force)
+        environment["HEXAGON_SDK_ROOT"] = hexagon_dir
+        environment["HEXAGON_TOOLS_ROOT"] = get_hexagon_tools_dir(hexagon_dir)
+    if args.opencl is not None:
+        opencl_dir = install_sdk(opencl_config, args.opencl, base_dir, args.force)
+        environment["OPENCL_SDK_ROOT"] = opencl_dir
+
+    set_user_environment(environment)
+    logger.info("SDK environment variables were updated. Start a new terminal before building.")
+
+
+def main():
+    logging.basicConfig(level=logging.INFO, format="%(message)s")
+    parser = argparse.ArgumentParser(description="Install Windows on Snapdragon SDKs for llama.cpp.")
+    parser.add_argument("--list-sdk-releases", action="store_true", help="List available Windows on Snapdragon SDK releases")
+    parser.add_argument("--sdk-base-dir", default=DEFAULT_SDK_BASE_DIR, help=r"SDK installation directory (default: C:\Qualcomm)")
+    parser.add_argument("--hexagon", nargs="?", const=SDK_CONFIGS[0]["default_version"], metavar="VERSION", help="Install the Hexagon SDK, optionally selecting a version")
+    parser.add_argument("--opencl", nargs="?", const=SDK_CONFIGS[1]["default_version"], metavar="VERSION", help="Install the OpenCL SDK, optionally selecting a version")
+    parser.add_argument("--force", action="store_true", help="Reinstall selected SDKs even when they already exist")
+    args = parser.parse_args()
+
+    if args.list_sdk_releases:
+        if args.sdk_base_dir != DEFAULT_SDK_BASE_DIR or args.hexagon is not None or args.opencl is not None or args.force:
+            parser.error("Installation options cannot be combined with --list-sdk-releases")
+        list_sdk_releases()
+        return
+    if args.hexagon is None and args.opencl is None:
+        parser.error("Select at least one SDK with --hexagon or --opencl")
+    if os.name != "nt":
+        parser.error("SDK setup must run on Windows")
+    setup_sdks(args)
+
+
+if __name__ == "__main__":
+    try:
+        main()
+    except KeyboardInterrupt:
+        logger.info("\nInterrupted by user.")
+        sys.exit(130)
+    except RuntimeError as err:
+        logger.error("Error: %s", err)
+        sys.exit(1)
@@ -1287,10 +1287,10 @@ struct ggml_tensor * llama_model_loader::create_tensor(
        return NULL;
    }

-    if ((flags & TENSOR_READ_LAZY) && use_mmap && tensor_read_lazy != LLAMA_TENSOR_READ_LAZY_OFF) {
+    if ((flags & TENSOR_READ_LAZY) && use_mmap && lazy_mode != LLAMA_LAZY_MODE_OFF) {
        // in auto mode, small tensors are cheap enough to keep resident
        constexpr size_t auto_lazy_min_size = 4ull * 1024 * 1024 * 1024;
-        if (tensor_read_lazy == LLAMA_TENSOR_READ_LAZY_ON || ggml_nbytes(cur) > auto_lazy_min_size) {
+        if (lazy_mode == LLAMA_LAZY_MODE_ON || ggml_nbytes(cur) > auto_lazy_min_size) {
            const auto & w = require_weight(tn.str().c_str());
            lazy_tensor_ranges[w.idx].emplace_back(w.offs, w.offs + ggml_nbytes(cur));

@@ -84,7 +84,7 @@ struct llama_model_loader {
    bool load_mtp;

    // set by the caller before the create_tensor() calls
-    enum llama_tensor_read_lazy tensor_read_lazy = LLAMA_TENSOR_READ_LAZY_OFF;
+    enum llama_lazy_mode lazy_mode = LLAMA_LAZY_MODE_OFF;

    llama_files files;
    llama_ftype ftype;
@@ -2681,7 +2681,7 @@ llama_model_params llama_model_default_params() {
        /*.n_gpu_layers                =*/ -1,
        /*.split_mode                  =*/ LLAMA_SPLIT_MODE_LAYER,
        /*.load_mode                   =*/ LLAMA_LOAD_MODE_AUTO,
-        /*.tensor_read_lazy            =*/ LLAMA_TENSOR_READ_LAZY_AUTO,
+        /*.lazy_mode                   =*/ LLAMA_LAZY_MODE_AUTO,
        /*.main_gpu                    =*/ 0,
        /*.tensor_split                =*/ nullptr,
        /*.progress_callback           =*/ nullptr,
@@ -318,7 +318,7 @@ static std::pair<int, llama_model *> llama_model_load(struct gguf_context * meta
        llama_model_loader ml(metadata, set_tensor_data, set_tensor_data_ud, fname, splits, file, params.load_mode,
            params.check_tensors, params.no_alloc, params.load_mtp, params.kv_overrides, params.tensor_buft_overrides);

-        ml.tensor_read_lazy = params.tensor_read_lazy;
+        ml.lazy_mode = params.lazy_mode;

        ml.print_info();
        std::unique_ptr<llama_model> model_ptr(llama_model_create(ml, params));
@@ -2360,9 +2360,12 @@ struct llama_model_qwen4exp : public llama_model_base {
                        int64_t   channels,
                            int   il);

+        ggml_tensor * build_inp_ple(
+  const llama_memory_hybrid_idx_context * mctx_hyb);
+
        ggml_tensor * build_ple(
             llm_graph_input_rs * inp,
-  const llama_memory_hybrid_idx_context * mctx_hyb,
+                    ggml_tensor * emb,
                    ggml_tensor * hidden,
                            int   il);

@@ -296,6 +296,7 @@ llama_model_qwen4exp::graph::graph(const llama_model & model, const llm_graph_pa

    ggml_tensor * inpL = build_inp_embd(model.tok_embd);
    cb(inpL, "model.input_embed", -1);
+    ggml_build_forward_expand(gf, inpL);

    auto * inp = build_inp_mem_hybrid();

@@ -312,6 +313,13 @@ llama_model_qwen4exp::graph::graph(const llama_model & model, const llm_graph_pa
    ggml_tensor * inp_pos     = build_inp_pos();
    ggml_tensor * inp_out_ids = build_inp_out_ids();

+    ggml_tensor * ple_emb = nullptr;
+    if (hparams.ple_n_heads > 0) {
+        ple_emb = build_inp_ple(mctx_hyb);
+        // make sure ple_emb and build_inp_embd are in the same graph split
+        ggml_build_forward_expand(gf, ple_emb);
+    }
+
    // the wide residual starts as hc identical copies of the embedding
    ggml_tensor * res_hc = ggml_repeat_4d(ctx0,
            ggml_reshape_3d(ctx0, inpL, n_embd, 1, n_tokens),
@@ -322,7 +330,7 @@ llama_model_qwen4exp::graph::graph(const llama_model & model, const llm_graph_pa
        res->t_layer_inp[il] = res_hc;

        if (hparams.is_ple(il)) {
-            res_hc = build_ple(inp->get_recr(), mctx_hyb, res_hc, il);
+            res_hc = build_ple(inp->get_recr(), ple_emb, res_hc, il);
        }

        ggml_tensor * inject = nullptr;
@@ -1090,13 +1098,8 @@ ggml_tensor * llama_model_qwen4exp::graph::build_conv_state_at(
    return conv_input;
 }

-ggml_tensor * llama_model_qwen4exp::graph::build_ple(
-        llm_graph_input_rs * inp,
-        const llama_memory_hybrid_idx_context * mctx_hyb,
-        ggml_tensor *        hidden,
-        int                  il) {
-    const int64_t hc      = hparams.dsv4_hc_mult;
-    const int64_t hc_dim  = hc * n_embd;
+ggml_tensor * llama_model_qwen4exp::graph::build_inp_ple(
+        const llama_memory_hybrid_idx_context * mctx_hyb) {
    const int64_t n_heads = hparams.ple_n_heads;

    // the attention cells see every ubatch regardless of the layer types
@@ -1111,7 +1114,18 @@ ggml_tensor * llama_model_qwen4exp::graph::build_ple(
    // gather then flatten the heads: get_rows lays the head dimension out slowest, as the reference does
    ggml_tensor * emb = ggml_get_rows(ctx0, model.per_layer_tok_embd, rows);
    emb = ggml_reshape_2d(ctx0, emb, hparams.ple_head_dim * n_heads, n_tokens);
-    cb(emb, "ple_embd", il);
+    cb(emb, "ple_embd", -1);
+
+    return emb;
+}
+
+ggml_tensor * llama_model_qwen4exp::graph::build_ple(
+        llm_graph_input_rs * inp,
+        ggml_tensor *        emb,
+        ggml_tensor *        hidden,
+        int                  il) {
+    const int64_t hc      = hparams.dsv4_hc_mult;
+    const int64_t hc_dim  = hc * n_embd;

    ggml_tensor * key   = build_lora_mm(model.layers[il].ple_key,   emb);
    ggml_tensor * value = build_lora_mm(model.layers[il].ple_value, emb);
@@ -8785,6 +8785,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
        test_cases.emplace_back(new test_conv_transpose_2d({3, 2, 3, 1}, {2, 2, 1, 3}, 1, kernel_type));
        test_cases.emplace_back(new test_conv_transpose_2d({10, 10, 9, 1}, {3, 3, 1, 9}, 2, kernel_type));
        test_cases.emplace_back(new test_conv_transpose_2d({129, 63, 35, 1}, {3, 3, 48, 35}, 1, kernel_type));
+        test_cases.emplace_back(new test_conv_transpose_2d({10, 10, 9, 2}, {3, 3, 1, 9}, 2, kernel_type)); // for multiple batches
    }

    test_cases.emplace_back(new test_count_equal(GGML_TYPE_F32, {4,  500, 1, 1}));
@@ -9366,7 +9367,12 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_BF16, GGML_TYPE_F32, 16, 16, b, 50, 200, 64));
    }

-    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 1, 1, false, 8, 16, 1));
+    // For issue 27873
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_IQ2_XXS, GGML_TYPE_F32, 1, 1, false, 1, 8192, 4096));
+
+    for (int k : {1, 63, 65}) {
+        test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 1, 1, false, 8, 16, k));
+    }
    test_cases.emplace_back(new test_mul_mat_id_fusion(GGML_TYPE_F16, GGML_TYPE_F32, 16, 16, false, 32, 32, 32, 3));

    // gpt-oss issue with Vulkan mmq_id
@@ -67,6 +67,7 @@ test parameters:
  -nkvo, --no-kv-offload <0|1>              (default: 0)
  -fa, --flash-attn <on|off|auto>           (default: auto)
  -dev, --device <dev0/dev1/...>            (default: auto)
+  --tensor-read-lazy <on|auto|off>          (default: auto)
  -mmp, --mmap <0|1>                        (DEPRECATED IN FAVOUR OF --load-mode)
  -dio, --direct-io <0|1>                   (DEPRECATED IN FAVOUR OF --load-mode)
  -embd, --embeddings <0|1>                 (default: 0)
@@ -271,6 +271,19 @@ static const char * split_mode_str(llama_split_mode mode) {
    }
 }

+static const char * lazy_mode_str(llama_lazy_mode mode) {
+    switch (mode) {
+        case LLAMA_LAZY_MODE_OFF:
+            return "off";
+        case LLAMA_LAZY_MODE_AUTO:
+            return "auto";
+        case LLAMA_LAZY_MODE_ON:
+            return "on";
+        default:
+            GGML_ABORT("invalid tensor read lazy mode");
+    }
+}
+
 static std::string pair_str(const std::pair<int, int> & p) {
    static char buf[32];
    snprintf(buf, sizeof(buf), "%d,%d", p.first, p.second);
@@ -341,6 +354,7 @@ struct cmd_params {
    std::vector<int>                 n_cpu_moe;
    std::vector<llama_split_mode>    split_mode;
    std::vector<llama_load_mode>     load_mode;
+    std::vector<llama_lazy_mode>     lazy_mode;
    std::vector<int>                 main_gpu;
    std::vector<bool>                no_kv_offload;
    std::vector<llama_flash_attn_type> flash_attn;
@@ -385,6 +399,7 @@ static const cmd_params cmd_params_defaults = {
    /* n_cpu_moe            */ { 0 },
    /* split_mode           */ { LLAMA_SPLIT_MODE_LAYER },
    /* load_mode            */ { LLAMA_LOAD_MODE_AUTO },
+    /* lazy_mode            */ { LLAMA_LAZY_MODE_AUTO },
    /* main_gpu             */ { 0 },
    /* no_kv_offload        */ { false },
    /* flash_attn           */ { LLAMA_FLASH_ATTN_TYPE_AUTO },
@@ -460,6 +475,7 @@ static void print_usage(int /* argc */, char ** argv) {
    printf("  -fa, --flash-attn <on|off|auto>                   (default: %s)\n", join(transform_to_str(cmd_params_defaults.flash_attn, llama_flash_attn_type_name), ",").c_str());
    printf("  -dev, --device <dev0/dev1/...>                    (default: auto)\n");
    printf("  -lm, --load-mode <auto|none|mmap|mlock|mmap+mlock|dio> (default: %s)\n", join(transform_to_str(cmd_params_defaults.load_mode, llama_load_mode_name), ",").c_str());
+    printf("  --tensor-read-lazy <on|auto|off>                  (default: %s)\n", join(transform_to_str(cmd_params_defaults.lazy_mode, lazy_mode_str), ",").c_str());
    printf("  -mmp, --mmap <0|1>                                (DEPRECATED IN FAVOUR OF --load-mode)\n");
    printf("  -dio, --direct-io <0|1>                           (DEPRECATED IN FAVOUR OF --load-mode)\n");
    printf("  -embd, --embeddings <0|1>                         (default: %s)\n", join(cmd_params_defaults.embeddings, ",").c_str());
@@ -786,6 +802,32 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
                    break;
                }
                params.load_mode.insert(params.load_mode.end(), modes.begin(), modes.end());
+            } else if (arg == "--tensor-read-lazy") {
+                if (++i >= argc) {
+                    invalid_param = true;
+                    break;
+                }
+                auto p = string_split<std::string>(argv[i], split_delim);
+
+                std::vector<llama_lazy_mode> modes;
+                for (const auto & m : p) {
+                    llama_lazy_mode mode;
+                    if (m == "on") {
+                        mode = LLAMA_LAZY_MODE_ON;
+                    } else if (m == "auto") {
+                        mode = LLAMA_LAZY_MODE_AUTO;
+                    } else if (m == "off") {
+                        mode = LLAMA_LAZY_MODE_OFF;
+                    } else {
+                        invalid_param = true;
+                        break;
+                    }
+                    modes.push_back(mode);
+                }
+                if (invalid_param) {
+                    break;
+                }
+                params.lazy_mode.insert(params.lazy_mode.end(), modes.begin(), modes.end());
            } else if (arg == "-mg" || arg == "--main-gpu") {
                if (++i >= argc) {
                    invalid_param = true;
@@ -1137,6 +1179,9 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
    if (params.load_mode.empty()) {
        params.load_mode = cmd_params_defaults.load_mode;
    }
+    if (params.lazy_mode.empty()) {
+        params.lazy_mode = cmd_params_defaults.lazy_mode;
+    }
    if (params.main_gpu.empty()) {
        params.main_gpu = cmd_params_defaults.main_gpu;
    }
@@ -1203,6 +1248,7 @@ struct cmd_params_instance {
    int                n_cpu_moe;
    llama_split_mode   split_mode;
    llama_load_mode    load_mode;
+    llama_lazy_mode    lazy_mode;
    int                main_gpu;
    bool               no_kv_offload;
    llama_flash_attn_type flash_attn;
@@ -1224,6 +1270,7 @@ struct cmd_params_instance {
        }
        mparams.split_mode    = split_mode;
        mparams.load_mode     = load_mode;
+        mparams.lazy_mode     = lazy_mode;
        mparams.main_gpu      = main_gpu;
        mparams.tensor_split  = tensor_split.data();
        mparams.no_host       = no_host;
@@ -1271,7 +1318,8 @@ struct cmd_params_instance {
        return model == other.model && n_gpu_layers == other.n_gpu_layers && n_cpu_moe == other.n_cpu_moe &&
               split_mode == other.split_mode &&
               main_gpu == other.main_gpu && tensor_split == other.tensor_split &&
-               load_mode == other.load_mode && devices == other.devices && no_host == other.no_host &&
+               load_mode == other.load_mode && lazy_mode == other.lazy_mode &&
+               devices == other.devices && no_host == other.no_host &&
               vec_tensor_buft_override_equal(tensor_buft_overrides, other.tensor_buft_overrides);
    }

@@ -1305,6 +1353,7 @@ static std::vector<cmd_params_instance> get_cmd_params_instances(const cmd_param
    for (const auto & ncmoe : params.n_cpu_moe)
    for (const auto & sm : params.split_mode)
    for (const auto & lm : params.load_mode)
+    for (const auto & lzm : params.lazy_mode)
    for (const auto & mg : params.main_gpu)
    for (const auto & devs : params.devices)
    for (const auto & ts : params.tensor_split)
@@ -1344,6 +1393,7 @@ static std::vector<cmd_params_instance> get_cmd_params_instances(const cmd_param
                /* .n_cpu_moe             = */ ncmoe,
                /* .split_mode            = */ sm,
                /* .load_mode             = */ lm,
+                /* .lazy_mode             = */ lzm,
                /* .main_gpu              = */ mg,
                /* .no_kv_offload         = */ nkvo,
                /* .flash_attn            = */ fa,
@@ -1380,6 +1430,7 @@ static std::vector<cmd_params_instance> get_cmd_params_instances(const cmd_param
                /* .n_cpu_moe             = */ ncmoe,
                /* .split_mode            = */ sm,
                /* .load_mode             = */ lm,
+                /* .lazy_mode             = */ lzm,
                /* .main_gpu              = */ mg,
                /* .no_kv_offload         = */ nkvo,
                /* .flash_attn            = */ fa,
@@ -1416,6 +1467,7 @@ static std::vector<cmd_params_instance> get_cmd_params_instances(const cmd_param
                /* .n_cpu_moe             = */ ncmoe,
                /* .split_mode            = */ sm,
                /* .load_mode             = */ lm,
+                /* .lazy_mode             = */ lzm,
                /* .main_gpu              = */ mg,
                /* .no_kv_offload         = */ nkvo,
                /* .flash_attn            = */ fa,
@@ -1457,6 +1509,7 @@ struct test {
    int                      n_cpu_moe;
    llama_split_mode         split_mode;
    llama_load_mode          load_mode;
+    llama_lazy_mode          lazy_mode;
    int                      main_gpu;
    bool                     no_kv_offload;
    llama_flash_attn_type    flash_attn;
@@ -1496,6 +1549,7 @@ struct test {
        n_cpu_moe      = inst.n_cpu_moe;
        split_mode     = inst.split_mode;
        load_mode      = inst.load_mode;
+        lazy_mode      = inst.lazy_mode;
        main_gpu       = inst.main_gpu;
        no_kv_offload  = inst.no_kv_offload;
        flash_attn     = inst.flash_attn;
@@ -1563,7 +1617,8 @@ struct test {
            "n_ubatch",       "n_threads",      "cpu_mask",      "cpu_strict",     "poll",
            "type_k",         "type_v",         "n_gpu_layers",  "n_cpu_moe",      "split_mode",
            "main_gpu",       "no_kv_offload",  "flash_attn",    "devices",        "tensor_split",
-            "tensor_buft_overrides",            "load_mode",     "embeddings",
+            "tensor_buft_overrides",            "load_mode",     "lazy_mode",
+            "embeddings",
            "no_op_offload",  "no_host",        "fit_target",    "fit_min_ctx",
            "n_prompt",       "n_gen",          "n_depth",
            "test_time",      "avg_ns",         "stddev_ns",     "avg_ts",         "stddev_ts"
@@ -1588,7 +1643,7 @@ struct test {
        if (field == "avg_ts" || field == "stddev_ts") {
            return FLOAT;
        }
-        if (field == "load_mode") {
+        if (field == "load_mode" || field == "lazy_mode") {
            return STRING;
        }
        return STRING;
@@ -1658,6 +1713,7 @@ struct test {
                                            tensor_split_str,
                                            tensor_buft_overrides_str,
                                            llama_load_mode_name(load_mode),
+                                            lazy_mode_str(lazy_mode),
                                            std::to_string(embeddings),
                                            std::to_string(no_op_offload),
                                            std::to_string(no_host),
@@ -1972,6 +2028,9 @@ struct markdown_printer : public printer {
        if (params.load_mode.size() > 1 || params.load_mode != cmd_params_defaults.load_mode) {
            fields.emplace_back("load_mode");
        }
+        if (params.lazy_mode.size() > 1 || params.lazy_mode != cmd_params_defaults.lazy_mode) {
+            fields.emplace_back("lazy_mode");
+        }
        if (params.embeddings.size() > 1 || params.embeddings != cmd_params_defaults.embeddings) {
            fields.emplace_back("embeddings");
        }
Author	SHA1	Message	Date
Nick FarrellandGitHub	cc83d7b482	sycl: make --fit respect --fit-target better (#27629 ) improve the --fit algorithm to take into account the actual peak required VRAM for a given context size on a SYCL backend. This includes both properly accounting for how much VRAM is required when the allocated context is fully used (which makes the reported context drop below what it did before, but stop it OOMing) as well as preventing some overly-conservative calculations which meant too much VRAM was being reserved. Tested on a Arc b70 with unsloth's qwen3.8 (Q4_K_XL), able to get 262144 context, fully usable, with q8_0 KV and MTP and 4k ubatch size using --fit-target 1	2026-08-29 05:00:09 -04:00
Jeff BolzandGitHub	c9ca51c1f6	vulkan: combine duplicated fastdiv functions, rename the one optimizing small divs (#27526 ) * vulkan: combine duplicated fastdiv functions, rename the one optimizing small divs * remove one more fastdiv	2026-08-29 10:59:48 +03:00
Jhen-Jie HongandGitHub	5ea1b124e7	metal : add fa-vec tunings for M1 Max (#27932 )	2026-08-29 15:12:23 +08:00
Jeff BolzandGitHub	77f132cb1d	vulkan: Change mul_mat_id to pad K rather than N (#27925 ) The N padding is needed for mul_mat, but not mul_mat_id. For mul_mat_id, we indirect the row index through a shared memory lookup table which avoids any OOB row coordinate. But that callback doesn't bounds check K, so we actually need K padding instead.	2026-08-29 10:09:24 +03:00
kurquhar GitHub Kristopher Urquhart	d7bd3bfcad	snapdragon: python SDK setup (Windows) (#27903 ) * port setup-build.ps1 to setup_sdk.py, to facilitate installation of Hexagon and OpenCL SDKs on Windows * rename setup_sdk.py -> setup-sdk.py * flake8 fix: print() -> logger.info() --------- Co-authored-by: Kristopher Urquhart <kurquhar@qti.qualcom.com>	2026-08-28 14:01:59 -07:00
Xuan-Son NguyenandGitHub	50f068ffff	bench: add --tensor-read-lazy (#27881 ) * bench: add --tensor-read-lazy * rm the alias * rename to LLAMA_LAZY_MODE_*	2026-08-28 20:51:05 +02:00
Xuan-Son NguyenandGitHub	6fe7498016	model: qwen4exp: reduce number of graph splits (#27880 )	2026-08-28 19:24:46 +02:00
Eric A Stalee GitHub Jeff Bolz	b387ddfd84	vulkan: fix missing view-alias dependencies in ggml_vk_graph_optimize (#27812 ) * vulkan: fix missing view-alias dependencies in ggml_vk_graph_optimize is_src_of doesn't treat two views of one tensor as dependent, so the optimizer reorders nodes across aliased reads and writes. Result: silently wrong tokens under greedy decoding, different output on every server start, and invalid speculative-decoding acceptance, with nothing logged. Hits Qwen3.8's recurrent state (and any model with view-aliased state) on AMD and NVIDIA Vulkan. CUDA is clean. Compare view_src bases on both sides. Fixes #27805 * vulkan: don't treat view/no-op nodes as aliasing dependencies Nodes whose op is NONE, RESHAPE, TRANSPOSE, VIEW or PERMUTE execute nothing, so aliasing through them is not a real dependency. The previous base comparison matched them anyway, which only costs the optimizer reordering freedom. Co-authored-by: Jeff Bolz <jbolz@nvidia.com> * vulkan: make the lambda parameter const and capture is_empty in is_src_of Code will not compile without these changes. is_src_of has an empty capture list, so is_empty was not visible inside it, and is_empty took a non-const pointer, while is_src_of receives const ones. Other call sites pass non-const pointers, which still convert as usual. --------- Co-authored-by: Jeff Bolz <jbolz@nvidia.com>	2026-08-28 19:12:33 +02:00
Tekin Ertekin GitHub Georgi Gerganov	a43c3986b4	ggml : fix conv_transpose_2d for multiple batches (#26132 ) * ggml : fix conv_transpose_2d for multiple batches ggml_compute_forward_conv_transpose_2d_impl only computed the first batch (ne[3] of the destination); every batch after the first was left as zero. Both the src1 permutation and the main compute loop now iterate over the batch dimension, and the work buffer size in ggml_graph_plan is scaled by the src1 batch count so the extra permuted batches fit. A multi-batch test case is added to test-backend-ops. Fixes ggml-org/ggml#1448 * metal : fix conv_transpose_2d for multiple batches The kernel only computed batch 0 of the input (src1->ne[3]); every output batch after the first was left as zero, so multi-batch conv_transpose_2d results diverged from the CPU reference. The grid now covers all batches (OW x OH x OC x N), the kernel decodes the batch from the grid z coordinate and offsets both the input and destination indices accordingly. nb3 is passed in the kernel args. Assisted-by: pi:llama.cpp/Qwen3.8-27B --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>	2026-08-28 20:09:08 +03:00