qwen4exp: fix sparse-attention block selection

Build QSA blocks per sequence in token order and select complete blocks
before expanding them to cache cells. Keep only the incomplete tail
unconditionally visible and rotate pooled keys with the first token's
full M-RoPE position.

This prevents unified-cache sequences from sharing pooled indexer keys
and avoids replacing padded tail entries with extra history tokens.

Synthetic Qwen4 architecture, exact mask, F16 and Q8_0 state,
sequence-copy, Metal, and AddressSanitizer checks pass.

Assisted-by: Codex

qwen4exp: support independent PLE embedding widths

Size the PLE key and value projections from the concatenated n-gram
embedding instead of assuming it matches the model hidden width.
Validate the head count before narrowing it to the stored type.

Add a synthetic PLE model with a 64-wide embedding and a 256-wide
hidden state, then verify inference and model roundtrip.

Assisted-by: Codex

qwen4exp: validate model metadata

Reject invalid GDN, hyper-connection, QSA, and PLE dimensions during
model loading instead of aborting later while building the graph.
Validate PLE array lengths before copying them into fixed storage.

The released configuration and synthetic Qwen4 architecture tests pass.

Assisted-by: Codex

qwen4exp: update indexer cache after sequence copies

Treat cached indexer keys as unrotated data and apply pending cache
updates alongside the attention and recurrent state. This copies indexer
data during non-unified cross-stream sequence copies without applying
RoPE shifts to raw keys.

Assisted-by: Codex

qwen4exp: enable recurrent state rollback

Assisted-by: Codex

qwen4exp: disable tensor split

Assisted-by: Codex

metal: align dynamic threadgroup memory

Assisted-by: Codex

metal: widen expert matmul thread index

Assisted-by: Codex

common/arg.cpp

@@ -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.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;  }
+            /**/ 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;  }
             else { throw std::invalid_argument("invalid value"); }
         }
     ).set_env("LLAMA_ARG_TENSOR_READ_LAZY"));

common/common.cpp

@@ -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.lazy_mode = params.lazy_mode;
+    mparams.tensor_read_lazy = params.tensor_read_lazy;
     mparams.tensor_split    = params.tensor_split;
     mparams.check_tensors   = params.check_tensors;
     mparams.use_extra_bufts = !params.no_extra_bufts;

common/common.h

@@ -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_lazy_mode lazy_mode = LLAMA_LAZY_MODE_AUTO; // on-demand reading of tensors marked by the arch
+    enum llama_tensor_read_lazy tensor_read_lazy = LLAMA_TENSOR_READ_LAZY_AUTO; // on-demand reading of tensors marked by the arch
 
     common_cpu_params cpuparams;
     common_cpu_params cpuparams_batch;

docs/backend/snapdragon/windows.md

@@ -24,18 +24,7 @@ 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 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.
+## Install the latest Adreno OpenCL SDK
 
 Either use the trimmed down version (optimized for CI) from
 

ggml/src/ggml-cpu/ggml-cpu.c

@@ -2936,13 +2936,12 @@ 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 * ne13;
+                        cur += ggml_type_size(node->src[0]->type) * ne10 * ne11 * ne12;
 
                     } break;
                 case GGML_OP_TOP_K:

ggml/src/ggml-cpu/ops.cpp

@@ -7267,21 +7267,18 @@ 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), for all batches
+        // permute source data (src1) from (Sw x Sh x Cin) to (Cin x Sw x Sh)
         {
             kernel_t * const wdata = (kernel_t *) params->wdata + nk;
-            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];
-                            }
+            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];
                         }
                     }
                 }
@@ -7308,27 +7305,24 @@ 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 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;
+        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);
                         }
+                        dst_data[(i11*stride + i01)*ne0 + i10*stride + i00] += v;
                     }
                 }
             }

ggml/src/ggml-metal/ggml-metal-device.m

@@ -800,7 +800,7 @@ void ggml_metal_encoder_set_buffer(ggml_metal_encoder_t encoder, struct ggml_met
 }
 
 void ggml_metal_encoder_set_threadgroup_memory_size(ggml_metal_encoder_t encoder, size_t size, int idx) {
-    [encoder->obj setThreadgroupMemoryLength:size atIndex:idx];
+    [encoder->obj setThreadgroupMemoryLength:GGML_PAD(size, 16) atIndex:idx];
 }
 
 void ggml_metal_encoder_dispatch_threadgroups(ggml_metal_encoder_t encoder, int tg0, int tg1, int tg2, int tptg0, int tptg1, int tptg2) {

ggml/src/ggml-metal/ggml-metal-impl.h

@@ -660,7 +660,6 @@ typedef struct {
     uint64_t nb0;
     uint64_t nb1;
     uint64_t nb2;
-    uint64_t nb3;
 } ggml_metal_kargs_conv_transpose_2d;
 
 typedef struct {

ggml/src/ggml-metal/ggml-metal-ops.cpp

@@ -4645,7 +4645,6 @@ 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,
@@ -4658,7 +4657,6 @@ 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);
@@ -4673,7 +4671,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 * N, KW, KH, 1);
+    ggml_metal_encoder_dispatch_threadgroups(enc, OW, OH, OC, KW, KH, 1);
 
     return 1;
 }

ggml/src/ggml-metal/kernels/conv.metal

@@ -366,8 +366,7 @@ kernel void kernel_conv_transpose_2d(
 
     const int64_t out_x = tgpig[0];
     const int64_t out_y = tgpig[1];
-    const int64_t batch = tgpig[2] / args.OC;
-    const int64_t out_c = tgpig[2] % args.OC;
+    const int64_t out_c = tgpig[2];
 
     const int64_t kw = tpitg[0];
     const int64_t kh = tpitg[1];
@@ -391,7 +390,7 @@ kernel void kernel_conv_transpose_2d(
 
         if (in_x >= args.IW) continue;
 
-        const int64_t input_idx = (args.IW * args.IH) * (args.IC * batch + in_c) + (args.IW) * in_y + in_x;
+        const int64_t input_idx = (args.IW * args.IH) * 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];
@@ -409,7 +408,7 @@ kernel void kernel_conv_transpose_2d(
             total += shared_sum[i];
         }
 
-        device float * dst_ptr = (device float *) (dst + batch*args.nb3 + out_c*args.nb2 + out_y * args.nb1 + out_x*args.nb0);
+        device float * dst_ptr = (device float *) (dst + out_x*args.nb0 + out_y * args.nb1 + out_c*args.nb2);
         dst_ptr[0] = total;
     }
 }

ggml/src/ggml-metal/kernels/mul_mm.metal

@@ -435,7 +435,7 @@ kernel void kernel_mul_mm_id(
         device       char * dst,
         threadgroup  char * shmem [[threadgroup(0)]],
         uint3  tgpig[[threadgroup_position_in_grid]],
-        ushort tiitg[[thread_index_in_threadgroup]],
+        uint   tiitg[[thread_index_in_threadgroup]],
         ushort tiisg[[thread_index_in_simdgroup]],
         ushort sgitg[[simdgroup_index_in_threadgroup]]) {
     threadgroup S0 * sa = (threadgroup S0 *)(shmem);

ggml/src/ggml-vulkan/ggml-vulkan.cpp

@@ -1348,8 +1348,6 @@ 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;
@@ -1430,10 +1428,6 @@ 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 {
@@ -1612,10 +1606,6 @@ 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;
@@ -5849,11 +5839,7 @@ 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);
 
-    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);
-    }
+    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;
@@ -8984,13 +8970,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) {
+        uint32_t padded_n) {
     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, padded_n };
     ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { a, b, d, ids, expert_count_buf }, pc, { m, nei1, n_as });
 }
 
@@ -10176,12 +10162,6 @@ 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;
@@ -10322,8 +10302,7 @@ 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;
 
-    const size_t expert_data_size = sizeof(uint32_t) *
-        (hoist_row_ids ? hoisted_row_id_words : n_as);
+    uint32_t expert_count_size = sizeof(uint32_t) * n_as;
 
     {
         if (
@@ -10339,8 +10318,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_data_size) {
-            ctx->prealloc_size_split_k = expert_data_size;
+        if (ctx->prealloc_size_split_k < expert_count_size) {
+            ctx->prealloc_size_split_k = expert_count_size;
             ggml_vk_preallocate_buffers(ctx, subctx);
         }
 
@@ -10406,23 +10385,18 @@ 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 = { ctx->prealloc_split_k, 0, expert_data_size };
+    vk_subbuffer expert_count_buf = ggml_vk_subbuffer(ctx, ctx->prealloc_split_k, 0);
     if (ctx->prealloc_split_k_need_sync) {
         ggml_vk_sync_buffers(ctx, subctx);
     }
     {
-        vk_op_count_experts_push_constants pc = { (uint32_t)nei0,
+        const std::vector<uint32_t> 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)n_as,
-                                           uint32_t(hoist_row_ids),
-                                           0, 0 };
-        init_pushconst_fastdiv(pc);
+                                           (uint32_t)(get_misalign_bytes(ctx, ids) / ggml_type_size(ids->type)) };
         ggml_vk_dispatch_pipeline(ctx, subctx, count_experts,
-            { vk_subbuffer{ d_ids, ids_buf_offset, ids_sz }, expert_count_buf }, pc,
-            { hoist_row_ids ? 1u : (uint32_t)n_as, 1, 1});
+            { vk_subbuffer{ d_ids, ids_buf_offset, ids_sz }, expert_count_buf }, pc, { (uint32_t)n_as, 1, 1});
     }
 
     if (x_non_contig) {
@@ -10491,7 +10465,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, hoist_row_ids
+        n_as, nei0, nei1, nbi1 / ggml_type_size(ids->type), ne11, padded_n
     );  // NOLINT
 
     if (x_non_contig || qx_needs_dequant) {
@@ -17800,32 +17774,20 @@ static void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph *
         return;
     }
 
-    auto const &is_empty = [](const ggml_tensor * node) -> bool {
+    auto const &is_empty = [](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 = [&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;
-        };
+    auto const &is_src_of = [](const ggml_tensor *dst, const ggml_tensor *src) -> bool {
         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
-        if (base(dst) == base(src)) {
+        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) {
             return true;
         }
         return false;

ggml/src/ggml-vulkan/vulkan-shaders/count_experts.comp

@@ -2,11 +2,6 @@
 
 #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
@@ -16,10 +11,6 @@ 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
@@ -30,98 +21,16 @@ 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 = fastdiv(idx, p.ne00mp, p.ne00L);
-        const uint i00 = idx - i01 * p.ne00;
+        const uint i01 = idx / p.ne00;
+        const uint i00 = idx % p.ne00;
         const uint a = data_a[p.a_offset + i01 * p.nb01 + i00 * p.nb00];
 
         count += uint(a == expert_id);

ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp

@@ -88,9 +88,6 @@ 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;
@@ -217,32 +214,28 @@ void main() {
     const uint loadstride_b = gl_WorkGroupSize.x * LOAD_VEC_B_EFF * LOAD_VEC_BATCH_B / BK;
 
 #ifdef MUL_MAT_ID
-    if (p.hoist_row_ids != 0) {
-        load_row_ids_hoisted(expert_idx, ic);
-    } else {
 #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);
-        }
+    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++;
+    _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++;
             }
         }
-
-        barrier();
-#endif
     }
 
+    barrier();
+#endif
+
     // Workgroup has no work
     if (ic * BN >= _ne1) return;
 #endif

ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_cm2.comp

@@ -67,10 +67,6 @@ 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;
 
 
@@ -229,23 +225,6 @@ 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() {
@@ -287,9 +266,7 @@ void main() {
     const uint ik = gl_WorkGroupID.x / blocks_m;
 
 #ifdef MUL_MAT_ID
-    if (p.hoist_row_ids != 0) {
-        load_row_ids_hoisted(expert_idx, ic);
-    } else if (bitCount(p.nei0) == 1) {
+    if (bitCount(p.nei0) == 1) {
         load_row_ids(expert_idx, true, ic);
     } else {
         load_row_ids(expert_idx, false, ic);

ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_id_funcs.glsl

@@ -71,19 +71,4 @@ 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

ggml/src/ggml-vulkan/vulkan-shaders/mul_mmq.comp

@@ -56,9 +56,6 @@ 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;
@@ -160,32 +157,28 @@ void main() {
     const uint loadstride_b = BLOCK_SIZE * LOAD_VEC_B / BK;
 
 #ifdef MUL_MAT_ID
-    if (p.hoist_row_ids != 0) {
-        load_row_ids_hoisted(expert_idx, ic);
-    } else {
 #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);
-        }
+    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++;
+    _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++;
             }
         }
-
-        barrier();
-#endif
     }
 
+    barrier();
+#endif
+
     // Workgroup has no work
     if (ic * BN >= _ne1) return;
 #endif

ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp

@@ -1039,7 +1039,6 @@ 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}) {

include/llama.h

@@ -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_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_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_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_lazy_mode lazy_mode; // on-demand reading of tensors marked by the arch
+        enum llama_tensor_read_lazy tensor_read_lazy; // 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;

scripts/snapdragon/build.py

@@ -11,8 +11,6 @@ import platform
 import shutil
 import logging
 
-from sdk import validate_windows_sdks
-
 logger = logging.getLogger("build")
 
 
@@ -67,13 +65,6 @@ 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:
@@ -129,6 +120,12 @@ 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...")
@@ -261,6 +258,3 @@ 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)

scripts/snapdragon/sdk.py

@@ -1,62 +0,0 @@
-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."
-        )

