bench: add --tensor-read-lazy (#27881 )

* bench: add --tensor-read-lazy * rm the alias * rename to LLAMA_LAZY_MODE_*
model: qwen4exp: reduce number of graph splits (#27880 )
2026-08-29 11:37:42 +02:00 · 2026-08-28 20:51:05 +02:00 · 2026-08-28 19:24:46 +02:00 · 2026-08-28 19:12:33 +02:00 · 2026-08-28 20:09:08 +03:00 · 2026-08-28 16:52:49 +02:00
26 changed files with 380 additions and 109 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;
@@ -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;
 }
@@ -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;
    }
 }
@@ -1348,6 +1348,8 @@ struct vk_mat_mat_id_push_constants {
    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;
 };
 struct vk_mat_vec_id_push_constants {
    uint32_t ncols;
@@ -1428,6 +1430,10 @@ struct vk_op_count_experts_push_constants {
    uint32_t nb00;
    uint32_t nb01;
    uint32_t a_offset;
+    uint32_t n_experts;
+    uint32_t hoist_row_ids;
+    uint32_t ne00mp;
+    uint32_t ne00L;
 };

 struct vk_op_glu_push_constants {
@@ -1606,6 +1612,10 @@ template <> void init_pushconst_fastdiv(vk_op_glu_push_constants &p) {
    init_fastdiv_values(p.ne20,                p.ne2_0mp,      p.ne2_0L);
 }

+template <> void init_pushconst_fastdiv(vk_op_count_experts_push_constants &p) {
+    init_fastdiv_values(p.ne00, p.ne00mp, p.ne00L);
+}
+
 struct vk_op_binary_push_constants {
    uint32_t ne;
    uint32_t ne00; uint32_t ne01; uint32_t ne02; uint32_t ne03; uint32_t nb00; uint32_t nb01; uint32_t nb02; uint32_t nb03;
@@ -5839,7 +5849,11 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {

    ggml_vk_create_pipeline(device, device->pipeline_count_equal_i32, "count_equal_i32", count_equal_i32_len, count_equal_i32_data, "main", 3, sizeof(vk_op_push_constants), {512, 1, 1}, { device->subgroup_size }, 1);

-    ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_len, count_experts_data, "main", 2, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true);
+    if (device->subgroup_arithmetic && device->subgroup_require_full_support) {
+        ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_subgroup_len, count_experts_subgroup_data, "main", 2, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true, true);
+    } else {
+        ggml_vk_create_pipeline(device, device->pipeline_count_experts, "count_experts", count_experts_len, count_experts_data, "main", 2, sizeof(vk_op_count_experts_push_constants), {1, 1, 1}, {}, 1, true);
+    }

    for (auto &s : device->pipeline_solve_tri_f32) {
        const vk_solve_tri_pipeline_state &state = s.first;
@@ -8970,13 +8984,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) {
+        uint32_t padded_n, 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 };
+                                              nei0, nei1, nbi1, ne11, padded_n, 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 });
 }

@@ -10162,6 +10176,12 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
    // const uint64_t ne23 = dst->ne[3];

    const uint64_t n_as = ne02;
+    // n_as counts, n_as offsets, one total, then one packed row id per (expert, token).
+    // Hoisting requires 16-bit indices for the packing and a table that fits one binding.
+    const uint64_t hoisted_row_id_words = 2 * n_as + 1 + nei0 * nei1;
+    const bool hoist_row_ids = n_as <= 256 && nei0 <= 0xffff && nei1 <= 0xffff &&
+                                hoisted_row_id_words * sizeof(uint32_t) <=
+                                    ctx->device->properties.limits.maxStorageBufferRange;

    ggml_backend_vk_buffer_context * dst_buf_ctx = (ggml_backend_vk_buffer_context *)dst->buffer->context;
    ggml_backend_vk_buffer_context * src0_buf_ctx = (ggml_backend_vk_buffer_context *)src0->buffer->context;
@@ -10302,7 +10322,8 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
    }
    vk_pipeline count_experts = ctx->device->pipeline_count_experts;

-    uint32_t expert_count_size = sizeof(uint32_t) * n_as;
+    const size_t expert_data_size = sizeof(uint32_t) *
+        (hoist_row_ids ? hoisted_row_id_words : n_as);

    {
        if (
@@ -10318,8 +10339,8 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
            ctx->prealloc_size_y = y_sz;
            ggml_vk_preallocate_buffers(ctx, subctx);
        }
-        if (ctx->prealloc_size_split_k < expert_count_size) {
-            ctx->prealloc_size_split_k = expert_count_size;
+        if (ctx->prealloc_size_split_k < expert_data_size) {
+            ctx->prealloc_size_split_k = expert_data_size;
            ggml_vk_preallocate_buffers(ctx, subctx);
        }

@@ -10385,18 +10406,23 @@ static void ggml_vk_mul_mat_id_q_f16(ggml_backend_vk_context * ctx, vk_context&
        }
    }
    // Count how many times each expert is used
-    vk_subbuffer expert_count_buf = ggml_vk_subbuffer(ctx, ctx->prealloc_split_k, 0);
+    vk_subbuffer expert_count_buf = { ctx->prealloc_split_k, 0, expert_data_size };
    if (ctx->prealloc_split_k_need_sync) {
        ggml_vk_sync_buffers(ctx, subctx);
    }
    {
-        const std::vector<uint32_t> pc = { (uint32_t)nei0,
+        vk_op_count_experts_push_constants pc = { (uint32_t)nei0,
                                           (uint32_t)nei1,
                                           (uint32_t)(nbi0 / ggml_type_size(ids->type)),
                                           (uint32_t)(nbi1 / ggml_type_size(ids->type)),
-                                           (uint32_t)(get_misalign_bytes(ctx, ids) / ggml_type_size(ids->type)) };
+                                           (uint32_t)(get_misalign_bytes(ctx, ids) / ggml_type_size(ids->type)),
+                                           (uint32_t)n_as,
+                                           uint32_t(hoist_row_ids),
+                                           0, 0 };
+        init_pushconst_fastdiv(pc);
        ggml_vk_dispatch_pipeline(ctx, subctx, count_experts,
-            { vk_subbuffer{ d_ids, ids_buf_offset, ids_sz }, expert_count_buf }, pc, { (uint32_t)n_as, 1, 1});
+            { vk_subbuffer{ d_ids, ids_buf_offset, ids_sz }, expert_count_buf }, pc,
+            { hoist_row_ids ? 1u : (uint32_t)n_as, 1, 1});
    }

    if (x_non_contig) {
@@ -10465,7 +10491,7 @@ 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
+        n_as, nei0, nei1, nbi1 / ggml_type_size(ids->type), ne11, padded_n, hoist_row_ids
    );  // NOLINT

    if (x_non_contig || qx_needs_dequant) {
@@ -17774,20 +17800,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;
@@ -2,6 +2,11 @@

 #extension GL_EXT_control_flow_attributes : enable

+#ifdef USE_SUBGROUPS
+#extension GL_KHR_shader_subgroup_basic : enable
+#extension GL_KHR_shader_subgroup_arithmetic : enable
+#endif
+
 #include "types.glsl"

 layout (push_constant) uniform parameter
@@ -11,6 +16,10 @@ layout (push_constant) uniform parameter
    uint32_t nb00;
    uint32_t nb01;
    uint32_t a_offset;
+    uint32_t n_experts;
+    uint32_t hoist_row_ids;
+    uint32_t ne00mp;
+    uint32_t ne00L;
 } p;

 #define BLOCK_SIZE 256
@@ -21,16 +30,98 @@ layout (binding = 0) readonly buffer A {uint data_a[];};
 layout (binding = 1) writeonly buffer D {uint data_d[];};

 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