scripts/snapdragon/setup-sdk.py

@@ -1,233 +0,0 @@
-#!/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)

src/llama-arch.cpp

@@ -1100,6 +1100,7 @@ bool llm_arch_supports_rs_rollback(const llm_arch & arch) {
     switch (arch) {
         case LLM_ARCH_QWEN35:
         case LLM_ARCH_QWEN35MOE:
+        case LLM_ARCH_QWEN4EXP:
         case LLM_ARCH_DEEPSEEK4:
         case LLM_ARCH_NEMOTRON_H:
         case LLM_ARCH_NEMOTRON_H_MOE:
@@ -1141,6 +1142,7 @@ bool llm_arch_supports_sm_tensor(const llm_arch & arch) {
         case LLM_ARCH_KIMI_LINEAR:
         case LLM_ARCH_BAILINGMOE3:
         case LLM_ARCH_KIMI_K3:
+        case LLM_ARCH_QWEN4EXP:
         case LLM_ARCH_QWEN3TTS:
             return false;
         default:

src/llama-context.cpp

@@ -231,10 +231,10 @@ llama_context::llama_context(
 
     cparams.fused_gdn_ar = true;
     cparams.fused_gdn_ch = true;
-    cparams.auto_fgdn    = false;
+    cparams.auto_fgdn    = true;
 
-    cparams.fused_lid = true;
-    cparams.auto_flid = false;
+    cparams.fused_lid    = true;
+    cparams.auto_flid    = true;
 
     cparams.fused_dsv4_hc_pre  = true;
     cparams.fused_dsv4_hc_comb = true;

src/llama-memory-hybrid-idx.cpp

@@ -9,7 +9,9 @@
 #include <cassert>
 #include <cmath>
 #include <iterator>
+#include <limits>
 #include <stdexcept>
+#include <tuple>
 
 //
 // llama_memory_hybrid_idx
@@ -47,6 +49,7 @@ llama_memory_hybrid_idx::llama_memory_hybrid_idx(
     hparams_idx(model.hparams),
     mem_idx(filter_idx == nullptr ? nullptr : [&] {
         // MQA with a single key head of indexer_head_size, as llama_kv_cache_dsa shapes its own
+        hparams_idx.rope_type = LLAMA_ROPE_TYPE_NONE;
         std::fill(hparams_idx.n_head_kv_arr.begin(), hparams_idx.n_head_kv_arr.end(), 1);
         hparams_idx.n_embd_head_k_full = model.hparams.indexer_head_size;
 
@@ -137,6 +140,8 @@ void llama_memory_hybrid_idx::clear(bool data) {
     if (mem_idx) {
         mem_idx->clear(data);
     }
+
+    qsa_histories.clear();
 }
 
 bool llama_memory_hybrid_idx::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_pos p1) {
@@ -149,15 +154,96 @@ bool llama_memory_hybrid_idx::seq_rm(llama_seq_id seq_id, llama_pos p0, llama_po
         mem_idx->seq_rm(seq_id, p0, p1);
     }
 
-    return get_mem_attn()->seq_rm(seq_id, p0, p1);
+    const bool res = get_mem_attn()->seq_rm(seq_id, p0, p1);
+    if (!res) {
+        return false;
+    }
+
+    auto remove = [&](qsa_history & history) {
+        history.erase(std::remove_if(history.begin(), history.end(), [&](const qsa_token & token) {
+            return (p0 < 0 || token.pos[0] >= p0) && (p1 < 0 || token.pos[0] < p1);
+        }), history.end());
+    };
+
+    if (seq_id < 0) {
+        for (auto & item : qsa_histories) {
+            remove(item.second);
+        }
+    } else {
+        auto it = qsa_histories.find(seq_id);
+        if (it != qsa_histories.end()) {
+            remove(it->second);
+            if (it->second.empty()) {
+                qsa_histories.erase(it);
+            }
+        }
+    }
+
+    return true;
 }
 
 void llama_memory_hybrid_idx::seq_cp(llama_seq_id seq_id_src, llama_seq_id seq_id_dst, llama_pos p0, llama_pos p1) {
+    if (seq_id_src == seq_id_dst) {
+        return;
+    }
+
+    qsa_history copied;
+    const auto & cells_src = get_mem_attn()->get_cells(seq_id_src);
+    const auto & cells_dst = get_mem_attn()->get_cells(seq_id_dst);
+    const bool replace = &cells_src != &cells_dst;
+    const auto src = qsa_histories.find(seq_id_src);
+    if (src != qsa_histories.end()) {
+        using pos_key = std::tuple<llama_pos, llama_pos, llama_pos>;
+        std::map<pos_key, std::vector<bool>> cells_by_pos;
+        for (uint32_t cell = 0; cell < cells_src.size(); ++cell) {
+            if (cells_src.is_empty(cell) || !cells_src.seq_has(cell, seq_id_src)) {
+                continue;
+            }
+
+            const llama_pos pos = cells_src.pos_get(cell);
+            if ((p0 >= 0 && pos < p0) || (p1 >= 0 && pos >= p1)) {
+                continue;
+            }
+
+            const auto & ext = cells_src.ext_get(cell);
+            cells_by_pos[{ pos, ext.y, ext.x }].push_back(!replace && cells_src.seq_has(cell, seq_id_dst));
+        }
+
+        std::map<pos_key, size_t> next_cell;
+        for (const auto & token : src->second) {
+            if ((p0 >= 0 && token.pos[0] < p0) || (p1 >= 0 && token.pos[0] >= p1)) {
+                continue;
+            }
+
+            const pos_key key = { token.pos[0], token.pos[1], token.pos[2] };
+            auto cells = cells_by_pos.find(key);
+            if (cells == cells_by_pos.end()) {
+                continue;
+            }
+
+            size_t & index = next_cell[key];
+            if (index < cells->second.size() && !cells->second[index++]) {
+                copied.push_back(token);
+            }
+        }
+    }
+
     llama_memory_hybrid::seq_cp(seq_id_src, seq_id_dst, p0, p1);
 
     if (mem_idx) {
         mem_idx->seq_cp(seq_id_src, seq_id_dst, p0, p1);
     }
+
+    if (replace) {
+        if (copied.empty()) {
+            qsa_histories.erase(seq_id_dst);
+        } else {
+            qsa_histories[seq_id_dst] = std::move(copied);
+        }
+    } else if (!copied.empty()) {
+        auto & dst = qsa_histories[seq_id_dst];
+        dst.insert(dst.end(), copied.begin(), copied.end());
+    }
 }
 
 void llama_memory_hybrid_idx::seq_keep(llama_seq_id seq_id) {
@@ -166,6 +252,13 @@ void llama_memory_hybrid_idx::seq_keep(llama_seq_id seq_id) {
     if (mem_idx) {
         mem_idx->seq_keep(seq_id);
     }
+
+    auto it = qsa_histories.find(seq_id);
+    qsa_history keep = it == qsa_histories.end() ? qsa_history{} : std::move(it->second);
+    qsa_histories.clear();
+    if (!keep.empty()) {
+        qsa_histories.emplace(seq_id, std::move(keep));
+    }
 }
 
 void llama_memory_hybrid_idx::seq_add(llama_seq_id seq_id, llama_pos p0, llama_pos p1, llama_pos shift) {
@@ -174,6 +267,15 @@ void llama_memory_hybrid_idx::seq_add(llama_seq_id seq_id, llama_pos p0, llama_p
     if (mem_idx) {
         mem_idx->seq_add(seq_id, p0, p1, shift);
     }
+
+    auto it = qsa_histories.find(seq_id);
+    if (it != qsa_histories.end()) {
+        for (auto & token : it->second) {
+            if ((p0 < 0 || token.pos[0] >= p0) && (p1 < 0 || token.pos[0] < p1)) {
+                token.pos[0] += shift;
+            }
+        }
+    }
 }
 
 void llama_memory_hybrid_idx::seq_div(llama_seq_id seq_id, llama_pos p0, llama_pos p1, int d) {
@@ -182,6 +284,15 @@ void llama_memory_hybrid_idx::seq_div(llama_seq_id seq_id, llama_pos p0, llama_p
     if (mem_idx) {
         mem_idx->seq_div(seq_id, p0, p1, d);
     }
+
+    auto it = qsa_histories.find(seq_id);
+    if (it != qsa_histories.end()) {
+        for (auto & token : it->second) {
+            if ((p0 < 0 || token.pos[0] >= p0) && (p1 < 0 || token.pos[0] < p1)) {
+                token.pos[0] /= d;
+            }
+        }
+    }
 }
 
 std::map<ggml_backend_buffer_type_t, size_t> llama_memory_hybrid_idx::memory_breakdown() const {
@@ -205,8 +316,28 @@ void llama_memory_hybrid_idx::state_write(llama_io_write_i & io, llama_seq_id se
         if (mem_idx) {
             mem_idx->state_write(io, seq_id, flags);
         }
-    }
 
+        uint32_t n_histories = 0;
+        if (seq_id < 0) {
+            n_histories = (uint32_t) qsa_histories.size();
+        } else if (qsa_histories.count(seq_id) != 0) {
+            n_histories = 1;
+        }
+        io.write(&n_histories, sizeof(n_histories));
+
+        for (const auto & item : qsa_histories) {
+            if (seq_id >= 0 && item.first != seq_id) {
+                continue;
+            }
+
+            io.write(&item.first, sizeof(item.first));
+            const uint64_t n_tokens = item.second.size();
+            io.write(&n_tokens, sizeof(n_tokens));
+            for (const auto & token : item.second) {
+                io.write(token.pos.data(), sizeof(token.pos));
+            }
+        }
+    }
 }
 
 void llama_memory_hybrid_idx::state_read(llama_io_read_i & io, llama_seq_id seq_id, llama_state_seq_flags flags) {
@@ -230,6 +361,34 @@ void llama_memory_hybrid_idx::state_read(llama_io_read_i & io, llama_seq_id seq_
             if (mem_idx) {
                 mem_idx->state_read_sinfo(io, seq_id, flags, nullptr, &sinfos_attn);
             }
+
+            uint32_t n_histories;
+            io.read(&n_histories, sizeof(n_histories));
+            if (n_histories > LLAMA_MAX_SEQ) {
+                throw std::runtime_error("invalid QSA history count");
+            }
+
+            if (seq_id < 0) {
+                qsa_histories.clear();
+            } else {
+                qsa_histories.erase(seq_id);
+            }
+
+            for (uint32_t ih = 0; ih < n_histories; ++ih) {
+                llama_seq_id stored_seq;
+                uint64_t n_tokens;
+                io.read(&stored_seq, sizeof(stored_seq));
+                io.read(&n_tokens, sizeof(n_tokens));
+                if (stored_seq < 0 || stored_seq >= LLAMA_MAX_SEQ || n_tokens > get_mem_attn()->get_size()) {
+                    throw std::runtime_error("invalid QSA history");
+                }
+
+                auto & history = qsa_histories[seq_id < 0 ? stored_seq : seq_id];
+                history.resize(n_tokens);
+                for (auto & token : history) {
+                    io.read(token.pos.data(), sizeof(token.pos));
+                }
+            }
         }
 
     } catch (...) {
@@ -255,12 +414,156 @@ void llama_memory_hybrid_idx::state_drop(llama_seq_id seq_id) {
     if (mem_idx) {
         mem_idx->seq_rm(seq_id, -1, -1);
     }
+
+    qsa_histories.erase(seq_id);
 }
 
 llama_kv_cache * llama_memory_hybrid_idx::get_mem_idx() const {
     return mem_idx.get();
 }
 
+void llama_memory_hybrid_idx::commit_qsa_tokens(const llama_ubatch & ubatch) {
+    if (!mem_idx) {
+        return;
+    }
+
+    for (uint32_t i = 0; i < ubatch.n_tokens; ++i) {
+        qsa_token token = {};
+        if (ubatch.token) {
+            token.pos = { ubatch.pos[i], ubatch.pos[i], ubatch.pos[i], 0 };
+        } else {
+            for (uint32_t ip = 0; ip < token.pos.size(); ++ip) {
+                token.pos[ip] = ip < ubatch.n_pos ? ubatch.pos[i + ip*ubatch.n_tokens] : ubatch.pos[i];
+            }
+        }
+
+        for (int32_t is = 0; is < ubatch.n_seq_id[i]; ++is) {
+            qsa_histories[ubatch.seq_id[i][is]].push_back(token);
+        }
+    }
+}
+
+void llama_memory_hybrid_idx::set_input_qsa(
+        ggml_tensor *       block_cells,
+        ggml_tensor *       block_pos,
+        ggml_tensor *       block_mask,
+        ggml_tensor *       selected,
+        const ggml_tensor * kq_mask,
+        const llama_ubatch * ubatch,
+        uint32_t            ratio,
+        uint32_t            block_topk) const {
+    GGML_ASSERT(ggml_backend_buffer_is_host(block_cells->buffer));
+    GGML_ASSERT(ggml_backend_buffer_is_host(block_pos->buffer));
+    GGML_ASSERT(ggml_backend_buffer_is_host(block_mask->buffer));
+    GGML_ASSERT(ggml_backend_buffer_is_host(selected->buffer));
+    GGML_ASSERT(ggml_backend_buffer_is_host(kq_mask->buffer));
+
+    const int64_t n_blocks = block_cells->ne[1];
+    const int64_t n_tokens = ubatch->n_tokens;
+    const int64_t n_pos    = block_pos->ne[1];
+    const int64_t n_kv     = selected->ne[0];
+
+    GGML_ASSERT(block_cells->type == GGML_TYPE_I32);
+    GGML_ASSERT(block_pos->type   == GGML_TYPE_I32);
+    GGML_ASSERT(block_mask->type  == GGML_TYPE_F32);
+    GGML_ASSERT(selected->type    == GGML_TYPE_F32);
+    GGML_ASSERT(block_cells->ne[0] == ratio && block_cells->ne[2] == n_tokens);
+    GGML_ASSERT(block_pos->ne[0] == n_blocks && block_pos->ne[2] == n_tokens);
+    GGML_ASSERT(block_mask->ne[0] == n_blocks && block_mask->ne[1] == n_tokens);
+    GGML_ASSERT(selected->ne[1] == n_tokens);
+    GGML_ASSERT(kq_mask->ne[0] == n_kv);
+
+    int32_t * cell_data     = (int32_t *) block_cells->data;
+    int32_t * pos_data      = (int32_t *) block_pos->data;
+    float *   mask_data     = (float *) block_mask->data;
+    float *   selected_data = (float *) selected->data;
+
+    std::fill(cell_data, cell_data + ggml_nelements(block_cells), 0);
+    std::fill(pos_data, pos_data + ggml_nelements(block_pos), 0);
+    std::fill(mask_data, mask_data + ggml_nelements(block_mask), -INFINITY);
+    std::fill(selected_data, selected_data + ggml_nelements(selected), 0.0f);
+
+    auto mask_visible = [&](int64_t query, uint32_t cell) {
+        const int64_t index = query*n_kv + cell;
+        if (kq_mask->type == GGML_TYPE_F16) {
+            return std::isfinite(ggml_fp16_to_fp32(((const ggml_fp16_t *) kq_mask->data)[index]));
+        }
+        return std::isfinite(((const float *) kq_mask->data)[index]);
+    };
+
+    for (int64_t iq = 0; iq < n_tokens; ++iq) {
+        const llama_seq_id seq_id = ubatch->seq_id[iq][0];
+        const auto found = qsa_histories.find(seq_id);
+        if (found == qsa_histories.end()) {
+            continue;
+        }
+
+        const auto & cells = get_mem_attn()->get_cells(seq_id);
+        using pos_key = std::tuple<llama_pos, llama_pos, llama_pos>;
+        std::map<pos_key, std::vector<uint32_t>> cells_by_pos;
+        for (uint32_t cell = 0; cell < cells.size() && cell < (uint32_t) n_kv; ++cell) {
+            if (cells.is_empty(cell) || !cells.seq_has(cell, seq_id)) {
+                continue;
+            }
+
+            const auto & ext = cells.ext_get(cell);
+            cells_by_pos[{ cells.pos_get(cell), ext.y, ext.x }].push_back(cell);
+        }
+
+        std::map<pos_key, size_t> next_cell;
+        std::vector<std::pair<const qsa_token *, uint32_t>> visible;
+        for (const auto & token : found->second) {
+            const pos_key key = { token.pos[0], token.pos[1], token.pos[2] };
+            auto cells = cells_by_pos.find(key);
+            if (cells == cells_by_pos.end()) {
+                continue;
+            }
+
+            size_t & index = next_cell[key];
+            if (index >= cells->second.size()) {
+                continue;
+            }
+
+            const uint32_t cell = cells->second[index++];
+            if (mask_visible(iq, cell)) {
+                visible.emplace_back(&token, cell);
+            }
+        }
+
+        const size_t n_complete = visible.size()/ratio;
+        const size_t n_write = std::min<size_t>(n_complete, n_blocks);
+        std::vector<uint8_t> used_cells(n_kv, 0);
+        for (size_t ib = 0; ib < n_write; ++ib) {
+            mask_data[iq*n_blocks + ib] = 0.0f;
+            for (uint32_t ir = 0; ir < ratio; ++ir) {
+                const uint32_t cell = visible[ib*ratio + ir].second;
+                cell_data[(iq*n_blocks + ib)*ratio + ir] = cell;
+                used_cells[cell] = 1;
+            }
+            for (int64_t ip = 0; ip < n_pos; ++ip) {
+                pos_data[(iq*n_pos + ip)*n_blocks + ib] = visible[ib*ratio].first->pos[ip];
+            }
+        }
+
+        uint32_t fallback_cell = 0;
+        for (size_t ib = n_write; ib < (size_t) n_blocks; ++ib) {
+            for (uint32_t ir = 0; ir < ratio; ++ir) {
+                while (fallback_cell < used_cells.size() && used_cells[fallback_cell]) {
+                    ++fallback_cell;
+                }
+                GGML_ASSERT(fallback_cell < used_cells.size());
+                cell_data[(iq*n_blocks + ib)*ratio + ir] = fallback_cell;
+                used_cells[fallback_cell++] = 1;
+            }
+        }
+
+        const size_t selected_start = n_complete <= block_topk ? 0 : n_complete*ratio;
+        for (size_t iv = selected_start; iv < visible.size(); ++iv) {
+            selected_data[iq*n_kv + visible[iv].second] = 1.0f;
+        }
+    }
+}
+
 //
 // llama_memory_hybrid_idx_context
 //
@@ -295,7 +598,9 @@ llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context(
                   llama_context * lctx,
                            bool   optimize) :
     llama_memory_hybrid_context(mem, lctx, optimize),
-    mem(mem) {}
+    mem(mem),
+    ctx_idx(mem->get_mem_idx() == nullptr ? nullptr :
+        mem->get_mem_idx()->init_update(lctx, optimize)) {}
 
 llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context(
         llama_memory_hybrid_idx * mem,
@@ -307,7 +612,8 @@ llama_memory_hybrid_idx_context::llama_memory_hybrid_idx_context(
     mem(mem),
     ns_ubatch(llama_memory_hybrid_idx_ns(sinfos_idx)),
     ctx_idx(mem->get_mem_idx() == nullptr ? nullptr :
-        new llama_kv_cache_context(mem->get_mem_idx(), std::move(sinfos_idx), ubatches)) {}
+        new llama_kv_cache_context(mem->get_mem_idx(), std::move(sinfos_idx), ubatches)),
+    has_ubatches(true) {}
 
 bool llama_memory_hybrid_idx_context::next() {
     if (ctx_idx) {
@@ -326,6 +632,10 @@ bool llama_memory_hybrid_idx_context::apply() {
         res = res & ctx_idx->apply();
     }
 
+    if (res && ctx_idx && has_ubatches) {
+        mem->commit_qsa_tokens(ctx_idx->get_ubatch());
+    }
+
     return res;
 }
 
@@ -340,126 +650,17 @@ uint32_t llama_memory_hybrid_idx_context::get_n_stream() const {
 }
 
 void llama_memory_hybrid_idx_context::set_input_qsa(
-        ggml_tensor * cell_blk,
-        ggml_tensor * blk_cells,
-        ggml_tensor * blk_pos,
-        ggml_tensor * bias,
+        ggml_tensor *       block_cells,
+        ggml_tensor *       block_pos,
+        ggml_tensor *       block_mask,
+        ggml_tensor *       selected,
+        const ggml_tensor * kq_mask,
         const llama_ubatch * ubatch,
-        uint32_t ratio,
-        bool blk_bias) const {
+        uint32_t            ratio,
+        uint32_t            block_topk) const {
     GGML_ASSERT(ratio > 0);
     GGML_ASSERT(mem != nullptr && mem->get_mem_idx() != nullptr);
 
-    GGML_ASSERT(ggml_backend_buffer_is_host(cell_blk->buffer));
-
-    const int64_t n_kv     = cell_blk->ne[0];
-    const int64_t n_ns     = cell_blk->ne[1];        // streams in this ubatch
-    const int64_t n_blocks = blk_pos->ne[0]/(4*n_ns);
-    const int64_t n_tokens = ubatch->n_tokens;
-    const int64_t r        = ratio;
-
-    GGML_ASSERT(n_tokens % n_ns == 0);
-    const int64_t n_tps = n_tokens/n_ns;             // tokens per stream
-
-    int32_t * dst_cell_blk  = (int32_t *) cell_blk->data;
-    int32_t * dst_blk_cells = (int32_t *) blk_cells->data;
-    int32_t * dst_blk_pos   = (int32_t *) blk_pos->data;
-    float   * dst_bias      = (float   *) bias->data;
-
-    // block b covers [b*ratio, (b+1)*ratio), so its first token is at b*ratio
-    // all mrope sections carry it: exact for text, approximate for images
-    for (int64_t sec = 0; sec < 4; ++sec) {
-        for (int64_t s = 0; s < n_ns; ++s) {
-            for (int64_t b = 0; b < n_blocks; ++b) {
-                dst_blk_pos[sec*(n_blocks*n_ns) + s*n_blocks + b] = (int32_t) (b*r);
-            }
-        }
-    }
-
-    // one pass per stream: cell j is a different token in each, so no mapping is shared
-    std::vector<int32_t> blk_of(n_kv);
-    std::vector<int32_t> filled(n_blocks);
-
-    for (int64_t s = 0; s < n_ns; ++s) {
-        // ubatch index s*n_tps belongs to this stream; ask which cells array it uses
-        const llama_seq_id seq_of_stream = ubatch->seq_id[s*n_tps][0];
-        const auto & cells = mem->get_mem_idx()->get_cells(seq_of_stream);
-
-        int32_t * cur_cell_blk  = dst_cell_blk  + s*n_kv;
-        int32_t * cur_blk_cells = dst_blk_cells + s*(r*n_blocks);
-
-        // an incomplete block cannot be pooled; the bias below forces those tail cells in
-        // -1 means no usable block, and block 0 only keeps the gather in range
-        std::fill(blk_of.begin(),  blk_of.end(),  -1);
-        std::fill(filled.begin(),  filled.end(),   0);
-        std::fill(cur_blk_cells, cur_blk_cells + r*n_blocks, 0);
-
-        // a cell no block covers needs its own -inf, which a per-block bias cannot carry
-        // every cache path keeps the position below the cell window, so this stays false
-        bool oor = false;
-
-        for (int64_t j = 0; j < n_kv; ++j) {
-            if (cells.is_empty(j)) {
-                continue;
-            }
-
-            const llama_pos p = cells.pos_get(j);
-            const int64_t   b = p/r;
-
-            if (b >= n_blocks) {
-                oor = true;
-                continue;
-            }
-
-            blk_of[j] = (int32_t) b;
-            cur_blk_cells[b*r + (p%r)] = (int32_t) j;
-            filled[b]++;
-        }
-
-        GGML_ASSERT((!blk_bias || !oor) && "qsa: cell position runs past the cell window");
-
-        // per-block mode keeps an unpooled cell's real block, so the block's own -inf reaches it
-        // per-cell mode carries that -inf itself and only needs the gather in range
-        for (int64_t j = 0; j < n_kv; ++j) {
-            if (blk_of[j] >= 0 && filled[blk_of[j]] < r && !blk_bias) {
-                blk_of[j] = -1;
-            }
-            cur_cell_blk[j] = blk_of[j] < 0 ? 0 : blk_of[j];
-        }
-
-        for (int64_t ii = 0; ii < n_tps; ++ii) {
-            const int64_t      i      = s*n_tps + ii;
-            const llama_seq_id seq_id = ubatch->seq_id[i][0];
-            const llama_pos    q      = ubatch->pos[i];
-
-            // the tail is an incomplete block and is always visible, as in the reference
-            const llama_pos tail_start = (q + 1)/r*r;
-
-            if (blk_bias) {
-                // a block sits wholly inside or outside the tail, so one value covers it
-                // the caller adds the attention mask, which drops empty, foreign and future cells
-                float * cur_blk_bias = dst_bias + i*n_blocks;
-
-                for (int64_t b = 0; b < n_blocks; ++b) {
-                    // finite, so it can never meet a -inf and produce a nan
-                    cur_blk_bias[b] = b*r >= tail_start ? 1e9f : (filled[b] < r ? -INFINITY : 0.0f);
-                }
-
-                continue;
-            }
-
-            float * cur_bias = dst_bias + i*n_kv;
-
-            for (int64_t j = 0; j < n_kv; ++j) {
-                float v = -INFINITY;
-
-                if (!cells.is_empty(j) && cells.seq_has(j, seq_id) && cells.pos_get(j) <= q) {
-                    // finite, so it can never meet a -inf and produce a nan
-                    v = cells.pos_get(j) >= tail_start ? 1e9f : (blk_of[j] < 0 ? -INFINITY : 0.0f);
-                }
-
-                cur_bias[j] = v;
-            }
-        }
-    }
+    mem->set_input_qsa(block_cells, block_pos, block_mask, selected,
+            kq_mask, ubatch, ratio, block_topk);
 }

src/llama-memory-hybrid-idx.h

@@ -2,6 +2,8 @@
 
 #include "llama-memory-hybrid.h"
 
+#include <array>
+#include <map>
 #include <memory>
 #include <vector>
 
@@ -75,7 +77,20 @@ public:
 
     llama_kv_cache * get_mem_idx() const;   // nullptr when the model carries no indexer
 
+    void set_input_qsa(ggml_tensor * block_cells, ggml_tensor * block_pos,
+                       ggml_tensor * block_mask, ggml_tensor * selected,
+                       const ggml_tensor * kq_mask, const llama_ubatch * ubatch,
+                       uint32_t ratio, uint32_t block_topk) const;
+
+    void commit_qsa_tokens(const llama_ubatch & ubatch);
+
 private:
+    struct qsa_token {
+        std::array<llama_pos, 4> pos;
+    };
+
+    using qsa_history = std::vector<qsa_token>;
+
     // forget seq_id (all of it if seq_id < 0) in every cache at once, so a failed restore cannot leave the caches out of step
     // seq_id < 0 drops the whole context, as the caches themselves do on a failed restore
     void state_drop(llama_seq_id seq_id);
@@ -85,6 +100,8 @@ private:
     llama_hparams hparams_idx;
 
     const std::unique_ptr<llama_kv_cache> mem_idx;
+
+    std::map<llama_seq_id, qsa_history> qsa_histories;
 };
 
 class llama_memory_hybrid_idx_context : public llama_memory_hybrid_context {
@@ -123,26 +140,20 @@ public:
     // llama_memory_hybrid_idx_context specific API
     //
 
-    // nullptr with no indexer, and for the update context, which builds no sparse graph
+    // nullptr with no indexer
     const llama_kv_cache_context * get_idx() const;
 
     // streams in the current slot info, the `ns` of get_k/get_v; 1 if unified
     uint32_t get_n_stream() const;
 
-    // block-compressed sparse attention (qwen4exp QSA) over the cells of the indexer cache.
-    // Blocks cut the position line, not the cell array, so no caller assumes a contiguous layout:
-    //   cell_blk  I32 [n_kv, ns]           block each cell belongs to
-    //   blk_cells I32 [ratio*n_blocks, ns] cells making up each block
-    //   blk_pos   I32 [4*n_blocks*ns]      mrope position rows of each block's first token
-    //   bias      F32 [n_kv, n_tokens/ns, ns] -inf where invisible, large where always visible
-    // blk_bias asks for the bias per block instead: [n_blocks, n_tokens/ns, ns]
-    // the caller then adds the attention mask, the only part of the bias that varies within a block
-    void set_input_qsa(ggml_tensor * cell_blk, ggml_tensor * blk_cells, ggml_tensor * blk_pos,
-                       ggml_tensor * bias, const llama_ubatch * ubatch, uint32_t ratio,
-                       bool blk_bias) const;
+    // QSA blocks follow each sequence's token order, not physical cells or scalar positions.
+    void set_input_qsa(ggml_tensor * block_cells, ggml_tensor * block_pos,
+                       ggml_tensor * block_mask, ggml_tensor * selected,
+                       const ggml_tensor * kq_mask, const llama_ubatch * ubatch,
+                       uint32_t ratio, uint32_t block_topk) const;
 
 private:
-    const llama_memory_hybrid_idx * mem = nullptr;
+    llama_memory_hybrid_idx * mem = nullptr;
 
     // streams per ubatch, read from the slot infos before ctx_idx takes them
     // declared first, so it is initialised while sinfos_idx is still intact
@@ -151,6 +162,8 @@ private:
     // null unless the model has an indexer and this is a batch or full context
     const llama_memory_context_ptr ctx_idx;
 
+    const bool has_ubatches = false;
+
     // mirrors the base class's ubatch cursor, which is private there
     size_t i_cur = 0;
 };

src/llama-model-loader.cpp

@@ -1287,10 +1287,10 @@ struct ggml_tensor * llama_model_loader::create_tensor(
         return NULL;
     }
 
-    if ((flags & TENSOR_READ_LAZY) && use_mmap && lazy_mode != LLAMA_LAZY_MODE_OFF) {
+    if ((flags & TENSOR_READ_LAZY) && use_mmap && tensor_read_lazy != LLAMA_TENSOR_READ_LAZY_OFF) {
         // in auto mode, small tensors are cheap enough to keep resident
         constexpr size_t auto_lazy_min_size = 4ull * 1024 * 1024 * 1024;
-        if (lazy_mode == LLAMA_LAZY_MODE_ON || ggml_nbytes(cur) > auto_lazy_min_size) {
+        if (tensor_read_lazy == LLAMA_TENSOR_READ_LAZY_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));
 

src/llama-model-loader.h

@@ -84,7 +84,7 @@ struct llama_model_loader {
     bool load_mtp;
 
     // set by the caller before the create_tensor() calls
-    enum llama_lazy_mode lazy_mode = LLAMA_LAZY_MODE_OFF;
+    enum llama_tensor_read_lazy tensor_read_lazy = LLAMA_TENSOR_READ_LAZY_OFF;
 
     llama_files files;
     llama_ftype ftype;

src/llama-model.cpp

@@ -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,
-        /*.lazy_mode                   =*/ LLAMA_LAZY_MODE_AUTO,
+        /*.tensor_read_lazy            =*/ LLAMA_TENSOR_READ_LAZY_AUTO,
         /*.main_gpu                    =*/ 0,
         /*.tensor_split                =*/ nullptr,
         /*.progress_callback           =*/ nullptr,

src/llama.cpp

@@ -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.lazy_mode = params.lazy_mode;
+        ml.tensor_read_lazy = params.tensor_read_lazy;
 
         ml.print_info();
         std::unique_ptr<llama_model> model_ptr(llama_model_create(ml, params));

src/models/models.h

@@ -2310,22 +2310,12 @@ struct llama_model_qwen4exp : public llama_model_base {
                             int * sections,
                             int   il);
 
-        // dense self-attention restricted to the cells that top_k names
-        ggml_tensor * build_attn_qsa(
-        llm_graph_input_attn_kv * inp,
-                    ggml_tensor * q_cur,
-                    ggml_tensor * k_cur,
-                    ggml_tensor * v_cur,
-                    ggml_tensor * top_k,
-                          float   kq_scale,
-                            int   il);
-
         // the QSA cache layout inputs do not depend on the layer, only on its compress ratio,
         // so the layers sharing a ratio share one input set
         std::map<uint32_t, llm_graph_input_qsa *> qsa_inps;
 
-        // QSA: token indices this layer's queries may attend to, or nullptr for dense
-        ggml_tensor * build_qsa_top_k(
+        // QSA mask for this layer, or nullptr for dense attention
+        ggml_tensor * build_qsa_mask(
   const llama_memory_hybrid_idx_context * mctx_hyb,
                     ggml_tensor * cur,
                     ggml_tensor * inp_pos,
@@ -2360,12 +2350,9 @@ 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,
-                    ggml_tensor * emb,
+  const llama_memory_hybrid_idx_context * mctx_hyb,
                     ggml_tensor * hidden,
                             int   il);
 

src/models/qwen4exp.cpp

@@ -18,21 +18,29 @@ void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) {
     ml.get_key(LLM_KV_SSM_STATE_SIZE,     hparams.ssm_d_state);
     ml.get_key(LLM_KV_SSM_TIME_STEP_RANK, hparams.ssm_dt_rank);
     ml.get_key(LLM_KV_SSM_GROUP_COUNT,    hparams.ssm_n_group);
-    GGML_ASSERT(hparams.ssm_d_conv  > 0 && hparams.ssm_d_inner > 0 && hparams.ssm_d_state > 0 &&
-                hparams.ssm_dt_rank > 0 && hparams.ssm_n_group > 0);
+    if (hparams.ssm_d_conv == 0 || hparams.ssm_d_inner == 0 || hparams.ssm_d_state == 0 ||
+            hparams.ssm_dt_rank == 0 || hparams.ssm_n_group == 0 ||
+            hparams.ssm_dt_rank % hparams.ssm_n_group != 0 ||
+            (uint64_t) hparams.ssm_d_state * hparams.ssm_dt_rank != hparams.ssm_d_inner) {
+        throw std::runtime_error("invalid Qwen4-Exp gated delta net dimensions");
+    }
 
     // HC; low_rank is qwen4exp-specific, DeepSeek-V4 leaves it absent (full rank)
     ml.get_key(LLM_KV_HYPER_CONNECTION_COUNT,    hparams.dsv4_hc_mult);
     ml.get_key(LLM_KV_HYPER_CONNECTION_LOW_RANK, hparams.hc_low_rank);
-    GGML_ASSERT(hparams.dsv4_hc_mult > 0 && hparams.hc_low_rank > 0);
+    if (hparams.n_embd == 0 || hparams.dsv4_hc_mult <= 1 || hparams.hc_low_rank == 0 ||
+            hparams.dsv4_hc_mult > UINT32_MAX/hparams.n_embd) {
+        throw std::runtime_error("invalid Qwen4-Exp hyper-connection dimensions");
+    }
     hparams.n_embd_out_impl = hparams.dsv4_hc_mult * hparams.n_embd;
 
     ml.get_key(LLM_KV_ATTENTION_INDEXER_HEAD_COUNT, hparams.indexer_n_head);
     ml.get_key(LLM_KV_ATTENTION_INDEXER_KEY_LENGTH, hparams.indexer_head_size);
     ml.get_key(LLM_KV_ATTENTION_INDEXER_TOP_K,      hparams.indexer_top_k);
-    GGML_ASSERT(hparams.indexer_n_head > 0
-             && hparams.indexer_head_size > 0
-             && hparams.indexer_top_k > 0);
+    if (hparams.indexer_n_head == 0 || hparams.indexer_head_size == 0 || hparams.indexer_top_k == 0 ||
+            hparams.n_rot_full > hparams.indexer_head_size) {
+        throw std::runtime_error("invalid Qwen4-Exp sparse-attention dimensions");
+    }
     ml.get_key_or_arr(LLM_KV_ATTENTION_COMPRESS_RATIOS, hparams.dsv4_compress_ratios, hparams.n_layer_all, false);
 
     // PLE n-gram hash embeddings; if the key group is absent every field stays zero
@@ -44,9 +52,11 @@ void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) {
     if (n_ple > 0) {
         std::vector<uint32_t> ple_layers;
         ml.get_arr(LLM_KV_PLE_LAYERS, ple_layers);
-        GGML_ASSERT(n_ple == 1 && "qwen4exp supports only one PLE layer");
+        if (n_ple != 1 || ple_layers.size() != n_ple) {
+            throw std::runtime_error("Qwen4-Exp supports exactly one PLE layer");
+        }
         for (uint32_t il : ple_layers) {
-            if (il >= hparams.n_layer_all) {
+            if (il >= hparams.n_layer()) {
                 throw std::runtime_error(format("PLE layer %u is out of range", il));
             }
             hparams.is_ple_impl.set(il);
@@ -59,15 +69,30 @@ void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) {
         // optional: files written before this key fall back to the EOS token
         ml.get_key(LLM_KV_PLE_IMAGE_TOKEN_ID,  hparams.ple_image_token_id, false);
         ml.get_key(LLM_KV_EMBEDDING_LENGTH_PER_LAYER, hparams.n_embd_per_layer);
-        GGML_ASSERT(hparams.ple_conv_kernel > 0 && hparams.n_embd_per_layer > 0);
+        if (hparams.ple_conv_kernel == 0 || hparams.n_embd_per_layer == 0) {
+            throw std::runtime_error("invalid Qwen4-Exp PLE dimensions");
+        }
 
-        hparams.ple_n_heads  = (hparams.ple_ngram_size - 1) * hparams.ple_heads_per_ngram;
-        hparams.ple_head_dim = hparams.n_embd_per_layer;
         if (hparams.ple_ngram_size < 2 || hparams.ple_ngram_size > LLAMA_MAX_PLE_NGRAM) {
             throw std::runtime_error(format("PLE n-gram size %u is out of range", hparams.ple_ngram_size));
         }
-        if (hparams.ple_n_heads == 0 || hparams.ple_n_heads > LLAMA_MAX_PLE_HEADS) {
-            throw std::runtime_error(format("PLE head count %u is out of range", hparams.ple_n_heads));
+
+        const uint64_t ple_n_heads = (uint64_t) (hparams.ple_ngram_size - 1) * hparams.ple_heads_per_ngram;
+        hparams.ple_head_dim = hparams.n_embd_per_layer;
+        if (ple_n_heads == 0 || ple_n_heads > LLAMA_MAX_PLE_HEADS) {
+            throw std::runtime_error(format("PLE head count %" PRIu64 " is out of range", ple_n_heads));
+        }
+        hparams.ple_n_heads = (uint32_t) ple_n_heads;
+
+        uint32_t n_multipliers = 0;
+        uint32_t n_offsets = 0;
+        uint32_t n_vocab_sizes = 0;
+        ml.get_arr_n(LLM_KV_PLE_LAYER_MULTIPLIERS, n_multipliers);
+        ml.get_arr_n(LLM_KV_PLE_HEAD_OFFSETS,      n_offsets);
+        ml.get_arr_n(LLM_KV_PLE_HEAD_VOCAB_SIZES,  n_vocab_sizes);
+        if (n_multipliers != hparams.ple_ngram_size ||
+                n_offsets != hparams.ple_n_heads || n_vocab_sizes != hparams.ple_n_heads) {
+            throw std::runtime_error("invalid Qwen4-Exp PLE metadata lengths");
         }
 
         ml.get_arr(LLM_KV_PLE_LAYER_MULTIPLIERS, hparams.ple_layer_multipliers);
@@ -93,12 +118,28 @@ void llama_model_qwen4exp::load_arch_hparams(llama_model_loader & ml) {
     if (!ml.get_key_or_arr(LLM_KV_ATTENTION_RECURRENT_LAYERS, hparams.is_recr_impl, hparams.n_layer_all, false)) {
         uint32_t full_attn_interval = 4;
         ml.get_key(LLM_KV_FULL_ATTENTION_INTERVAL, full_attn_interval, false);
-        GGML_ASSERT(full_attn_interval > 0);
+        if (full_attn_interval == 0) {
+            throw std::runtime_error("invalid Qwen4-Exp full-attention interval");
+        }
         for (uint32_t i = 0; i < hparams.n_layer_all; ++i) {
             hparams.is_recr_impl[i] = (i < hparams.n_layer()) && ((i + 1) % full_attn_interval != 0);
         }
     }
 
+    for (uint32_t il = 0; il < hparams.n_layer(); ++il) {
+        const uint32_t ratio = hparams.dsv4_compress_ratios[il];
+        if (hparams.is_recr(il)) {
+            if (ratio != 0) {
+                throw std::runtime_error(format("Qwen4-Exp recurrent layer %u has a QSA compression ratio", il));
+            }
+        } else if (ratio == 0 || hparams.indexer_top_k % ratio != 0) {
+            throw std::runtime_error(format("invalid Qwen4-Exp QSA compression ratio %u at layer %u", ratio, il));
+        }
+        if (hparams.is_ple(il) && !hparams.is_recr(il)) {
+            throw std::runtime_error(format("Qwen4-Exp PLE layer %u is not recurrent", il));
+        }
+    }
+
     switch (hparams.n_layer()) {
         case 48: type = LLM_TYPE_A3B; break;
         default: type = LLM_TYPE_UNKNOWN;
@@ -127,8 +168,15 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
     // flat [ple_head_dim, n_rows] gather target; n_rows is padded, so read it back
     if (hparams.ple_n_heads > 0) {
         const std::string ple_name = tn(LLM_TENSOR_PER_LAYER_TOKEN_EMBD, "weight").str();
-        const auto & ple_w = ml.require_weight(ple_name.c_str());
-        const int64_t ple_rows = ple_w.tensor->ne[1];
+        const auto * ple_w = ml.get_weight(ple_name.c_str());
+        int64_t ple_rows = 0;
+        for (uint32_t h = 0; h < hparams.ple_n_heads; ++h) {
+            ple_rows = std::max<int64_t>(ple_rows,
+                    (int64_t) hparams.ple_head_offsets[h] + hparams.ple_head_vocab_sizes[h]);
+        }
+        if (ple_w) {
+            ple_rows = ple_w->tensor->ne[1];
+        }
 
         // sanity check
         for (uint32_t h = 0; h < hparams.ple_n_heads; ++h) {
@@ -190,8 +238,9 @@ void llama_model_qwen4exp::load_arch_tensors(llama_model_loader & ml) {
         }
 
         if (hparams.is_ple(il)) {
-            layer.ple_key        = create_tensor(tn(LLM_TENSOR_PLE_KEY,        "weight", il), { n_embd, hc_dim }, 0);
-            layer.ple_value      = create_tensor(tn(LLM_TENSOR_PLE_VALUE,      "weight", il), { n_embd, n_embd }, 0);
+            const int64_t ple_dim = (int64_t) hparams.ple_head_dim * hparams.ple_n_heads;
+            layer.ple_key        = create_tensor(tn(LLM_TENSOR_PLE_KEY,        "weight", il), { ple_dim, hc_dim }, 0);
+            layer.ple_value      = create_tensor(tn(LLM_TENSOR_PLE_VALUE,      "weight", il), { ple_dim, n_embd }, 0);
             layer.ple_norm_key   = create_tensor(tn(LLM_TENSOR_PLE_NORM_KEY,   "weight", il), { hc_dim }, 0);
             layer.ple_norm_query = create_tensor(tn(LLM_TENSOR_PLE_NORM_QUERY, "weight", il), { hc_dim }, 0);
             layer.ple_norm_conv  = create_tensor(tn(LLM_TENSOR_PLE_NORM_CONV,  "weight", il), { hc_dim }, 0);
@@ -296,7 +345,6 @@ 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();
 
@@ -313,13 +361,6 @@ 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),
@@ -330,7 +371,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(), ple_emb, res_hc, il);
+            res_hc = build_ple(inp->get_recr(), mctx_hyb, res_hc, il);
         }
 
         ggml_tensor * inject = nullptr;
@@ -424,13 +465,17 @@ ggml_tensor * llama_model_qwen4exp::graph::build_norm_gated(
 // one mean-pooled indexer key scores each block; set_input resolves the cache layout
 class llama_model_qwen4exp::llm_graph_input_qsa : public llm_graph_input_i {
 public:
-    llm_graph_input_qsa(const llama_memory_hybrid_idx_context * mctx, uint32_t ratio, bool blk_bias) :
-        mctx(mctx), ratio(ratio), blk_bias(blk_bias) {}
+    llm_graph_input_qsa(
+            const llama_memory_hybrid_idx_context * mctx,
+            ggml_tensor * kq_mask,
+            uint32_t ratio,
+            uint32_t block_topk) :
+        mctx(mctx), kq_mask(kq_mask), ratio(ratio), block_topk(block_topk) {}
     virtual ~llm_graph_input_qsa() = default;
 
     void set_input(const llama_ubatch * ubatch) override {
-        mctx->get_idx()->set_input_k_idxs(k_idxs, ubatch);
-        mctx->set_input_qsa(cell_blk, blk_cells, blk_pos, bias, ubatch, ratio, blk_bias);
+        mctx->set_input_qsa(block_cells, block_pos, block_mask, selected,
+                kq_mask, ubatch, ratio, block_topk);
     }
 
     bool can_reuse(const llm_graph_params & params) override {
@@ -442,39 +487,33 @@ public:
         }
 
         const int64_t n_kv     = idx->get_n_kv();
-        const int64_t n_stream = mctx->get_n_stream();
-        const int64_t n_blocks = (n_kv + ratio - 1)/ratio;
+        const int64_t n_blocks = n_kv/ratio;
 
         bool res = true;
 
-        res &= params.ubatch.n_tokens % n_stream == 0;
-
-        res &= k_idxs->ne[0]    == params.ubatch.n_tokens;
-        res &= cell_blk->ne[0]  == n_kv;
-        res &= cell_blk->ne[1]  == n_stream;
-        res &= blk_cells->ne[0] == (int64_t) ratio*n_blocks;
-        res &= blk_pos->ne[0]   == 4*n_blocks*n_stream;
-        res &= bias->ne[0]      == (blk_bias ? n_blocks : n_kv);
-        res &= bias->ne[1]      == params.ubatch.n_tokens/n_stream;
+        res &= n_kv > (int64_t) block_topk*ratio + ratio - 1;
+        res &= block_cells->ne[1] == n_blocks;
+        res &= block_cells->ne[2] == params.ubatch.n_tokens;
+        res &= block_pos->ne[2]   == params.ubatch.n_tokens;
+        res &= block_mask->ne[1]  == params.ubatch.n_tokens;
+        res &= selected->ne[0]    == n_kv;
+        res &= selected->ne[1]    == params.ubatch.n_tokens;
 
         return res;
     }
 
-    // per stream: a cell index names a different token in each stream
-    ggml_tensor * k_idxs    = nullptr;   // I32 [n_tokens]
-    ggml_tensor * cell_blk  = nullptr;   // I32 [n_kv, n_stream]
-    ggml_tensor * blk_cells = nullptr;   // I32 [ratio*n_blocks, n_stream]
-    ggml_tensor * blk_pos   = nullptr;   // I32 [4*n_blocks*n_stream]
-    ggml_tensor * bias      = nullptr;   // F32 [n_blocks or n_kv, n_tokens/n_stream, n_stream]
+    ggml_tensor * block_cells = nullptr;
+    ggml_tensor * block_pos   = nullptr;
+    ggml_tensor * block_mask  = nullptr;
+    ggml_tensor * selected    = nullptr;
 
     const llama_memory_hybrid_idx_context * mctx;
+    ggml_tensor * kq_mask;
     const uint32_t ratio;
-
-    // the per-cell half of the bias is the attention mask, so only the per-block half is uploaded
-    const bool blk_bias;
+    const uint32_t block_topk;
 };
 
-ggml_tensor * llama_model_qwen4exp::graph::build_qsa_top_k(
+ggml_tensor * llama_model_qwen4exp::graph::build_qsa_mask(
         const llama_memory_hybrid_idx_context * mctx_hyb,
         ggml_tensor *                           cur,
         ggml_tensor *                           inp_pos,
@@ -490,21 +529,13 @@ ggml_tensor * llama_model_qwen4exp::graph::build_qsa_top_k(
 
     GGML_ASSERT(r > 0);
 
-    const int64_t n_blocks = (n_kv + r - 1)/r;
+    const int64_t n_blocks = n_kv/r;
+    const int64_t block_topk = hparams.indexer_top_k/r;
 
-    // build_attn_qsa and the KQ mask need the tokens to divide evenly across the streams
     const int64_t n_stream = mctx_hyb->get_n_stream();
     GGML_ASSERT(n_tokens % n_stream == 0);
     const int64_t n_tps = n_tokens/n_stream;
 
-    // only the "which block is visible" half of the bias varies per block
-    // the rest is the visible/not test the attention mask already carries, so upload the per-block half only: 1/ratio of the cells
-    // alibi writes distances instead of a mask and non-causal keeps future cells, so both opt out
-    // the mask also holds an mrope rule for the query's own position, but only 2d image positions can differ there
-    const bool blk_bias = kq_mask != nullptr &&
-        kq_mask->ne[0] == n_kv && kq_mask->ne[1] == n_tps && kq_mask->ne[3] == n_stream &&
-        cparams.causal_attn && !hparams.use_alibi;
-
     // nothing above depends on the layer, so the layers sharing a ratio share one input set
     llm_graph_input_qsa * inp = nullptr;
 
@@ -512,59 +543,29 @@ ggml_tensor * llama_model_qwen4exp::graph::build_qsa_top_k(
     if (it != qsa_inps.end()) {
         inp = it->second;
     } else {
-        auto qsa = std::make_unique<llm_graph_input_qsa>(mctx_hyb, (uint32_t) r, blk_bias);
+        auto qsa = std::make_unique<llm_graph_input_qsa>(
+                mctx_hyb, kq_mask, (uint32_t) r, (uint32_t) block_topk);
 
-        qsa->k_idxs    = mctx_idx->build_input_k_idxs(ctx0, ubatch);
-        qsa->cell_blk  = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, n_kv, n_stream);
-        qsa->blk_cells = ggml_new_tensor_2d(ctx0, GGML_TYPE_I32, r*n_blocks, n_stream);
-        qsa->blk_pos   = ggml_new_tensor_1d(ctx0, GGML_TYPE_I32, 4*n_blocks*n_stream);
-        qsa->bias      = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, blk_bias ? n_blocks : n_kv, n_tps, n_stream);
+        qsa->block_cells = ggml_new_tensor_3d(ctx0, GGML_TYPE_I32, r, n_blocks, n_tokens);
+        qsa->block_pos   = ggml_new_tensor_3d(ctx0, GGML_TYPE_I32, n_blocks, 4, n_tokens);
+        qsa->block_mask  = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_blocks, n_tokens);
+        qsa->selected    = ggml_new_tensor_2d(ctx0, GGML_TYPE_F32, n_kv, n_tokens);
 
-        ggml_set_input(qsa->cell_blk);
-        ggml_set_input(qsa->blk_cells);
-        ggml_set_input(qsa->blk_pos);
-        ggml_set_input(qsa->bias);
+        ggml_set_input(qsa->block_cells);
+        ggml_set_input(qsa->block_pos);
+        ggml_set_input(qsa->block_mask);
+        ggml_set_input(qsa->selected);
 
         inp = qsa.get();
         res->add_input(std::move(qsa));
         qsa_inps.emplace((uint32_t) r, inp);
     }
 
-    // cached indexer keys are raw: pooling precedes norm and rotation, so apply neither
-    ggml_tensor * k_raw = build_lora_mm(model.layers[il].index_k_proj, cur);
-    k_raw = ggml_reshape_3d(ctx0, k_raw, idx_dim, 1, n_tokens);
-    cb(k_raw, "indexer_k_raw", il);
-
-    ggml_build_forward_expand(gf, mctx_idx->cpy_k(ctx0, k_raw, inp->k_idxs, il));
+    kq_mask = inp->kq_mask;
 
     // one key head, so rows are contiguous. get_k gives [idx_dim, n_head_kv, n_kv, n_stream].
     ggml_tensor * k_all = mctx_idx->get_k(ctx0, il);
     k_all = ggml_view_3d(ctx0, k_all, idx_dim, n_kv, n_stream, k_all->nb[2], k_all->nb[3], 0);
-
-    // gathers per stream: blk_cells row s indexes stream s's own cells
-    ggml_tensor * members = ggml_get_rows(ctx0, k_all, inp->blk_cells);
-    members = ggml_reshape_4d(ctx0, members, idx_dim, r, n_blocks, n_stream);
-
-    // mean over the block members; r is small, so summing slices beats a transpose plus sum_rows
-    ggml_tensor * pooled = nullptr;
-    for (int64_t i = 0; i < r; ++i) {
-        ggml_tensor * slice = ggml_cont(ctx0,
-                ggml_view_3d(ctx0, members, idx_dim, n_blocks, n_stream,
-                        members->nb[2], members->nb[3], i*members->nb[1]));
-        pooled = pooled ? ggml_add(ctx0, pooled, slice) : slice;
-    }
-    pooled = ggml_scale(ctx0, pooled, 1.0f/(float) r);
-    cb(pooled, "indexer_k_pooled", il);
-
-    // rope wants [n_dims, n_head, n_tokens]: lay every stream's blocks flat, split after.
-    pooled = ggml_reshape_3d(ctx0, pooled, idx_dim, 1, n_blocks*n_stream);
-    pooled = build_norm(pooled, model.layers[il].index_k_norm, nullptr, LLM_NORM_RMS, il);
-    pooled = ggml_rope_multi(ctx0, pooled, inp->blk_pos, nullptr,
-            n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale,
-            ext_factor, attn_factor, beta_fast, beta_slow);
-    pooled = ggml_reshape_3d(ctx0, pooled, idx_dim, n_blocks, n_stream);
-    cb(pooled, "indexer_k", il);
-
     ggml_tensor * q = build_lora_mm(model.layers[il].index_q_proj, cur);
     q = ggml_reshape_3d(ctx0, q, idx_dim, n_idx_h, n_tokens);
     q = build_norm(q, model.layers[il].index_q_norm, nullptr, LLM_NORM_RMS, il);
@@ -573,128 +574,86 @@ ggml_tensor * llama_model_qwen4exp::graph::build_qsa_top_k(
             ext_factor, attn_factor, beta_fast, beta_slow);
     cb(q, "indexer_q", il);
 
-    // rectify each head dot product before the sum, as in the DeepSeek lightning indexer
-    // mul_mat matches ne[2], so the queries of stream s only meet the blocks of stream s
-    ggml_tensor * score = ggml_mul_mat(ctx0, pooled,
-            ggml_reshape_3d(ctx0, ggml_cont(ctx0, q), idx_dim, n_idx_h*n_tps, n_stream));
-    score = ggml_reshape_4d(ctx0, score, n_blocks, n_idx_h, n_tps, n_stream);
-    score = ggml_relu(ctx0, score);
-    score = ggml_cont(ctx0, ggml_permute(ctx0, score, 1, 0, 2, 3));
-    score = ggml_sum_rows(ctx0, score);
-    score = ggml_reshape_3d(ctx0, score, n_blocks, n_tps, n_stream);
-    cb(score, "indexer_score", il);
+    const ggml_type activation_type = mctx_idx->type_k();
+    std::vector<ggml_tensor *> selected_streams;
+    selected_streams.reserve(n_stream);
 
-    // one value per block, so it is cheaper to bias here than after the cells are expanded
-    if (blk_bias) {
-        score = ggml_add(ctx0, score, inp->bias);
+    for (int64_t is = 0; is < n_stream; ++is) {
+        ggml_tensor * cache = ggml_view_2d(ctx0, k_all, idx_dim, n_kv,
+                k_all->nb[1], is*k_all->nb[2]);
+        ggml_tensor * block_cells = ggml_view_3d(ctx0, inp->block_cells, r, n_blocks, n_tps,
+                inp->block_cells->nb[1], inp->block_cells->nb[2], is*n_tps*inp->block_cells->nb[2]);
+        ggml_tensor * block_keys = ggml_get_rows(ctx0, cache,
+                ggml_reshape_1d(ctx0, block_cells, r*n_blocks*n_tps));
+        block_keys = ggml_reshape_4d(ctx0, block_keys, idx_dim, r, n_blocks, n_tps);
+        block_keys = ggml_cont(ctx0, ggml_transpose(ctx0, block_keys));
+        block_keys = ggml_mean(ctx0, block_keys);
+        block_keys = ggml_cont(ctx0, ggml_transpose(ctx0, block_keys));
+        block_keys = ggml_reshape_3d(ctx0, block_keys, idx_dim, 1, n_blocks*n_tps);
+
+        if (block_keys->type != activation_type) {
+            block_keys = ggml_cast(ctx0, block_keys, activation_type);
+            block_keys = ggml_cast(ctx0, block_keys, GGML_TYPE_F32);
+        }
+        block_keys = build_norm(block_keys, model.layers[il].index_k_norm, nullptr, LLM_NORM_RMS, il);
+
+        ggml_tensor * block_pos = ggml_view_3d(ctx0, inp->block_pos, n_blocks, 4, n_tps,
+                inp->block_pos->nb[1], inp->block_pos->nb[2], is*n_tps*inp->block_pos->nb[2]);
+        block_keys = ggml_rope_multi(ctx0, block_keys,
+                ggml_reshape_1d(ctx0, block_pos, n_blocks*4*n_tps), nullptr,
+                n_rot, sections, rope_type, n_ctx_orig, freq_base, freq_scale,
+                ext_factor, attn_factor, beta_fast, beta_slow);
+        cb(block_keys, "indexer_k", il);
+
+        block_keys = ggml_reshape_4d(ctx0, block_keys, idx_dim, n_blocks, 1, n_tps);
+        ggml_tensor * query = ggml_view_3d(ctx0, q, idx_dim, n_idx_h, n_tps,
+                q->nb[1], q->nb[2], is*n_tps*q->nb[2]);
+        query = ggml_reshape_4d(ctx0, query, idx_dim, n_idx_h, 1, n_tps);
+
+        ggml_tensor * scores = ggml_mul_mat(ctx0, block_keys, query);
+        ggml_mul_mat_set_prec(scores, GGML_PREC_F32);
+        scores = ggml_relu(ctx0, scores);
+        scores = ggml_sum_rows(ctx0, ggml_cont(ctx0, ggml_permute(ctx0, scores, 1, 0, 2, 3)));
+        scores = ggml_scale(ctx0, ggml_reshape_2d(ctx0, scores, n_blocks, n_tps),
+                1.0f/sqrtf((float) idx_dim));
+
+        ggml_tensor * block_mask = ggml_view_2d(ctx0, inp->block_mask, n_blocks, n_tps,
+                inp->block_mask->nb[1], is*n_tps*inp->block_mask->nb[1]);
+        scores = ggml_add(ctx0, scores, block_mask);
+        cb(scores, "indexer_score", il);
+
+        ggml_tensor * top_blocks = ggml_top_k(ctx0, scores, block_topk);
+        ggml_tensor * top_cells = ggml_get_rows(ctx0, block_cells, top_blocks);
+        top_cells = ggml_reshape_2d(ctx0, top_cells, r*block_topk, n_tps);
+
+        ggml_tensor * base_selected = ggml_view_2d(ctx0, inp->selected, n_kv, n_tps,
+                inp->selected->nb[1], is*n_tps*inp->selected->nb[1]);
+        base_selected = ggml_reshape_3d(ctx0, base_selected, 1, n_kv, n_tps);
+        ggml_tensor * selected_top = ggml_fill(ctx0, base_selected, 0.0f);
+        ggml_tensor * ones = ggml_new_tensor_3d(ctx0, GGML_TYPE_F32, 1, r*block_topk, n_tps);
+        ones = ggml_fill(ctx0, ones, 1.0f);
+        selected_top = ggml_set_rows(ctx0, selected_top, ones, top_cells);
+        ggml_tensor * selected_stream = ggml_clamp(
+                ctx0, ggml_add(ctx0, base_selected, selected_top), 0.0f, 1.0f);
+        selected_streams.push_back(ggml_reshape_2d(ctx0, selected_stream, n_kv, n_tps));
     }
 
-    // every token of a block gets the block score; the budget is whole blocks, so top-k cuts on a block boundary
-    ggml_tensor * expanded = ggml_get_rows(ctx0,
-            ggml_cont(ctx0, ggml_permute(ctx0, score, 1, 0, 2, 3)), inp->cell_blk);
-    expanded = ggml_cont(ctx0, ggml_permute(ctx0, expanded, 1, 0, 2, 3));
-
-    if (blk_bias) {
-        // flash attention keeps the mask in f16; the scores are f32
-        ggml_tensor * mask = kq_mask->type == GGML_TYPE_F32 ? kq_mask : ggml_cast(ctx0, kq_mask, GGML_TYPE_F32);
-        expanded = ggml_add(ctx0, expanded, ggml_reshape_3d(ctx0, mask, n_kv, n_tps, n_stream));
-    } else {
-        expanded = ggml_add(ctx0, expanded, inp->bias);
+    ggml_tensor * selected = selected_streams[0];
+    for (int64_t is = 1; is < n_stream; ++is) {
+        selected = ggml_concat(ctx0, selected, selected_streams[is], 1);
     }
-    cb(expanded, "indexer_score_tokens", il);
+    selected = ggml_scale_bias(ctx0, selected, 1e30f, -1e30f);
 
-    // the reference returns indexer_top_k + compress_ratio - 1: whole blocks plus the tail
-    const int64_t width = std::min<int64_t>(n_kv, (int64_t) hparams.indexer_top_k + r - 1);
-
-    ggml_tensor * top_k = ggml_cont(ctx0, ggml_top_k(ctx0, expanded, width));
-
-    // build_attn_qsa reads [n_top_k, n_batch, 1, n_stream], matching the KQ mask.
-    top_k = ggml_reshape_4d(ctx0, top_k, width, n_tps, 1, n_stream);
-    cb(top_k, "indexer_top_k", il);
-
-    return top_k;
-}
-
-// Dense GQA self-attention restricted to the cells that top_k names.
-// The mask build below copies the MLA sparse path in llm_graph_context::build_attn.
-ggml_tensor * llama_model_qwen4exp::graph::build_attn_qsa(
-        llm_graph_input_attn_kv * inp,
-        ggml_tensor *             q_cur,
-        ggml_tensor *             k_cur,
-        ggml_tensor *             v_cur,
-        ggml_tensor *             top_k,
-        float                     kq_scale,
-        int                       il) {
-    // rotate q/k/v before they reach a quantized cache, as the dense path does. the indexer
-    // has already scored with its own query in build_qsa_top_k, so top_k is unaffected.
-    if (inp->self_k_rot) {
-        q_cur = llama_mul_mat_hadamard(ctx0, q_cur, inp->self_k_rot);
-        k_cur = llama_mul_mat_hadamard(ctx0, k_cur, inp->self_k_rot);
+    ggml_tensor * base_mask = ggml_reshape_2d(ctx0, kq_mask, n_kv, n_tokens);
+    if (base_mask->type != GGML_TYPE_F32) {
+        base_mask = ggml_cast(ctx0, base_mask, GGML_TYPE_F32);
     }
-
-    if (inp->self_v_rot) {
-        v_cur = llama_mul_mat_hadamard(ctx0, v_cur, inp->self_v_rot);
+    ggml_tensor * mask = ggml_add(ctx0, base_mask, selected);
+    if (cparams.flash_attn) {
+        mask = ggml_cast(ctx0, mask, GGML_TYPE_F16);
     }
-
-    // these nodes are added to the graph together so that they are not reordered
-    // by doing so, the number of splits in the graph is reduced
-    // expand k later to enable rope fusion which directly writes into k-v cache
-    ggml_build_forward_expand(gf, q_cur);
-    ggml_build_forward_expand(gf, v_cur);
-    ggml_build_forward_expand(gf, k_cur);
-
-    const auto * mctx_cur = inp->mctx;
-
-    // store to KV cache
-    {
-        const auto & k_idxs = inp->get_k_idxs();
-        const auto & v_idxs = inp->get_v_idxs();
-
-        ggml_build_forward_expand(gf, mctx_cur->cpy_k(ctx0, k_cur, k_idxs, il));
-        ggml_build_forward_expand(gf, mctx_cur->cpy_v(ctx0, v_cur, v_idxs, il));
-    }
-
-    ggml_tensor * kq_mask = inp->get_kq_mask();
-
-    // prepare new kq mask - starts filled with -INFINITY
-    ggml_tensor * kq_mask_all = ggml_fill(ctx0, kq_mask, -INFINITY);
-
-    // reshape KQ mask into tensor with rows of size 1:
-    // [n_kv, n_batch, 1, n_stream] -> [1, n_kv, n_batch, n_stream]
-    kq_mask_all = ggml_view_4d(ctx0, kq_mask_all, 1, kq_mask_all->ne[0], kq_mask_all->ne[1], kq_mask_all->ne[3], kq_mask_all->nb[0], kq_mask_all->nb[1], kq_mask_all->nb[2], 0);
-
-    // reshape top_k indices: [n_top_k, n_batch, 1, n_stream] -> [n_top_k, n_batch, n_stream, 1]
-    ggml_tensor * top_k_3d = ggml_view_4d(ctx0, top_k, top_k->ne[0], top_k->ne[1], top_k->ne[3], 1, top_k->nb[1], top_k->nb[2], top_k->ne[3]*top_k->nb[3], 0);
-
-    // prepare zero-filled tensor with rows of size 1: [1, n_top_k, n_batch, n_stream]
-    // this will be our source of zero values for unmasking top k mask elements
-    ggml_tensor * zeros = ggml_new_tensor_4d(ctx0, GGML_TYPE_F32, 1, top_k_3d->ne[0], top_k_3d->ne[1], top_k_3d->ne[2]);
-    zeros = ggml_fill(ctx0, zeros, 0.0f);
-
-    // modify KQ mask by unmasking elements that are in top_k indices
-    // ggml_set_rows([1, n_kv, n_batch, n_stream], [1, n_top_k, n_batch, n_stream], [n_top_k, n_batch, n_stream, 1])
-    ggml_tensor * kq_mask_top_k = ggml_set_rows(ctx0, kq_mask_all, zeros, top_k_3d);
-
-    // reshape to restore the original shape of KQ mask:
-    // [1, n_kv, n_batch, n_stream] -> [n_kv, n_batch, 1, n_stream]
-    kq_mask_top_k = ggml_view_4d(ctx0, kq_mask_top_k, kq_mask_top_k->ne[1], kq_mask_top_k->ne[2], 1, kq_mask_top_k->ne[3], kq_mask_top_k->nb[2], kq_mask_top_k->nb[3], kq_mask_top_k->nb[3], 0);
-
-    // combine with the original kq mask
-    kq_mask_top_k = ggml_add(ctx0, kq_mask_top_k, kq_mask);
-
-    ggml_tensor * q = q_cur;
-    ggml_tensor * k = mctx_cur->get_k(ctx0, il);
-    ggml_tensor * v = mctx_cur->get_v(ctx0, il);
-
-    ggml_tensor * cur = build_attn_mha(q, k, v, nullptr, kq_mask_top_k, nullptr, nullptr, kq_scale, il);
-    cb(cur, "kqv_out", il);
-
-    // the rotation is its own inverse, so undo it on the value side of the output
-    if (inp->self_v_rot) {
-        cur = llama_mul_mat_hadamard(ctx0, cur, inp->self_v_rot);
-    }
-
-    return cur;
+    cb(mask, "qsa_mask", il);
+    return mask;
 }
 
 ggml_tensor * llama_model_qwen4exp::graph::build_layer_attn(
@@ -707,10 +666,23 @@ ggml_tensor * llama_model_qwen4exp::graph::build_layer_attn(
     const int64_t n_embd_head = hparams.n_embd_head_v();
     GGML_ASSERT(n_embd_head == hparams.n_embd_head_k());
 
-    // indexer reads the same block input as q/k/v; no cache or no ratio means dense
-    const bool qsa = mctx_hyb->get_idx() != nullptr && hparams.dsv4_compress_ratios[il] > 0;
+    const llama_kv_cache_context * mctx_idx = mctx_hyb->get_idx();
+    if (mctx_idx) {
+        ggml_tensor * index_k = build_lora_mm(model.layers[il].index_k_proj, cur);
+        index_k = ggml_reshape_3d(ctx0, index_k, hparams.indexer_head_size, 1, n_tokens);
+        cb(index_k, "indexer_k_raw", il);
+        ggml_build_forward_expand(gf, mctx_idx->cpy_k(ctx0, index_k, inp->get_k_idxs(), il));
+    }
 
-    ggml_tensor * top_k = qsa ? build_qsa_top_k(mctx_hyb, cur, inp_pos, inp->get_kq_mask(), sections, il) : nullptr;
+    const int64_t ratio = hparams.dsv4_compress_ratios[il];
+    const bool qsa = mctx_idx && ratio > 0 &&
+            mctx_idx->get_n_kv() > (int64_t) hparams.indexer_top_k + ratio - 1;
+    if (qsa) {
+        inp->self_kq_mask_cnv = build_qsa_mask(
+                mctx_hyb, cur, inp_pos, inp->self_kq_mask, sections, il);
+    } else {
+        inp->self_kq_mask_cnv = inp->self_kq_mask;
+    }
 
     // Qwen3Next uses a single Q projection that outputs query + gate
     ggml_tensor * Qcur_full = build_lora_mm(model.layers[il].wq, cur, model.layers[il].wq_s); // [ (n_embd_head * 2) * n_head, n_tokens ]
@@ -762,13 +734,9 @@ ggml_tensor * llama_model_qwen4exp::graph::build_layer_attn(
 
     const float kq_scale = hparams.f_attention_scale == 0.0f ? 1.0f / sqrtf(float(n_embd_head)) : hparams.f_attention_scale;
 
-    if (top_k) {
-        cur = build_attn_qsa(inp, Qcur, Kcur, Vcur, top_k, kq_scale, il);
-    } else {
-        cur = build_attn(inp,
-                    nullptr, nullptr, nullptr,
-                    Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il);
-    }
+    cur = build_attn(inp,
+                nullptr, nullptr, nullptr,
+                Qcur, Kcur, Vcur, nullptr, nullptr, nullptr, kq_scale, il);
     cb(cur, "attn_pregate", il);
 
     ggml_tensor * gate_sigmoid = ggml_sigmoid(ctx0, gate);
@@ -1098,8 +1066,13 @@ ggml_tensor * llama_model_qwen4exp::graph::build_conv_state_at(
     return conv_input;
 }
 
-ggml_tensor * llama_model_qwen4exp::graph::build_inp_ple(
-        const llama_memory_hybrid_idx_context * mctx_hyb) {
+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;
     const int64_t n_heads = hparams.ple_n_heads;
 
     // the attention cells see every ubatch regardless of the layer types
@@ -1114,18 +1087,7 @@ ggml_tensor * llama_model_qwen4exp::graph::build_inp_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", -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;
+    cb(emb, "ple_embd", il);
 
     ggml_tensor * key   = build_lora_mm(model.layers[il].ple_key,   emb);
     ggml_tensor * value = build_lora_mm(model.layers[il].ple_value, emb);

tests/CMakeLists.txt

@@ -220,6 +220,16 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS)
         FIXTURES_REQUIRED generate-models
     )
 
+    llama_test(
+        test-recurrent-state-rollback
+        NAME test-recurrent-state-rollback-qwen4exp
+        LABEL main
+        ARGS -m "${MODEL_DIR}/qwen4exp-moe.gguf"
+    )
+    set_tests_properties(test-recurrent-state-rollback-qwen4exp PROPERTIES
+        FIXTURES_REQUIRED generate-models
+    )
+
     llama_test(
         test-recurrent-state-rollback
         NAME test-recurrent-state-rollback-nemotron-h

tests/test-backend-ops.cpp

@@ -8785,7 +8785,6 @@ 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}));
@@ -9359,6 +9358,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     // test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 512, 262144, 9216, {1, 1}, {1, 1}));
 
     // test large experts*tokens
+    test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 512, 10, false, 64, 512, 64));
     for (bool b : {false, true}) {
         test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 16, 16, b, 32, 1024, 16));
         test_cases.emplace_back(new test_mul_mat_id(GGML_TYPE_F16, GGML_TYPE_F32, 2, 2, b, 32, 8192, 64));

tests/test-llama-archs.cpp

@@ -79,7 +79,7 @@ static std::vector<llama_token> get_tokens(const uint32_t n_tokens, const uint32
     return ret;
 }
 
-static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
+static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe, const bool qwen_ple = false) {
     gguf_context_ptr ret(gguf_init_empty());
     llama_model_saver ms(arch, ret.get());
     const uint32_t n_ctx = 256;
@@ -252,8 +252,23 @@ static gguf_context_ptr get_gguf_ctx(const llm_arch arch, const bool moe) {
     if (arch == LLM_ARCH_QWEN4EXP) {
         ms.add_kv(LLM_KV_HYPER_CONNECTION_COUNT,    uint32_t(4));
         ms.add_kv(LLM_KV_HYPER_CONNECTION_LOW_RANK, uint32_t(8));
-        // without this the QSA layers fall back to dense and go uncovered
-        ms.add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS, std::vector<uint32_t>(n_layer, 4));
+        std::vector<uint32_t> ratios(n_layer, 0);
+        for (uint32_t il = 1; il < n_layer; il += 2) {
+            ratios[il] = 4;
+        }
+        ms.add_kv(LLM_KV_ATTENTION_COMPRESS_RATIOS, ratios);
+
+        if (qwen_ple) {
+            ms.add_kv(LLM_KV_PLE_LAYERS,                std::vector<uint32_t>({0}));
+            ms.add_kv(LLM_KV_PLE_NGRAM_SIZE,             uint32_t(2));
+            ms.add_kv(LLM_KV_PLE_HEADS_PER_NGRAM,        uint32_t(1));
+            ms.add_kv(LLM_KV_PLE_CONV_KERNEL,            uint32_t(2));
+            ms.add_kv(LLM_KV_PLE_EOS_TOKEN_ID,           uint32_t(1));
+            ms.add_kv(LLM_KV_EMBEDDING_LENGTH_PER_LAYER, uint32_t(64));
+            ms.add_kv(LLM_KV_PLE_LAYER_MULTIPLIERS,      std::vector<uint64_t>({1, 3}));
+            ms.add_kv(LLM_KV_PLE_HEAD_OFFSETS,           std::vector<uint64_t>({0}));
+            ms.add_kv(LLM_KV_PLE_HEAD_VOCAB_SIZES,       std::vector<uint64_t>({16}));
+        }
     }
 
     // minimax-m3 keeps one indexer head per GQA head; the rest use a fixed 64 to match the fused
@@ -347,7 +362,8 @@ static bool silent_model_load_progress(float /*progress*/, void * /*user_data*/)
 
 static std::pair<llama_model_ptr, llama_context_ptr> get_model_and_ctx(
         struct gguf_context * gguf_ctx, FILE * file, const size_t seed, const std::vector<ggml_backend_dev_t> & devs,
-        const llama_split_mode split_mode = LLAMA_SPLIT_MODE_LAYER, bool encode = false) {
+        const llama_split_mode split_mode = LLAMA_SPLIT_MODE_LAYER, bool encode = false,
+        ggml_backend_sched_eval_callback cb_eval = nullptr, void * cb_eval_user_data = nullptr) {
     GGML_ASSERT((gguf_ctx == nullptr) != (file == nullptr));
     llama_model_params model_params = llama_model_default_params();
     model_params.progress_callback = silent_model_load_progress;
@@ -360,6 +376,8 @@ static std::pair<llama_model_ptr, llama_context_ptr> get_model_and_ctx(
     ctx_params.n_ctx = 0;
     ctx_params.n_threads = 4;
     ctx_params.n_threads_batch = 4;
+    ctx_params.cb_eval = cb_eval;
+    ctx_params.cb_eval_user_data = cb_eval_user_data;
     if (!encode) {
         ctx_params.n_ubatch = 64;
     }
@@ -412,6 +430,54 @@ static std::vector<float> get_logits(
     return ret;
 }
 
+struct qwen4_qsa_mask_check {
+    int64_t n_seen = 0;
+    int64_t n_tokens_seen = 0;
+    bool ok = true;
+};
+
+static bool check_qwen4_qsa_mask(ggml_tensor * tensor, bool ask, void * user_data) {
+    if (strncmp(tensor->name, "qsa_mask", 8) != 0) {
+        return false;
+    }
+    if (ask) {
+        return true;
+    }
+
+    auto * check = (qwen4_qsa_mask_check *) user_data;
+    const int64_t n_kv = tensor->ne[0];
+    const int64_t n_tokens = tensor->ne[1];
+    std::vector<float> mask(ggml_nelements(tensor));
+    if (tensor->type == GGML_TYPE_F32) {
+        ggml_backend_tensor_get(tensor, mask.data(), 0, ggml_nbytes(tensor));
+    } else {
+        GGML_ASSERT(tensor->type == GGML_TYPE_F16);
+        std::vector<ggml_fp16_t> mask_f16(ggml_nelements(tensor));
+        ggml_backend_tensor_get(tensor, mask_f16.data(), 0, ggml_nbytes(tensor));
+        for (size_t i = 0; i < mask.size(); ++i) {
+            mask[i] = ggml_fp16_to_fp32(mask_f16[i]);
+        }
+    }
+
+    for (int64_t it = 0; it < n_tokens; ++it) {
+        const int64_t n_visible = check->n_tokens_seen + it + 1;
+        const int64_t n_complete = n_visible/4;
+        const int64_t expected = n_complete <= 2 ? n_visible : 8 + n_visible%4;
+        int64_t actual = 0;
+        for (int64_t ikv = 0; ikv < n_kv; ++ikv) {
+            actual += mask[it*n_kv + ikv] > -1e20f;
+        }
+        if (actual != expected) {
+            fprintf(stderr, "Qwen4 QSA mask row %lld selects %lld tokens, expected %lld\n",
+                    (long long) (check->n_tokens_seen + it), (long long) actual, (long long) expected);
+        }
+        check->ok = check->ok && actual == expected;
+    }
+    check->n_tokens_seen += n_tokens;
+    check->n_seen++;
+    return true;
+}
+
 static bool moe_mandatory(const llm_arch arch) {
     switch (arch) {
         case LLM_ARCH_LLAMA4:
@@ -686,6 +752,32 @@ static int test_backends(const llm_arch target_arch, const size_t seed, const gg
             if (arch == LLM_ARCH_BAILINGMOE3) {
                 GGML_ASSERT(gguf_remove_key(gguf_ctx.get(), "bailingmoe3.kda.safe_gate") >= 0);
             }
+            if (arch == LLM_ARCH_QWEN4EXP) {
+                qwen4_qsa_mask_check check;
+                auto model_and_ctx = get_model_and_ctx(
+                        gguf_ctx.get(), nullptr, seed, {}, LLAMA_SPLIT_MODE_LAYER, false,
+                        check_qwen4_qsa_mask, &check);
+                get_logits(model_and_ctx.first.get(), model_and_ctx.second.get(), tokens);
+                GGML_ASSERT(check.ok && check.n_seen > 0);
+
+                gguf_context_ptr gguf_ctx_ple = get_gguf_ctx(arch, moe, true);
+                auto model_and_ctx_ple = get_model_and_ctx(gguf_ctx_ple.get(), nullptr, seed, {});
+                const std::vector<float> logits_ple = get_logits(
+                        model_and_ctx_ple.first.get(), model_and_ctx_ple.second.get(), tokens);
+
+                FILE * file_ple = tmpfile();
+                GGML_ASSERT(file_ple);
+                llama_model_saver saver_ple(model_and_ctx_ple.first.get());
+                saver_ple.add_kv_from_model();
+                saver_ple.add_tensors_from_model();
+                saver_ple.save(file_ple);
+                rewind(file_ple);
+
+                auto model_and_ctx_ple_saved = get_model_and_ctx(nullptr, file_ple, seed, {});
+                const std::vector<float> logits_ple_saved = get_logits(
+                        model_and_ctx_ple_saved.first.get(), model_and_ctx_ple_saved.second.get(), tokens);
+                GGML_ASSERT(logits_ple == logits_ple_saved);
+            }
             std::pair<llama_model_ptr, llama_context_ptr> model_and_ctx_cpu;
             std::vector<float> logits_cpu;
             for (device_config & dc : dev_configs) {

tools/llama-bench/README.md

@@ -67,7 +67,6 @@ 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)

tools/llama-bench/llama-bench.cpp

@@ -271,19 +271,6 @@ 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);
@@ -354,7 +341,6 @@ 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;
@@ -399,7 +385,6 @@ 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 },
@@ -475,7 +460,6 @@ 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());
@@ -802,32 +786,6 @@ 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;
@@ -1179,9 +1137,6 @@ 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;
     }
@@ -1248,7 +1203,6 @@ 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;
@@ -1270,7 +1224,6 @@ 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;
@@ -1318,8 +1271,7 @@ 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 && lazy_mode == other.lazy_mode &&
-               devices == other.devices && no_host == other.no_host &&
+               load_mode == other.load_mode && devices == other.devices && no_host == other.no_host &&
                vec_tensor_buft_override_equal(tensor_buft_overrides, other.tensor_buft_overrides);
     }
 
@@ -1353,7 +1305,6 @@ 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)
@@ -1393,7 +1344,6 @@ 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,
@@ -1430,7 +1380,6 @@ 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,
@@ -1467,7 +1416,6 @@ 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,
@@ -1509,7 +1457,6 @@ 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;
@@ -1549,7 +1496,6 @@ 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;
@@ -1617,8 +1563,7 @@ 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",     "lazy_mode",
-            "embeddings",
+            "tensor_buft_overrides",            "load_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"
@@ -1643,7 +1588,7 @@ struct test {
         if (field == "avg_ts" || field == "stddev_ts") {
             return FLOAT;
         }
-        if (field == "load_mode" || field == "lazy_mode") {
+        if (field == "load_mode") {
             return STRING;
         }
         return STRING;
@@ -1713,7 +1658,6 @@ 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),
@@ -2028,9 +1972,6 @@ 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");
         }