+//   [2*n_experts]                 total row count
+//   [2*n_experts + 1,         )   row ids grouped by expert, packed as (i01 << 16) | (i00 & 0xffff)
+// Otherwise only data_d[expert_id] is written, holding that expert's row count.
 void main() {
    const uint expert_id = gl_WorkGroupID.x;
    const uint num_elements = p.ne00 * p.ne01;
    const uint tid = gl_LocalInvocationID.x;

+    if (p.hoist_row_ids != 0) {
+        if (tid < p.n_experts) {
+            vals[tid] = 0;
+        }
+        barrier();
+
+        for (uint idx = tid; idx < num_elements; idx += BLOCK_SIZE) {
+            const uint i01 = fastdiv(idx, p.ne00mp, p.ne00L);
+            const uint i00 = idx - i01 * p.ne00;
+            const uint expert = data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00];
+            if (expert < p.n_experts) {
+                atomicAdd(vals[expert], 1);
+            }
+        }
+        barrier();
+
+#ifdef USE_SUBGROUPS
+        if (gl_SubgroupID == 0) {
+            // pad the trip count so the subgroup ops stay in uniform control flow
+            const uint n_experts_padded = (p.n_experts + gl_SubgroupSize - 1) & ~(gl_SubgroupSize - 1);
+            uint base = 0;
+            for (uint expert = gl_SubgroupInvocationID; expert < n_experts_padded; expert += gl_SubgroupSize) {
+                const bool in_range = expert < p.n_experts;
+                const uint count = in_range ? vals[expert] : 0;
+                const uint offset = base + subgroupExclusiveAdd(count);
+                if (in_range) {
+                    data_d[expert] = count;
+                    data_d[p.n_experts + expert] = offset;
+                    offsets[expert] = offset;
+                    cursors[expert] = 0;
+                }
+                base += subgroupAdd(count);
+            }
+            if (subgroupElect()) {
+                data_d[2 * p.n_experts] = base;
+            }
+        }
+#else
+        if (tid == 0) {
+            uint offset = 0;
+            for (uint expert = 0; expert < p.n_experts; ++expert) {
+                const uint count = vals[expert];
+                data_d[expert] = count;
+                data_d[p.n_experts + expert] = offset;
+                offsets[expert] = offset;
+                cursors[expert] = 0;
+                offset += count;
+            }
+            data_d[2 * p.n_experts] = offset;
+        }
+#endif
+        barrier();
+
+        for (uint idx = tid; idx < num_elements; idx += BLOCK_SIZE) {
+            const uint i01 = fastdiv(idx, p.ne00mp, p.ne00L);
+            const uint i00 = idx - i01 * p.ne00;
+            const uint expert = data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00];
+            if (expert < p.n_experts) {
+                const uint row = atomicAdd(cursors[expert], 1);
+                const uint packed_row_id = (i01 << 16) | (i00 & 0xffffu);
+                data_d[2 * p.n_experts + 1 + offsets[expert] + row] = packed_row_id;
+            }
+        }
+        return;
+    }
+
    uint count = 0;
    for (uint idx = tid; idx < num_elements; idx += BLOCK_SIZE) {
-        const uint i01 = idx / p.ne00;
-        const uint i00 = idx % p.ne00;
+        const uint i01 = fastdiv(idx, p.ne00mp, p.ne00L);
+        const uint i00 = idx - i01 * p.ne00;
        const uint a = data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00];

        count += uint(a == expert_id);
@@ -88,6 +88,9 @@ layout (push_constant) uniform parameter
    uint nei1;
    uint nbi1;
    uint ne11;
+    uint padded_N;
+    uint n_experts;
+    uint hoist_row_ids;
 #else
    uint base_work_group_z;
    uint num_batches;
@@ -214,27 +217,31 @@ void main() {
    const uint loadstride_b = gl_WorkGroupSize.x * LOAD_VEC_B_EFF * LOAD_VEC_BATCH_B / BK;

 #ifdef MUL_MAT_ID
-#ifdef MUL_MAT_ID_USE_SUBGROUPS
-    if (bitCount(p.nei0) == 1) {
-        load_row_ids(expert_idx, true, ic);
+    if (p.hoist_row_ids != 0) {
+        load_row_ids_hoisted(expert_idx, ic);
    } else {
-        load_row_ids(expert_idx, false, ic);
-    }
+#ifdef MUL_MAT_ID_USE_SUBGROUPS
+        if (bitCount(p.nei0) == 1) {
+            load_row_ids(expert_idx, true, ic);
+        } else {
+            load_row_ids(expert_idx, false, ic);
+        }
 #else
-    _ne1 = 0;
-    for (uint ii1 = 0; ii1 < p.nei1 && _ne1 < (ic + 1) * BN; ii1++) {
-        for (uint ii0 = 0; ii0 < p.nei0 && _ne1 < (ic + 1) * BN; ii0++) {
-            if (data_ids[ii1*p.nbi1 + ii0] == expert_idx) {
-                if (_ne1 >= ic * BN) {
-                    row_ids[_ne1 - ic * BN] = u16vec2(ii0, ii1);
+        _ne1 = 0;
+        for (uint ii1 = 0; ii1 < p.nei1 && _ne1 < (ic + 1) * BN; ii1++) {
+            for (uint ii0 = 0; ii0 < p.nei0 && _ne1 < (ic + 1) * BN; ii0++) {
+                if (data_ids[ii1*p.nbi1 + ii0] == expert_idx) {
+                    if (_ne1 >= ic * BN) {
+                        row_ids[_ne1 - ic * BN] = u16vec2(ii0, ii1);
+                    }
+                    _ne1++;
                }
-                _ne1++;
            }
        }
-    }

-    barrier();
+        barrier();
 #endif
+    }

    // Workgroup has no work
    if (ic * BN >= _ne1) return;
@@ -67,6 +67,10 @@ layout (push_constant) uniform parameter
 #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;


@@ -225,6 +229,23 @@ void load_row_ids(uint expert_idx, bool nei0_is_pow2, uint ic) {
    }
    barrier();
 }
+
+void load_row_ids_hoisted(uint expert_idx, uint ic) {
+    _ne1 = uint(data_expert_count[expert_idx]);
+
+    const uint tile_begin = ic * BN;
+    const uint tile_count = tile_begin < _ne1 ? min(BN, _ne1 - tile_begin) : 0;
+    const uint expert_offset = uint(data_expert_count[p.n_experts + expert_idx]);
+    const uint row_ids_offset = 2 * p.n_experts + 1 + expert_offset + tile_begin;
+
+    for (uint i = gl_LocalInvocationIndex; i < tile_count; i += BLOCK_SIZE) {
+        const uint packed_row_id = uint(data_expert_count[row_ids_offset + i]);
+        const uint ii0 = packed_row_id & 0xffffu;
+        const uint ii1 = packed_row_id >> 16;
+        row_ids[i] = u16vec4(fastmod(ii0, p.ne11), ii1, ii0, 0);
+    }
+    barrier();
+}
 #endif

 void main() {
@@ -266,7 +287,9 @@ void main() {
    const uint ik = gl_WorkGroupID.x / blocks_m;

 #ifdef MUL_MAT_ID
-    if (bitCount(p.nei0) == 1) {
+    if (p.hoist_row_ids != 0) {
+        load_row_ids_hoisted(expert_idx, ic);
+    } else if (bitCount(p.nei0) == 1) {
        load_row_ids(expert_idx, true, ic);
    } else {
        load_row_ids(expert_idx, false, ic);
@@ -71,4 +71,19 @@ void load_row_ids(uint expert_idx, bool nei0_is_pow2, uint ic) {
    barrier();
 }
 #endif // MUL_MAT_ID_USE_SUBGROUPS
+
+void load_row_ids_hoisted(uint expert_idx, uint ic) {
+    _ne1 = uint(data_expert_count[expert_idx]);
+
+    const uint tile_begin = ic * BN;
+    const uint tile_count = tile_begin < _ne1 ? min(BN, _ne1 - tile_begin) : 0;
+    const uint expert_offset = uint(data_expert_count[p.n_experts + expert_idx]);
+    const uint row_ids_offset = 2 * p.n_experts + 1 + expert_offset + tile_begin;
+
+    for (uint i = gl_LocalInvocationIndex; i < tile_count; i += BLOCK_SIZE) {
+        const uint packed_row_id = uint(data_expert_count[row_ids_offset + i]);
+        row_ids[i] = u16vec2(packed_row_id & 0xffffu, packed_row_id >> 16);
+    }
+    barrier();
+}
 #endif // MUL_MAT_ID
@@ -56,6 +56,9 @@ layout (push_constant) uniform parameter
    uint nei1;
    uint nbi1;
    uint ne11;
+    uint padded_N;
+    uint n_experts;
+    uint hoist_row_ids;
 #else
    uint base_work_group_z;
    uint num_batches;
@@ -157,27 +160,31 @@ void main() {
    const uint loadstride_b = BLOCK_SIZE * LOAD_VEC_B / BK;

 #ifdef MUL_MAT_ID
-#ifdef MUL_MAT_ID_USE_SUBGROUPS
-    if (bitCount(p.nei0) == 1) {
-        load_row_ids(expert_idx, true, ic);
+    if (p.hoist_row_ids != 0) {
+        load_row_ids_hoisted(expert_idx, ic);
    } else {
-        load_row_ids(expert_idx, false, ic);
-    }
+#ifdef MUL_MAT_ID_USE_SUBGROUPS
+        if (bitCount(p.nei0) == 1) {
+            load_row_ids(expert_idx, true, ic);
+        } else {
+            load_row_ids(expert_idx, false, ic);
+        }
 #else
-    _ne1 = 0;
-    for (uint ii1 = 0; ii1 < p.nei1 && _ne1 < (ic + 1) * BN; ii1++) {
-        for (uint ii0 = 0; ii0 < p.nei0 && _ne1 < (ic + 1) * BN; ii0++) {
-            if (data_ids[ii1*p.nbi1 + ii0] == expert_idx) {
-                if (_ne1 >= ic * BN) {
-                    row_ids[_ne1 - ic * BN] = u16vec2(ii0, ii1);
+        _ne1 = 0;
+        for (uint ii1 = 0; ii1 < p.nei1 && _ne1 < (ic + 1) * BN; ii1++) {
+            for (uint ii0 = 0; ii0 < p.nei0 && _ne1 < (ic + 1) * BN; ii0++) {
+                if (data_ids[ii1*p.nbi1 + ii0] == expert_idx) {
+                    if (_ne1 >= ic * BN) {
+                        row_ids[_ne1 - ic * BN] = u16vec2(ii0, ii1);
+                    }
+                    _ne1++;
                }
-                _ne1++;
            }
        }
-    }

-    barrier();
+        barrier();
 #endif
+    }

    // Workgroup has no work
    if (ic * BN >= _ne1) return;
@@ -1039,6 +1039,7 @@ void process_shaders() {
    string_to_spv("cumsum_multipass2_f32", "cumsum_multipass2.comp", merge_maps(base_dict, {{"A_TYPE", "float"}, {"D_TYPE", "float"}}));

    string_to_spv("count_experts", "count_experts.comp", merge_maps(base_dict, {{"A_TYPE", "uint"}, {"D_TYPE", "uint"}}));
+    string_to_spv("count_experts_subgroup", "count_experts.comp", merge_maps(base_dict, {{"A_TYPE", "uint"}, {"D_TYPE", "uint"}, {"USE_SUBGROUPS", "1"}}));

    for (std::string dim_str : {"", "_3d"}) {
        for (bool bda : {false, true}) {
@@ -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;
@@ -231,10 +231,10 @@ llama_context::llama_context(

    cparams.fused_gdn_ar = true;
    cparams.fused_gdn_ch = true;
-    cparams.auto_fgdn    = true;
+    cparams.auto_fgdn    = false;

-    cparams.fused_lid    = true;
-    cparams.auto_flid    = true;
+    cparams.fused_lid = true;
+    cparams.auto_flid = false;

    cparams.fused_dsv4_hc_pre  = true;
    cparams.fused_dsv4_hc_comb = true;
@@ -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}));
@@ -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
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
ravel7524 GitHub Jeff Bolz	90c26fcd4b	Vulkan: add hoisting support for row IDs and expert count in shaders (#26686 ) * vulkan: add hoisting support for row IDs and expert count in shaders * use hoisted row ids in coopmat2 * vulkan: address review feedback on count_experts - use vk_op_count_experts_push_constants instead of a raw uint vector - apply the fastdiv trick to the ne00 div/mod in count_experts - compute the per-expert offsets with subgroupExclusiveAdd when the device supports it, keeping the serial path as fallback - document the data_d layout and the hoisted_row_id_words bound - drop a leftover debug print in ggml_vk_matmul_id * vulkan: use init_pushconst_fastdiv for count_experts push constants * vulkan: refine comments for row ID hoisting and data layout in count_experts shader * Whitespace --------- Co-authored-by: Jeff Bolz <jbolz@nvidia.com>	2026-08-28 16:52:49 +02:00
Georgi GerganovandGitHub	8663224818	context : disable non-fused GDN and LID ops (#27877 )	2026-08-28 16:34:26 +03:00