spec: fix failed to decode mtmd chunk with DFlash (#28587)

* speculative: fix failed to decode mtmd chunk with DFlash

When using DFlash w/ vision models, the drafter memory fails to
allocate new tokens because images report a fixed offset. Stop copying
them to allow the drafter to continue.

* address PR feedback

limit M-RoPE skip to images only, allow audio to pass through. Clean up
comments to align to the updated implementation

.github/workflows/build-riscv.yml

@@ -106,83 +106,3 @@ jobs:
           wget https://huggingface.co/karpathy/tinyllamas/resolve/main/stories260K/stories260K.bin
           ./bin/llama-convert-llama2c-to-ggml --copy-vocab-from-model ./tok512.bin --llama2c-model stories260K.bin --llama2c-output-model stories260K.gguf
           ./bin/llama-completion -m stories260K.gguf -p "One day, Lily met a Shoggoth" -n 500 -c 256
-
-  ubuntu-riscv64-native-sanitizer:
-    runs-on: ubuntu-24.04-riscv
-
-    continue-on-error: true
-
-    strategy:
-      matrix:
-        sanitizer: [ADDRESS, THREAD, UNDEFINED]
-        build_type: [Debug]
-
-    steps:
-      - name: Install dependencies
-        run: |
-          # Set gcc-14 and g++-14 as the default compilers
-          sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-14 100
-          sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-14 100
-
-          git lfs install
-
-      - name: GCC version check
-        run: |
-          gcc --version
-          g++ --version
-
-      - name: Clone
-        id: checkout
-        uses: actions/checkout@v6
-
-      # note: sparing some ccache since these jobs run on dedicated runners that are not part of the organitzation
-      #- name: ccache
-      #  uses: ggml-org/ccache-action@v1.2.24
-      #  with:
-      #    key: riscv-ubuntu-native-sanitizer-${{ matrix.sanitizer }}-${{ matrix.build_type }}
-      #    evict-old-files: 1d
-      #    save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
-
-      - name: Build
-        id: cmake_build
-        if: ${{ matrix.sanitizer != 'THREAD' }}
-        run: |
-          cmake -B build \
-            -DLLAMA_OPENSSL=OFF \
-            -DCMAKE_BUILD_TYPE=${{ matrix.build_type }} \
-            -DGGML_OPENMP=ON \
-            -DLLAMA_BUILD_EXAMPLES=ON \
-            -DLLAMA_BUILD_TOOLS=ON \
-            -DLLAMA_BUILD_TESTS=OFF \
-            -DCMAKE_C_COMPILER_LAUNCHER=ccache \
-            -DCMAKE_CXX_COMPILER_LAUNCHER=ccache \
-            -DLLAMA_SANITIZE_${{ matrix.sanitizer }}=ON \
-            -DCMAKE_C_COMPILER=riscv64-linux-gnu-gcc-14 \
-            -DCMAKE_CXX_COMPILER=riscv64-linux-gnu-g++-14
-
-          cmake --build build --config ${{ matrix.build_type }} -j $(nproc)
-
-      - name: Build (no OpenMP)
-        id: cmake_build_no_openmp
-        if: ${{ matrix.sanitizer == 'THREAD' }}
-        run: |
-          cmake -B build \
-            -DLLAMA_OPENSSL=OFF \
-            -DCMAKE_BUILD_TYPE=${{ matrix.build_type }} \
-            -DGGML_OPENMP=OFF \
-            -DLLAMA_BUILD_EXAMPLES=ON \
-            -DLLAMA_BUILD_TOOLS=ON \
-            -DLLAMA_BUILD_TESTS=OFF \
-            -DCMAKE_C_COMPILER_LAUNCHER=ccache \
-            -DCMAKE_CXX_COMPILER_LAUNCHER=ccache \
-            -DLLAMA_SANITIZE_${{ matrix.sanitizer }}=ON \
-            -DCMAKE_C_COMPILER=riscv64-linux-gnu-gcc-14 \
-            -DCMAKE_CXX_COMPILER=riscv64-linux-gnu-g++-14
-
-          cmake --build build --config ${{ matrix.build_type }} -j $(nproc)
-
-      - name: Test
-        id: cmake_test
-        run: |
-          cd build
-          ctest -L main --verbose --timeout 900

.github/workflows/build-sanitize.yml

@@ -101,8 +101,6 @@ jobs:
 
       - name: Test
         id: cmake_test
-        # skip run in Debug - very slow
-        if: ${{ matrix.sanitizer != 'UNDEFINED' }}
         run: |
           cd build
           ctest -L main -E tokenizer --verbose --timeout 900

common/speculative.cpp

@@ -1094,8 +1094,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
 
         // Target prefill may contain token IDs or multimodal embeddings. Both
         // produce the target-layer features used to seed the draft KV cache, so
-        // skipping the embedding batches leaves a hole in the draft's cache and
-        // the next injection fails to initialize.
+        // embeddings are injected too, except the pinned ones skipped below.
         // TODO: revisit after https://github.com/ggml-org/llama.cpp/pull/24669 is merged
         const bool has_tokens     = batch_in.token != nullptr;
         const bool has_embeddings = batch_in.embd  != nullptr;
@@ -1131,6 +1130,13 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
             }
             const int32_t n_rows = i_batch_end[seq_id] - i_batch_beg[seq_id] + 1;
 
+            // an M-RoPE image pins all its rows to one position, so a windowed draft
+            // cache cannot free cells for it - skip it, the draft can jump over the gap
+            const bool pos_pinned = batch_in.pos[i_batch_beg[seq_id]] == batch_in.pos[i_batch_end[seq_id]];
+            if (has_embeddings && n_rows > 1 && pos_pinned) {
+                continue;
+            }
+
             for (int32_t offset = 0; offset < n_rows; offset += n_ubatch) {
                 const int32_t n_chunk = std::min(n_ubatch, n_rows - offset);
 

conversion/nemotron.py

@@ -216,14 +216,14 @@ class NemotronHModel(GraniteHybridModel):
         hparams = kwargs.pop("hparams", None)
         if hparams is None:
             hparams = ModelBase.load_hparams(args[0], self.is_mistral_format)
-        has_moe_params = (
-            "num_experts_per_tok" in hparams
-            or (isinstance(hparams.get("llm_config"), dict) and "num_experts_per_tok" in hparams["llm_config"])
-        )
+        llm_config = {**hparams, **(hparams.get("llm_config") or {})}
+
+        has_moe_params = "num_experts_per_tok" in llm_config
+        layers_block_type = llm_config.get("layers_block_type")
+
         if has_moe_params:
             self.model_arch = gguf.MODEL_ARCH.NEMOTRON_H_MOE
             self.is_moe = True
-        layers_block_type = hparams.get("layers_block_type")
         if layers_block_type is not None:
             hparams["num_hidden_layers"] = len(layers_block_type)
 

docs/build-s390x.md

@@ -243,6 +243,7 @@ IBM VXE/VXE2 SIMD acceleration depends on the BLAS implementation. It is strongl
 | FP32       | ✅           | ✅    | ❓     |
 | FP16       | ✅           | ✅    | ❓     |
 | BF16       | ✅           | ✅    | ❓     |
+| Q1_0       | ✅           | ❓    | ❓     |
 | Q4_0       | ✅           | ❓    | ❓     |
 | Q4_1       | ✅           | ❓    | ❓     |
 | MXFP4      | ✅           | ❓    | ❓     |
@@ -272,4 +273,4 @@ IBM VXE/VXE2 SIMD acceleration depends on the BLAS implementation. It is strongl
 -   🚫 - acceleration unavailable, will still run using scalar implementation
 -   ❓ - acceleration unknown, please contribute if you can test it yourself
 
-Last Updated by **Aaron Teo (aaron.teo1@ibm.com)** on Feb 15, 2026.
+Last Updated by **Aaron Teo (aaron.teo1@ibm.com)** on Sep 8, 2026.

ggml/src/ggml-cpu/CMakeLists.txt

@@ -520,7 +520,9 @@ function(ggml_add_cpu_backend_variant_impl tag_name)
     elseif (GGML_SYSTEM_ARCH STREQUAL "s390x")
         message(STATUS "s390x detected")
         list(APPEND GGML_CPU_SOURCES
-            ggml-cpu/arch/s390/quants.c)
+            ggml-cpu/arch/s390/quants.c
+            ggml-cpu/arch/s390/repack.cpp
+            )
 
         # for native compilation
         if (GGML_NATIVE)

ggml/src/ggml-cpu/arch-fallback.h

@@ -247,7 +247,6 @@
 // quants.c
 #define quantize_row_q8_K_generic quantize_row_q8_K
 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0
-#define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K
 #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K
@@ -260,11 +259,9 @@
 #define ggml_vec_dot_iq1_s_q8_K_generic ggml_vec_dot_iq1_s_q8_K
 #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
 // repack.cpp
-#define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4
 #define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8
 #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4
 #define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8
-#define ggml_gemv_q4_0_4x4_q8_0_generic ggml_gemv_q4_0_4x4_q8_0
 #define ggml_gemv_q4_0_4x8_q8_0_generic ggml_gemv_q4_0_4x8_q8_0
 #define ggml_gemv_q4_0_8x8_q8_0_generic ggml_gemv_q4_0_8x8_q8_0
 #define ggml_gemv_q2_K_8x8_q8_K_generic ggml_gemv_q2_K_8x8_q8_K
@@ -280,7 +277,6 @@
 #define ggml_gemv_mxfp4_8x8_q8_0_generic ggml_gemv_mxfp4_8x8_q8_0
 #define ggml_gemv_q8_0_4x4_q8_0_generic ggml_gemv_q8_0_4x4_q8_0
 #define ggml_gemv_q8_0_4x8_q8_0_generic ggml_gemv_q8_0_4x8_q8_0
-#define ggml_gemm_q4_0_4x4_q8_0_generic ggml_gemm_q4_0_4x4_q8_0
 #define ggml_gemm_q4_0_4x8_q8_0_generic ggml_gemm_q4_0_4x8_q8_0
 #define ggml_gemm_q4_0_8x8_q8_0_generic ggml_gemm_q4_0_8x8_q8_0
 #define ggml_gemm_q2_K_8x8_q8_K_generic ggml_gemm_q2_K_8x8_q8_K

ggml/src/ggml-cpu/arch/s390/quants.c

@@ -146,6 +146,74 @@ void quantize_row_q8_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, i
 
 //===================================== Dot products =================================
 
+void ggml_vec_dot_q1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
+    const int qk = QK1_0;  // 128
+    const int nb = n / qk;
+
+    assert(n % qk == 0);
+    assert(nrc == 1);
+    UNUSED(nrc);
+    UNUSED(bx);
+    UNUSED(by);
+    UNUSED(bs);
+
+    const block_q1_0 * GGML_RESTRICT x = vx;
+    const block_q8_0 * GGML_RESTRICT y = vy;
+
+#if defined(__VXE__) || defined(__VXE2__)
+    float32x4_t v_sumf = vec_splats(0.0f);
+
+    const uint8x16_t v_zero = vec_splats((uint8_t)0x00);  // zero
+    const uint8x16_t v_bias = vec_splats((uint8_t)0x80);  // bias from signed to unsigned
+                                                          // v ^ 0x80 == v + 128
+
+    const uint8x16_t v_idx = (const uint8x16_t){ 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1 };
+    const uint8x16_t v_bit = (const uint8x16_t){ 1, 2, 4, 8, 16, 32, 64, 128, 1, 2, 4, 8, 16, 32, 64, 128 };
+
+    for (int i = 0; i < nb; ++i) {
+        const uint8x16_t  v_x  = vec_xl(0, (const uint8_t *)x[i].qs);
+        const float32x4_t v_xd = vec_splats(GGML_CPU_FP16_TO_FP32(x[i].d));
+
+        for (int k = 0; k < 4; ++k) {
+            // sub-block k holds elements 32k .. 32k+31
+            const block_q8_0 * GGML_RESTRICT yb = &y[i*4 + k];
+            const float32x4_t v_yd = vec_splats(GGML_CPU_FP16_TO_FP32(yb->d));
+
+            const uint8x16_t v_xrl = vec_perm(v_x, v_x, vec_add(v_idx, vec_splats((uint8_t)(k*4 + 0))));
+            const uint8x16_t v_xrh = vec_perm(v_x, v_x, vec_add(v_idx, vec_splats((uint8_t)(k*4 + 2))));
+
+            // isolate each lane's bit, then set all ones where that bit is clear, the -d case
+            const int8x16_t v_ml = (int8x16_t)vec_cmpeq(vec_and(v_xrl, v_bit), v_zero);
+            const int8x16_t v_mh = (int8x16_t)vec_cmpeq(vec_and(v_xrh, v_bit), v_zero);
+
+            const int8x16_t v_yl = vec_xl(0,       (const int8_t *)yb->qs);
+            const int8x16_t v_yh = vec_xl(QK8_0/2, (const int8_t *)yb->qs);
+
+            // weights are only +1 or -1, so negate y
+            const int8x16_t v_ysl = vec_sub(vec_xor(v_yl, v_ml), v_ml);
+            const int8x16_t v_ysh = vec_sub(vec_xor(v_yh, v_mh), v_mh);
+
+            // bias to unsigned, then vec_sum4 adds each group of 4 bytes into one word
+            const uint32x4_t v_p = vec_add(vec_sum4(vec_xor((uint8x16_t)v_ysl, v_bias), v_zero),
+                                           vec_sum4(vec_xor((uint8x16_t)v_ysh, v_bias), v_zero));
+
+            // each word summed 8 biased bytes, so take back 8 * 128
+            const int32x4_t v_xy = vec_sub((int32x4_t)v_p, vec_splats((int32_t)1024));
+
+            // apply both block scales and add into the running total
+            v_sumf = vec_madd(vec_float(v_xy), vec_mul(v_xd, v_yd), v_sumf);
+        }
+    }
+
+    *s = vec_hsum_f32x4(v_sumf);
+#else
+    UNUSED(nb);
+    UNUSED(x);
+    UNUSED(y);
+    ggml_vec_dot_q1_0_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc);
+#endif
+}
+
 void ggml_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
     const int qk = QK8_0;
     const int nb = n / qk;

ggml/src/ggml-cpu/arch/s390/repack.cpp

@@ -0,0 +1,223 @@
+#define GGML_COMMON_IMPL_CPP
+#define GGML_COMMON_DECL_CPP
+#include "ggml-common.h"
+#include "ggml-backend-impl.h"
+
+#include "ggml-impl.h"
+#include "ggml-cpu.h"
+#include "ggml-cpu-impl.h"
+#include "simd-mappings.h"
+#include "traits.h"
+
+#include <cmath>
+#include <cstring>
+#include <cassert>
+
+#define GGML_CPU_CLANG_WORKAROUND
+#include "../../repack.h"
+
+#define UNUSED GGML_UNUSED
+
+void ggml_quantize_mat_q8_0_4x4(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) {
+    assert(QK8_0 == 32);
+    assert(k % QK8_0 == 0);
+    const int nb = k / QK8_0;
+
+    block_q8_0x4 * GGML_RESTRICT y = (block_q8_0x4 *) vy;
+
+#if defined(__VXE__) || defined(__VXE2__)
+    float32x4_t v_src[4][8];
+    float       id[4];
+
+    for (int i = 0; i < nb; i++) {
+        float32x4_t v_asrc[8];
+        float32x4_t v_amax[8];
+
+        for (int row_iter = 0; row_iter < 4; row_iter++) {
+            for (int j = 0; j < 8; j++) v_src[row_iter][j] = vec_xl(0, x + row_iter * k + i * 32 + 4 * j);
+            for (int j = 0; j < 8; j++) v_asrc[j] = vec_abs(v_src[row_iter][j]);
+
+            for (int j = 0; j < 4; j++) v_amax[2 * j] = vec_max(v_asrc[2 * j], v_asrc[2 * j + 1]);
+            for (int j = 0; j < 2; j++) v_amax[4 * j] = vec_max(v_amax[4 * j], v_amax[4 * j + 2]);
+            for (int j = 0; j < 1; j++) v_amax[8 * j] = vec_max(v_amax[8 * j], v_amax[8 * j + 4]);
+
+            const float amax = MAX(MAX(vec_extract(v_amax[0], 0), vec_extract(v_amax[0], 1)),
+                                   MAX(vec_extract(v_amax[0], 2), vec_extract(v_amax[0], 3)));
+
+            const float d = amax / ((1 << 7) - 1);
+            id[row_iter] = d ? 1.0f / d : 0.0f;
+
+            y[i].d[row_iter] = GGML_CPU_FP32_TO_FP16(d);
+        }
+
+        for (int j = 0; j < 8; j++) {
+            /* Uses non-default rounding for vec_signed or vec_round */
+            const int32x4_t v_qs0 = vec_signed(__builtin_s390_vfisb(vec_mul(v_src[0][j], id[0]), 4, 1));
+            const int32x4_t v_qs1 = vec_signed(__builtin_s390_vfisb(vec_mul(v_src[1][j], id[1]), 4, 1));
+            const int32x4_t v_qs2 = vec_signed(__builtin_s390_vfisb(vec_mul(v_src[2][j], id[2]), 4, 1));
+            const int32x4_t v_qs3 = vec_signed(__builtin_s390_vfisb(vec_mul(v_src[3][j], id[3]), 4, 1));
+
+            const int16x8_t v_qs01 = vec_packs(v_qs0, v_qs1);
+            const int16x8_t v_qs23 = vec_packs(v_qs2, v_qs3);
+
+            vec_xst(vec_packs(v_qs01, v_qs23), 0, y[i].qs + 16 * j);
+        }
+    }
+#else
+    UNUSED(nb);
+    UNUSED(y);
+    ggml_quantize_mat_q8_0_4x4_generic(x, vy, k);
+#endif
+}
+
+static inline int16x8_t vxe_dot_acc(const int8x16_t v_x, const int8x16_t v_y, const int16x8_t v_acc) {
+    return vec_meadd(v_x, v_y, vec_moadd(v_x, v_y, v_acc));
+}
+
+static inline int8x16_t vxe_splat_granule(const int8_t * qs) {
+    uint32_t g;
+    memcpy(&g, qs, sizeof(g));
+    return (int8x16_t)vec_splats(g);
+}
+
+static inline int32x4_t vxe_fold(const int16x8_t v_sumi) {
+    const int16x8_t v_ones = vec_splats((int16_t)1);
+    return vec_add(vec_mule(v_sumi, v_ones), vec_mulo(v_sumi, v_ones));
+}
+
+void ggml_gemv_q4_0_4x4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int nr, int nc) {
+    const int qk = QK8_0;
+    const int nb = n / qk;
+    const int ncols_interleaved = 4;
+
+    assert(nr == 1);
+    assert(n % qk == 0);
+    assert(nc % ncols_interleaved == 0);
+
+    UNUSED(bs);
+    UNUSED(nr);
+
+#if defined(__VXE__) || defined(__VXE2__)
+    const block_q8_0 * a_ptr = (const block_q8_0 *) vy;
+    float * res_ptr = s;
+
+    for (int x = 0; x < nc / ncols_interleaved; x++) {
+        const block_q4_0x4 * b_ptr = (const block_q4_0x4 *) vx + (x * nb);
+
+        float32x4_t v_sumf = vec_splats(0.0f);
+
+        for (int l = 0; l < nb; l++) {
+            const int8_t * x_qs = b_ptr[l].qs;
+
+            const int8x16_t v_x0 = vec_xl( 0, x_qs);
+            const int8x16_t v_x1 = vec_xl(16, x_qs);
+            const int8x16_t v_x2 = vec_xl(32, x_qs);
+            const int8x16_t v_x3 = vec_xl(48, x_qs);
+
+            const int8x16_t v_x0l = vec_sra(vec_sl(v_x0, 4), 4);
+            const int8x16_t v_x1l = vec_sra(vec_sl(v_x1, 4), 4);
+            const int8x16_t v_x2l = vec_sra(vec_sl(v_x2, 4), 4);
+            const int8x16_t v_x3l = vec_sra(vec_sl(v_x3, 4), 4);
+
+            const int8x16_t v_x0h = vec_sra(v_x0, 4);
+            const int8x16_t v_x1h = vec_sra(v_x1, 4);
+            const int8x16_t v_x2h = vec_sra(v_x2, 4);
+            const int8x16_t v_x3h = vec_sra(v_x3, 4);
+
+            const int8_t * y_lo = a_ptr[l].qs;
+            const int8_t * y_hi = y_lo + qk / 2;
+
+            int16x8_t v_sumi = vec_splats((int16_t)0);
+
+            v_sumi = vxe_dot_acc(v_x0l, vxe_splat_granule(y_lo +  0), v_sumi);
+            v_sumi = vxe_dot_acc(v_x1l, vxe_splat_granule(y_lo +  4), v_sumi);
+            v_sumi = vxe_dot_acc(v_x2l, vxe_splat_granule(y_lo +  8), v_sumi);
+            v_sumi = vxe_dot_acc(v_x3l, vxe_splat_granule(y_lo + 12), v_sumi);
+
+            v_sumi = vxe_dot_acc(v_x0h, vxe_splat_granule(y_hi +  0), v_sumi);
+            v_sumi = vxe_dot_acc(v_x1h, vxe_splat_granule(y_hi +  4), v_sumi);
+            v_sumi = vxe_dot_acc(v_x2h, vxe_splat_granule(y_hi +  8), v_sumi);
+            v_sumi = vxe_dot_acc(v_x3h, vxe_splat_granule(y_hi + 12), v_sumi);
+
+            const float32x4_t v_yd = vec_splats(GGML_CPU_FP16_TO_FP32(a_ptr[l].d));
+            const float32x4_t v_xd = __lzs_f16cx4_load(b_ptr[l].d);
+            const float32x4_t v_d  = vec_mul(v_yd, v_xd);
+
+            v_sumf = vec_madd(vec_float(vxe_fold(v_sumi)), v_d, v_sumf);
+        }
+
+        vec_xst(v_sumf, 0, res_ptr + x * ncols_interleaved);
+    }
+#else
+    UNUSED(nb);
+    UNUSED(ncols_interleaved);
+    ggml_gemv_q4_0_4x4_q8_0_generic(n, s, bs, vx, vy, nr, nc);
+#endif
+}
+
+void ggml_gemm_q4_0_4x4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int nr, int nc) {
+    const int qk = QK8_0;
+    const int nb = n / qk;
+    const int ncols_interleaved = 4;
+
+    assert(nr % 4 == 0);
+    assert(n % qk == 0);
+    assert(nc % ncols_interleaved == 0);
+
+#if defined(__VXE__) || defined(__VXE2__)
+    for (int y = 0; y < nr / 4; y++) {
+        const block_q8_0x4 * a_ptr = (const block_q8_0x4 *) vy + (y * nb);
+
+        for (int x = 0; x < nc / ncols_interleaved; x++) {
+            const block_q4_0x4 * b_ptr = (const block_q4_0x4 *) vx + (x * nb);
+
+            float32x4_t v_sumf[4];
+            for (int m = 0; m < 4; m++) {
+                v_sumf[m] = vec_splats(0.0f);
+            }
+
+            for (int l = 0; l < nb; l++) {
+                int16x8_t v_sumi0 = vec_splats((int16_t)0);
+                int16x8_t v_sumi1 = vec_splats((int16_t)0);
+                int16x8_t v_sumi2 = vec_splats((int16_t)0);
+                int16x8_t v_sumi3 = vec_splats((int16_t)0);
+
+                for (int k = 0; k < 4; k++) {
+                    const int8x16_t v_x  = vec_xl(0, b_ptr[l].qs + 16 * k);
+                    const int8x16_t v_xl = vec_sra(vec_sl(v_x, 4), 4);
+                    const int8x16_t v_xh = vec_sra(v_x, 4);
+
+                    const int8_t * y_lo = a_ptr[l].qs + 16 * k;
+                    const int8_t * y_hi = y_lo + qk / 2 * 4;
+
+                    v_sumi0 = vxe_dot_acc(v_xl, vxe_splat_granule(y_lo +  0), v_sumi0);
+                    v_sumi1 = vxe_dot_acc(v_xl, vxe_splat_granule(y_lo +  4), v_sumi1);
+                    v_sumi2 = vxe_dot_acc(v_xl, vxe_splat_granule(y_lo +  8), v_sumi2);
+                    v_sumi3 = vxe_dot_acc(v_xl, vxe_splat_granule(y_lo + 12), v_sumi3);
+
+                    v_sumi0 = vxe_dot_acc(v_xh, vxe_splat_granule(y_hi +  0), v_sumi0);
+                    v_sumi1 = vxe_dot_acc(v_xh, vxe_splat_granule(y_hi +  4), v_sumi1);
+                    v_sumi2 = vxe_dot_acc(v_xh, vxe_splat_granule(y_hi +  8), v_sumi2);
+                    v_sumi3 = vxe_dot_acc(v_xh, vxe_splat_granule(y_hi + 12), v_sumi3);
+                }
+
+                const float32x4_t v_yd = __lzs_f16cx4_load(a_ptr[l].d);
+                const float32x4_t v_xd = __lzs_f16cx4_load(b_ptr[l].d);
+
+                v_sumf[0] = vec_madd(vec_float(vxe_fold(v_sumi0)), vec_mul(v_xd, vec_splat(v_yd, 0)), v_sumf[0]);
+                v_sumf[1] = vec_madd(vec_float(vxe_fold(v_sumi1)), vec_mul(v_xd, vec_splat(v_yd, 1)), v_sumf[1]);
+                v_sumf[2] = vec_madd(vec_float(vxe_fold(v_sumi2)), vec_mul(v_xd, vec_splat(v_yd, 2)), v_sumf[2]);
+                v_sumf[3] = vec_madd(vec_float(vxe_fold(v_sumi3)), vec_mul(v_xd, vec_splat(v_yd, 3)), v_sumf[3]);
+            }
+
+            for (int m = 0; m < 4; m++) {
+                vec_xst(v_sumf[m], 0, s + (y * 4 + m) * bs + x * ncols_interleaved);
+            }
+        }
+    }
+#else
+    UNUSED(nb);
+    UNUSED(ncols_interleaved);
+    ggml_gemm_q4_0_4x4_q8_0_generic(n, s, bs, vx, vy, nr, nc);
+#endif
+}

ggml/src/ggml-cpu/repack.cpp

@@ -4586,6 +4586,11 @@ static const ggml::cpu::tensor_traits * ggml_repack_get_optimal_repack_type(cons
                 return &q4_0_4x4_q8_0;
             }
         }
+        if (ggml_cpu_has_vxe()) {
+            if (cur->ne[1] % 4 == 0) {
+                return &q4_0_4x4_q8_0;
+            }
+        }
         if (ggml_cpu_has_riscv_v()) {
             #if defined __riscv_zvfh
             switch (__riscv_vlenb() * 8) {

ggml/src/ggml-hexagon/ggml-hexagon.cpp

@@ -56,6 +56,7 @@
 #include "htp/unary-ops.h"
 #include "htp/get-rows-ops.h"
 #include "htp/set-rows-ops.h"
+#include "htp/rope-ops.h"
 #include "htp_iface.h"
 #include "htp-drv.h"
 
@@ -299,6 +300,12 @@ static void ggml_hexagon_precompute_set_rows_params(
     struct htp_set_rows_kernel_params * kparams
 );
 
+static void ggml_hexagon_precompute_rope_params(
+    const struct ggml_hexagon_session * sess,
+    const struct ggml_tensor * op,
+    struct htp_rope_kernel_params * kparams
+);
+
 static void ggml_hexagon_precompute_fused_mmnx_params(
     const struct ggml_hexagon_session * sess,
     const struct ggml_tensor * src0,
@@ -4148,6 +4155,36 @@ static void ggml_hexagon_precompute_set_rows_params(
     kparams->vtcm_size = vtcm_layout.total_bytes;
 }
 
+static void ggml_hexagon_precompute_rope_params(
+    const struct ggml_hexagon_session * sess,
+    const struct ggml_tensor * op,
+    struct htp_rope_kernel_params * kparams
+) {
+    memset(kparams, 0, sizeof(*kparams));
+
+    const struct ggml_tensor * src0 = op->src[0];
+    const struct ggml_tensor * dst  = op;
+
+    const uint32_t src0_nrows = src0->ne[1] * src0->ne[2] * src0->ne[3];
+    const uint32_t n_threads  = (std::min)((uint32_t) sess->n_threads, src0_nrows);
+
+    struct htp_rope_vtcm_layout layout;
+    htp_rope_vtcm_layout_build(&layout, src0->ne[0], n_threads);
+
+    kparams->n_threads              = n_threads;
+    kparams->src0_nrows             = src0_nrows;
+    kparams->src0_nrows_per_thread  = (src0_nrows + n_threads - 1) / n_threads;
+    kparams->vtcm_size              = (uint32_t) layout.total_bytes;
+    kparams->spad_per_thread        = (uint32_t) layout.bytes_per_thread;
+    kparams->theta_cache_offset     = (uint32_t) layout.theta_cache_size_aligned;
+    kparams->src0_row_size_aligned  = (uint32_t) layout.src0_row_size_aligned;
+
+    if (src0_nrows > 0) {
+        kparams->div_ne2_ne1 = init_fastdiv_values(dst->ne[2] * dst->ne[1]);
+        kparams->div_ne1     = init_fastdiv_values(dst->ne[1]);
+    }
+}
+
 static void ggml_hexagon_precompute_fused_mmnx_params(
     const struct ggml_hexagon_session * sess,
     const struct ggml_tensor * src0, // W0
@@ -4706,56 +4743,82 @@ static bool ggml_hexagon_supported_argsort(const struct ggml_hexagon_session * s
 }
 
 static bool ggml_hexagon_supported_rope(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
-    const int32_t * op_params = &op->op_params[0];
-
-    // ggml_rope_set_offset: HVX kernels need a VLEN-aligned window start (32 f32 elems)
-    if (op_params[15] % 32 != 0) {
-        return false;
-    }
-
-    int mode = op_params[2];
-
-    // n_dims == ne0/2, so the rotation spans the full row
-    if (mode == GGML_ROPE_TYPE_VISION) {
-        const int n_dims = op_params[1];
-        if (n_dims != (int) (op->src[0]->ne[0] / 2)) {
-            return false;
-        }
-    }
-    if (mode & 1) {
-        return false;
-    }
-
     const struct ggml_tensor * src0 = op->src[0];
     const struct ggml_tensor * src1 = op->src[1];
     const struct ggml_tensor * src2 = op->src[2];
     const struct ggml_tensor * dst  = op;
 
-    if (src0->type != GGML_TYPE_F32) {
-        return false;  // FIXME: add support for GGML_TYPE_F16 for src0
-    }
-    if (dst->type != GGML_TYPE_F32) {
+    if (!ggml_are_same_shape(src0, dst)) {
         return false;
     }
-    if (src1->type != GGML_TYPE_I32) {
+
+    if (src0->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32 || src1->type != GGML_TYPE_I32) {
         return false;
     }
-    if (src2) {
-        if (src2->type != GGML_TYPE_F32) {
-            return false;
-        }
-        int n_dims = op_params[1];
-        if (src2->ne[0] < (n_dims / 2)) {
+
+    if (src0->ne[0] <= 0) {
+        return false;
+    }
+
+    const uint32_t src0_nrows = src0->ne[1] * src0->ne[2] * src0->ne[3];
+    if (src0_nrows == 0) {
+        return false;
+    }
+
+    const int32_t * op_params = &op->op_params[0];
+    const int n_dims = op_params[1];
+    const int mode   = op_params[2];
+    const int n_offs = op_params[15];
+
+    if (n_dims <= 0 || n_dims % 2 != 0) {
+        return false;
+    }
+
+    // ggml_rope_set_offset: HVX kernels need a VLEN-aligned window start (32 f32 elems)
+    if (n_offs < 0 || (n_offs % 32 != 0) || (n_offs + n_dims > src0->ne[0])) {
+        return false;
+    }
+
+    float freq_base;
+    memcpy(&freq_base, op_params + 5, sizeof(float));
+    if (freq_base <= 0.0f) {
+        return false;
+    }
+
+    if (mode != GGML_ROPE_TYPE_NORMAL &&
+        mode != GGML_ROPE_TYPE_NEOX &&
+        mode != GGML_ROPE_TYPE_MROPE &&
+        mode != GGML_ROPE_TYPE_VISION &&
+        mode != GGML_ROPE_TYPE_IMROPE) {
+        return false;
+    }
+
+    const bool is_mrope = (mode & GGML_ROPE_TYPE_MROPE) != 0;
+
+    // n_dims == ne0/2, so the rotation spans the full row
+    if (mode == GGML_ROPE_TYPE_VISION) {
+        if (n_dims != (int) (src0->ne[0] / 2) || n_offs != 0) {
             return false;
         }
     }
 
-    if (src2) {
-        if (!ggml_is_contiguous(src1) || !ggml_is_contiguous(src2)) {
+    if (is_mrope) {
+        const int32_t * sections = op_params + 11;
+        if (sections[0] <= 0 && sections[1] <= 0 && sections[2] <= 0) {
             return false;
         }
-    } else {
-        if (!ggml_is_contiguous(src1)) {
+    }
+
+    const int64_t min_pos_len = (is_mrope || mode == GGML_ROPE_TYPE_VISION) ? src0->ne[2] * 4 : src0->ne[2];
+    if (src1->ne[0] < min_pos_len || !ggml_is_contiguous(src1)) {
+        return false;
+    }
+
+    if (src2) {
+        if (src2->type != GGML_TYPE_F32 || !ggml_is_contiguous(src2)) {
+            return false;
+        }
+        if (src2->ne[0] < (n_dims / 2)) {
             return false;
         }
     }
@@ -4768,9 +4831,16 @@ static bool ggml_hexagon_supported_rope(const struct ggml_hexagon_session * sess
     if (src0->nb[1] < src0->ne[0] * sizeof(float) || dst->nb[1] < dst->ne[0] * sizeof(float)) {
         return false;
     }
-    return true;
 
-    GGML_UNUSED(sess);
+    const uint32_t n_threads = (std::min)((uint32_t) sess->n_threads, src0_nrows);
+
+    struct htp_rope_vtcm_layout layout;
+    htp_rope_vtcm_layout_build(&layout, src0->ne[0], n_threads);
+    if (layout.total_bytes > sess->vtcm_size) {
+        return false;
+    }
+
+    return true;
 }
 
 static bool ggml_hexagon_supported_ssm_conv(const struct ggml_hexagon_session * sess, const struct ggml_tensor * op) {
@@ -5206,6 +5276,11 @@ static ggml_status ggml_backend_hexagon_graph_compute(ggml_backend_t backend, gg
                     node.node->src[0], node.node->src[1], node.dst(),
                     (struct htp_set_rows_kernel_params *)node.kernel_params
                 );
+            } else if (node.opcode == HTP_OP_ROPE) {
+                ggml_hexagon_precompute_rope_params(sess,
+                    node.node,
+                    (struct htp_rope_kernel_params *)node.kernel_params
+                );
             }
             computed_nodes.push_back(std::move(node));
         }

ggml/src/ggml-hexagon/htp/hvx-sin-cos.h

@@ -4,87 +4,75 @@
 #include "hvx-base.h"
 #include "hvx-floor.h"
 
-static inline HVX_Vector hvx_vec_cos_f32(HVX_Vector x) {
-    HVX_Vector const_inv_pi = hvx_vec_splat_f32(0.3183098861837907f);
-    HVX_Vector const_half   = hvx_vec_splat_f32(0.5f);
-    HVX_Vector const_pi     = hvx_vec_splat_f32(3.141592653589793f);
-    HVX_Vector const_one    = hvx_vec_splat_f32(1.0f);
+// Range-reduce x to y in [-pi/2, pi/2] and the quadrant sign (-1)^n.
+// Floor/truncate need IEEE bits, so convert qf32 back to sf before them.
+static inline void hvx_vec_sincos_reduce_f32(HVX_Vector x, HVX_Vector * y, HVX_Vector * sign) {
+    HVX_Vector const_inv_pi  = hvx_vec_splat_f32(0.3183098861837907f);
+    HVX_Vector const_half    = hvx_vec_splat_f32(0.5f);
+    HVX_Vector const_pi      = hvx_vec_splat_f32(3.141592653589793f);
+    HVX_Vector const_one     = hvx_vec_splat_f32(1.0f);
     HVX_Vector const_neg_one = hvx_vec_splat_f32(-1.0f);
+    HVX_Vector const_one_i   = Q6_V_vsplat_R(1);
 
-    // n = floor(x * (1/pi) + 0.5)
-    HVX_Vector n_float = hvx_vec_floor_f32(hvx_vec_add_f32_f32(hvx_vec_mul_f32_f32(x, const_inv_pi), const_half));
+    HVX_Vector x_over_pi = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(x, const_inv_pi));
+    x_over_pi = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(x_over_pi, const_half));
 
-    // y = x - n * pi
-    HVX_Vector y = hvx_vec_sub_f32_f32(x, hvx_vec_mul_f32_f32(n_float, const_pi));
+    HVX_Vector n_float = hvx_vec_floor_f32(x_over_pi);
+    HVX_Vector n_int   = hvx_vec_truncate_f32(n_float);
 
-    // Sign determination: if n is odd, sign is -1.0f, else 1.0f
-    // half_n = n * 0.5f
-    HVX_Vector half_n = hvx_vec_mul_f32_f32(n_float, const_half);
-    // floor_half_n = floor(half_n)
-    HVX_Vector floor_half_n = hvx_vec_floor_f32(half_n);
-    // is_odd = half_n > floor_half_n
-    HVX_VectorPred is_odd = Q6_Q_vcmp_gt_VsfVsf(half_n, floor_half_n);
-    // sign = vmux(is_odd, -1.0f, 1.0f)
-    HVX_Vector sign = Q6_V_vmux_QVV(is_odd, const_neg_one, const_one);
+    HVX_Vector n_pi = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(n_float, const_pi));
+    *y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vsub_VsfVsf(x, n_pi));
 
-    // z = y^2
-    HVX_Vector z = hvx_vec_mul_f32_f32(y, y);
+    HVX_VectorPred is_odd = Q6_Q_vcmp_eq_VwVw(Q6_V_vand_VV(n_int, const_one_i), const_one_i);
+    *sign = Q6_V_vmux_QVV(is_odd, const_neg_one, const_one);
+}
+
+static inline void hvx_vec_sincos_f32(HVX_Vector x, HVX_Vector * vcos, HVX_Vector * vsin) {
+    HVX_Vector y;
+    HVX_Vector sign;
+    hvx_vec_sincos_reduce_f32(x, &y, &sign);
+
+    HVX_Vector z = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(y, y));
 
-    // Chebyshev approximation for cos(y)
     HVX_Vector c4 = hvx_vec_splat_f32(2.3557242013849433e-05f);
     HVX_Vector c3 = hvx_vec_splat_f32(-0.0013871428263450528f);
     HVX_Vector c2 = hvx_vec_splat_f32(0.041665895266688284f);
     HVX_Vector c1 = hvx_vec_splat_f32(-0.4999999360426369f);
     HVX_Vector c0 = hvx_vec_splat_f32(0.9999999999071725f);
 
-    HVX_Vector cos_y = hvx_vec_add_f32_f32(c3, hvx_vec_mul_f32_f32(z, c4));
-    cos_y = hvx_vec_add_f32_f32(c2, hvx_vec_mul_f32_f32(z, cos_y));
-    cos_y = hvx_vec_add_f32_f32(c1, hvx_vec_mul_f32_f32(z, cos_y));
-    cos_y = hvx_vec_add_f32_f32(c0, hvx_vec_mul_f32_f32(z, cos_y));
+    HVX_Vector cos_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(c3, Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(z, c4))));
+    cos_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(c2, Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(z, cos_y))));
+    cos_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(c1, Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(z, cos_y))));
+    cos_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(c0, Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(z, cos_y))));
 
-    return hvx_vec_mul_f32_f32(cos_y, sign);
-}
-
-static inline HVX_Vector hvx_vec_sin_f32(HVX_Vector x) {
-    HVX_Vector const_inv_pi = hvx_vec_splat_f32(0.3183098861837907f);
-    HVX_Vector const_half   = hvx_vec_splat_f32(0.5f);
-    HVX_Vector const_pi     = hvx_vec_splat_f32(3.141592653589793f);
-    HVX_Vector const_one    = hvx_vec_splat_f32(1.0f);
-    HVX_Vector const_neg_one = hvx_vec_splat_f32(-1.0f);
-
-    // n = floor(x * (1/pi) + 0.5)
-    HVX_Vector n_float = hvx_vec_floor_f32(hvx_vec_add_f32_f32(hvx_vec_mul_f32_f32(x, const_inv_pi), const_half));
-
-    // y = x - n * pi
-    HVX_Vector y = hvx_vec_sub_f32_f32(x, hvx_vec_mul_f32_f32(n_float, const_pi));
-
-    // Sign determination: if n is odd, sign is -1.0f, else 1.0f
-    // half_n = n * 0.5f
-    HVX_Vector half_n = hvx_vec_mul_f32_f32(n_float, const_half);
-    // floor_half_n = floor(half_n)
-    HVX_Vector floor_half_n = hvx_vec_floor_f32(half_n);
-    // is_odd = half_n > floor_half_n
-    HVX_VectorPred is_odd = Q6_Q_vcmp_gt_VsfVsf(half_n, floor_half_n);
-    // sign = vmux(is_odd, -1.0f, 1.0f)
-    HVX_Vector sign = Q6_V_vmux_QVV(is_odd, const_neg_one, const_one);
-
-    // z = y^2
-    HVX_Vector z = hvx_vec_mul_f32_f32(y, y);
-
-    // Chebyshev approximation for sin(y)
     HVX_Vector s4 = hvx_vec_splat_f32(2.642186986152672e-06f);
     HVX_Vector s3 = hvx_vec_splat_f32(-0.00019825318964070864f);
     HVX_Vector s2 = hvx_vec_splat_f32(0.00833326283319605f);
     HVX_Vector s1 = hvx_vec_splat_f32(-0.16666666082087775f);
     HVX_Vector s0 = hvx_vec_splat_f32(0.999999999915155f);
 
-    HVX_Vector sin_y = hvx_vec_add_f32_f32(s3, hvx_vec_mul_f32_f32(z, s4));
-    sin_y = hvx_vec_add_f32_f32(s2, hvx_vec_mul_f32_f32(z, sin_y));
-    sin_y = hvx_vec_add_f32_f32(s1, hvx_vec_mul_f32_f32(z, sin_y));
-    sin_y = hvx_vec_add_f32_f32(s0, hvx_vec_mul_f32_f32(z, sin_y));
-    sin_y = hvx_vec_mul_f32_f32(y, sin_y);
+    HVX_Vector sin_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(s3, Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(z, s4))));
+    sin_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(s2, Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(z, sin_y))));
+    sin_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(s1, Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(z, sin_y))));
+    sin_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_VsfVsf(s0, Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(z, sin_y))));
+    sin_y = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(y, sin_y));
 
-    return hvx_vec_mul_f32_f32(sin_y, sign);
+    *vcos = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(cos_y, sign));
+    *vsin = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vmpy_VsfVsf(sin_y, sign));
+}
+
+static inline HVX_Vector hvx_vec_cos_f32(HVX_Vector x) {
+    HVX_Vector vcos;
+    HVX_Vector vsin;
+    hvx_vec_sincos_f32(x, &vcos, &vsin);
+    return vcos;
+}
+
+static inline HVX_Vector hvx_vec_sin_f32(HVX_Vector x) {
+    HVX_Vector vcos;
+    HVX_Vector vsin;
+    hvx_vec_sincos_f32(x, &vcos, &vsin);
+    return vsin;
 }
 
 #endif /* HVX_SIN_COS_H */

ggml/src/ggml-hexagon/htp/rope-ops.c

@@ -17,8 +17,8 @@
 #include "ggml-common.h"
 #include "htp-ctx.h"
 #include "htp-ops.h"
-#include "htp-ops.h"
 #include "htp-tensor.h"
+#include "rope-ops.h"
 
 // Redefined the rope type constants as we can't include ggml.h
 #define HTP_ROPE_TYPE_NORMAL 0
@@ -27,9 +27,6 @@
 #define HTP_ROPE_TYPE_VISION 24
 #define HTP_ROPE_TYPE_IMROPE 40
 
-#define HTP_ROPE_SPAD_NROWS  16
-#define HTP_ROPE_SPAD_BLOCK  (HTP_ROPE_SPAD_NROWS/2)
-
 #define htp_rope_preamble              \
     const uint32_t ne00 = src0->ne[0]; \
     const uint32_t ne01 = src0->ne[1]; \
@@ -65,26 +62,27 @@ struct htp_rope_context {
     float beta_fast;
     float beta_slow;
     float theta_scale;
+    float theta_scale_32;
+    float theta_powers[32];
     float corr_dims[2];
 
     uint32_t src0_nrows_per_thread;
-    size_t spad_stride;
 
     struct htp_ops_context * octx;
 
+    uint8_t * vtcm_base;
+    size_t    spad_per_thread;
+    size_t    theta_cache_offset;
+
     size_t src0_row_size;
     size_t src0_row_stride;
     size_t dst_row_size;
     size_t dst_row_stride;
     size_t src0_row_size_aligned;
-    size_t dst_row_size_aligned;
-    size_t theta_cache_offset;
     uint32_t src0_nrows;
 
     struct fastdiv_values div_ne2_ne1;
     struct fastdiv_values div_ne1;
-
-    uint64_t t_start;
 };
 
 static float rope_yarn_ramp(const float low, const float high, const int i0) {
@@ -112,94 +110,80 @@ static inline void rope_yarn_one(float theta, float freq_scale, float * corr_dim
         mscale_final  *= 1.0f + 0.1f * logf(1.0f / freq_scale);
     }
 
-    cache[i0 + 0] = cosf(theta_final) * mscale_final;
-    cache[i0 + 1] = sinf(theta_final) * mscale_final;
+    const uint32_t b = i0 / 64;
+    const uint32_t k = (i0 % 64) / 2;
+    cache[b * 64 + k]      = cosf(theta_final) * mscale_final;
+    cache[b * 64 + 32 + k] = sinf(theta_final) * mscale_final;
+}
+
+// 32 thetas -> 32 deinterleaved pairs [cos[32] | sin[32]] at cache[i0].
+static inline void rope_cache_hvx_32(float * cache, uint32_t i0,
+                                     HVX_Vector v_theta,
+                                     const float * freq_factors,
+                                     HVX_Vector v_freq_scale,
+                                     HVX_Vector v_mscale) {
+    if (freq_factors) {
+        HVX_Vector v_ff = hvx_vmemu(freq_factors + i0 / 2);
+        v_theta = hvx_vec_mul_f32_f32(v_theta, hvx_vec_inverse_f32(v_ff));
+    }
+
+    HVX_Vector v_theta_final = hvx_vec_mul_f32_f32(v_theta, v_freq_scale);
+    HVX_Vector vcos;
+    HVX_Vector vsin;
+    hvx_vec_sincos_f32(v_theta_final, &vcos, &vsin);
+    vcos = hvx_vec_mul_f32_f32(vcos, v_mscale);
+    vsin = hvx_vec_mul_f32_f32(vsin, v_mscale);
+
+    if (((uintptr_t) (cache + i0)) % 128 == 0) {
+        hvx_vmem(cache + i0 + 0)  = vcos;
+        hvx_vmem(cache + i0 + 32) = vsin;
+    } else {
+        hvx_vec_store_u(cache + i0 + 0,  32 * sizeof(float), vcos);
+        hvx_vec_store_u(cache + i0 + 32, 32 * sizeof(float), vsin);
+    }
 }
 
 static __attribute__((noinline)) void rope_cache_init(const float    theta_base,
                             const float    freq_scale,
                             const float *  freq_factors,
                             float *        corr_dims,
-                            const uint32_t ne0,
+                            const uint32_t n_cache,
                             const float    ext_factor,
                             const float    mscale,
                             float *        cache,
-                            const float    theta_scale) {
+                            const float    theta_scale,
+                            const float *  theta_powers,
+                            const float    theta_scale_32) {
     // ref: https://github.com/jquesnelle/yarn/blob/master/scaled_rope/LlamaYaRNScaledRotaryEmbedding.py
-#if __HVX_ARCH__ >= 79
-    const bool is_v79_or_newer = true;
-#else
-    const bool is_v79_or_newer = false;
-#endif
-
-    if (is_v79_or_newer && ext_factor == 0.0f) {
+    if (ext_factor == 0.0f) {
         // Fast path: fully vectorized
         // We process 32 pairs (64 elements) per iteration.
-        const uint32_t n_blocks = ne0 / 64;
-
-        // Initialize theta scale powers: [1.0f, theta_scale, theta_scale^2, ..., theta_scale^31]
-        float __attribute__((aligned(128))) theta_powers[32];
-        theta_powers[0] = 1.0f;
-        for (int j = 1; j < 32; j++) {
-            theta_powers[j] = theta_powers[j - 1] * theta_scale;
-        }
-        HVX_Vector v_theta_powers = hvx_vmem(theta_powers);
+        const uint32_t n_blocks = n_cache / 64;
 
+        HVX_Vector v_theta_powers = hvx_vmemu(theta_powers);
         HVX_Vector v_freq_scale = hvx_vec_splat_f32(freq_scale);
         HVX_Vector v_mscale = hvx_vec_splat_f32(mscale);
 
-        // Base theta starts at theta_base
         float theta_block = theta_base;
-        // The scale factor for the next block is theta_scale^32
-        float theta_scale_32 = 1.0f;
-        for (int j = 0; j < 32; j++) {
-            theta_scale_32 *= theta_scale;
-        }
 
         for (uint32_t b = 0; b < n_blocks; b++) {
             uint32_t i0 = b * 64;
             HVX_Vector v_theta_base = hvx_vec_splat_f32(theta_block);
             HVX_Vector v_theta = hvx_vec_mul_f32_f32(v_theta_base, v_theta_powers);
-
-            if (freq_factors) {
-                // Load 32 elements of freq_factors
-                HVX_Vector v_ff = hvx_vmemu(freq_factors + i0 / 2);
-                HVX_Vector v_inv_ff = hvx_vec_inverse_f32(v_ff);
-                v_theta = hvx_vec_mul_f32_f32(v_theta, v_inv_ff);
-            }
-
-            HVX_Vector v_theta_final = hvx_vec_mul_f32_f32(v_theta, v_freq_scale);
-
-            HVX_Vector vcos = hvx_vec_cos_f32(v_theta_final);
-            HVX_Vector vsin = hvx_vec_sin_f32(v_theta_final);
-
-            vcos = hvx_vec_mul_f32_f32(vcos, v_mscale);
-            vsin = hvx_vec_mul_f32_f32(vsin, v_mscale);
-
-            HVX_VectorPair vstore = Q6_W_vshuff_VVR(vsin, vcos, -4);
-
-            if (((uintptr_t)cache) % 128 == 0) {
-                hvx_vmem(cache + i0 + 0)  = Q6_V_lo_W(vstore);
-                hvx_vmem(cache + i0 + 32) = Q6_V_hi_W(vstore);
-            } else {
-                hvx_vec_store_u(cache + i0 + 0,  32 * sizeof(float), Q6_V_lo_W(vstore));
-                hvx_vec_store_u(cache + i0 + 32, 32 * sizeof(float), Q6_V_hi_W(vstore));
-            }
-
+            rope_cache_hvx_32(cache, i0, v_theta, freq_factors, v_freq_scale, v_mscale);
             theta_block *= theta_scale_32;
         }
 
         // Leftovers
         float theta = theta_block;
-        for (uint32_t i0 = n_blocks * 64; i0 < ne0; i0 += 2) {
+        for (uint32_t i0 = n_blocks * 64; i0 < n_cache; i0 += 2) {
             const float ff = freq_factors ? freq_factors[i0 / 2] : 1.0f;
             rope_yarn_one(theta / ff, freq_scale, corr_dims, i0, ext_factor, mscale, cache);
             theta *= theta_scale;
         }
     } else {
-        // Fallback to original scalar loop
         float theta = theta_base;
-        for (uint32_t i0 = 0; i0 < ne0; i0 += 2) {
+        for (uint32_t i0 = 0; i0 < n_cache; i0 += 2) {
             const float ff = freq_factors ? freq_factors[i0 / 2] : 1.0f;
             rope_yarn_one(theta / ff, freq_scale, corr_dims, i0, ext_factor, mscale, cache);
             theta *= theta_scale;
@@ -207,6 +191,72 @@ static __attribute__((noinline)) void rope_cache_init(const float    theta_base,
     }
 }
 
+static inline float mrope_pick_theta(float theta_t, float theta_h, float theta_w, float theta_e,
+                                     int sector, const int32_t sections[4], int sec_w, int sec_e,
+                                     bool is_imrope) {
+    if (is_imrope) {
+        if      (sector % 3 == 0 && sector < 3 * sections[0]) { return theta_t; }
+        else if (sector % 3 == 1 && sector < 3 * sections[1]) { return theta_h; }
+        else if (sector % 3 == 2 && sector < 3 * sections[2]) { return theta_w; }
+        else                                                   { return theta_e; }
+    }
+    if      (sector < sections[0]) { return theta_t; }
+    else if (sector < sec_w)       { return theta_h; }
+    else if (sector < sec_e)       { return theta_w; }
+    else                           { return theta_e; }
+}
+
+// lane j is 1 when (j % 3) == rem
+static const float __attribute__((aligned(128))) mrope_mod3_eq0[32] = {
+    1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0
+};
+static const float __attribute__((aligned(128))) mrope_mod3_eq1[32] = {
+    0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1
+};
+static const float __attribute__((aligned(128))) mrope_mod3_eq2[32] = {
+    0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0,1,0,0
+};
+
+static const float __attribute__((aligned(128))) mrope_k_ramp[32] = {
+    0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,
+    16,17,18,19,20,21,22,23,24,25,26,27,28,29,30,31
+};
+
+static inline HVX_VectorPred mrope_mask_eq1(const float * m) {
+    return Q6_Q_vcmp_gt_VsfVsf(hvx_vmemu(m), Q6_V_vzero());
+}
+
+// IMROPE without wrap: theta[k] = pos[k % 3] * scale^k
+static inline HVX_Vector mrope_thetas_imrope_mod3(float pos_t, float pos_h, float pos_w,
+                                                  uint32_t k0, HVX_Vector v_powers, float scale_block) {
+    const int r = (int) (k0 % 3);
+    const float * mt = (r == 0) ? mrope_mod3_eq0 : (r == 1) ? mrope_mod3_eq2 : mrope_mod3_eq1;
+    const float * mh = (r == 0) ? mrope_mod3_eq1 : (r == 1) ? mrope_mod3_eq0 : mrope_mod3_eq2;
+
+    HVX_Vector v = hvx_vec_splat_f32(pos_w);
+    v = Q6_V_vmux_QVV(mrope_mask_eq1(mh), hvx_vec_splat_f32(pos_h), v);
+    v = Q6_V_vmux_QVV(mrope_mask_eq1(mt), hvx_vec_splat_f32(pos_t), v);
+    v = hvx_vec_mul_f32_f32(v, v_powers);
+    return hvx_vec_mul_f32_f32(v, hvx_vec_splat_f32(scale_block));
+}
+
+// Contiguous MROPE without wrap: theta[k] = pos[section(k)] * scale^k
+static inline HVX_Vector mrope_thetas_contig(float pos_t, float pos_h, float pos_w, float pos_e,
+                                             uint32_t k0, int s0, int sec_w, int sec_e,
+                                             HVX_Vector v_powers, float scale_block) {
+    HVX_Vector v_k = hvx_vec_add_f32_f32(hvx_vec_splat_f32((float) k0), hvx_vmemu(mrope_k_ramp));
+    HVX_VectorPred lt_s0 = Q6_Q_vcmp_gt_VsfVsf(hvx_vec_splat_f32((float) s0), v_k);
+    HVX_VectorPred lt_sw = Q6_Q_vcmp_gt_VsfVsf(hvx_vec_splat_f32((float) sec_w), v_k);
+    HVX_VectorPred lt_se = Q6_Q_vcmp_gt_VsfVsf(hvx_vec_splat_f32((float) sec_e), v_k);
+
+    HVX_Vector v = hvx_vec_splat_f32(pos_e);
+    v = Q6_V_vmux_QVV(lt_se, hvx_vec_splat_f32(pos_w), v);
+    v = Q6_V_vmux_QVV(lt_sw, hvx_vec_splat_f32(pos_h), v);
+    v = Q6_V_vmux_QVV(lt_s0, hvx_vec_splat_f32(pos_t), v);
+    v = hvx_vec_mul_f32_f32(v, v_powers);
+    return hvx_vec_mul_f32_f32(v, hvx_vec_splat_f32(scale_block));
+}
+
 // pos_t/h/w/e: the four position ids for this sequence step (t=time, h=height, w=width, e=extra).
 // sections[4]: number of head dims assigned to each position component.
 static __attribute__((noinline)) void mrope_cache_init(const float    pos_t,
@@ -219,23 +269,71 @@ static __attribute__((noinline)) void mrope_cache_init(const float    pos_t,
                              const float    freq_scale,
                              const float *  freq_factors,
                              float *        corr_dims,
-                             const uint32_t ne0,
+                             const uint32_t n_cache,
                              const float    ext_factor,
                              const float    mscale,
                              float *        cache,
-                             const float    theta_scale) {
+                             const float    theta_scale,
+                             const float *  theta_powers,
+                             const float    theta_scale_32) {
     const int sect_dims = sections[0] + sections[1] + sections[2] + sections[3];
     const int sec_w     = sections[0] + sections[1];
     const int sec_e     = sec_w + sections[2];
+    const uint32_t n_pairs = n_cache / 2;
+
+    const bool no_wrap = (sect_dims > 0) && (n_pairs <= (uint32_t) sect_dims);
+    const bool imrope_mod3 = is_imrope && !indep_sects && no_wrap
+        && sections[0] > 0 && sections[1] > 0 && sections[2] > 0
+        && n_pairs <= (uint32_t) (3 * sections[0])
+        && n_pairs <= (uint32_t) (3 * sections[1])
+        && n_pairs <= (uint32_t) (3 * sections[2]);
+    const bool contig = !is_imrope && !indep_sects && no_wrap;
+
+    if (ext_factor == 0.0f && (imrope_mod3 || contig)) {
+        HVX_Vector v_powers     = hvx_vmemu(theta_powers);
+        HVX_Vector v_freq_scale = hvx_vec_splat_f32(freq_scale);
+        HVX_Vector v_mscale     = hvx_vec_splat_f32(mscale);
+        float scale_block = 1.0f;
+        const uint32_t n_blocks = n_cache / 64;
+
+        for (uint32_t b = 0; b < n_blocks; b++) {
+            const uint32_t i0 = b * 64;
+            const uint32_t k0 = b * 32;
+            HVX_Vector v_theta = imrope_mod3
+                ? mrope_thetas_imrope_mod3(pos_t, pos_h, pos_w, k0, v_powers, scale_block)
+                : mrope_thetas_contig(pos_t, pos_h, pos_w, pos_e, k0, sections[0], sec_w, sec_e,
+                                      v_powers, scale_block);
+            rope_cache_hvx_32(cache, i0, v_theta, freq_factors, v_freq_scale, v_mscale);
+            scale_block *= theta_scale_32;
+        }
+
+        float theta_k = scale_block;
+        for (uint32_t k = n_blocks * 32; k < n_pairs; k++) {
+            const uint32_t i0 = 2 * k;
+            const float pos = mrope_pick_theta(pos_t, pos_h, pos_w, pos_e,
+                                               (int) k, sections, sec_w, sec_e, is_imrope);
+            const float ff = freq_factors ? freq_factors[k] : 1.0f;
+            rope_yarn_one(pos * theta_k / ff, freq_scale, corr_dims, i0, ext_factor, mscale, cache);
+            theta_k *= theta_scale;
+        }
+        return;
+    }
 
     float theta_t = pos_t;
     float theta_h = pos_h;
     float theta_w = pos_w;
     float theta_e = pos_e;
 
-    for (uint32_t i0 = 0; i0 < ne0; i0 += 2) {
-        const float ff     = freq_factors ? freq_factors[i0 / 2] : 1.0f;
-        const int   sector = (i0 / 2) % sect_dims;
+    const bool use_hvx = (ext_factor == 0.0f);
+    float __attribute__((aligned(128))) thetas[32];
+    uint32_t n_thetas = 0;
+    uint32_t block_i0 = 0;
+
+    HVX_Vector v_freq_scale = hvx_vec_splat_f32(freq_scale);
+    HVX_Vector v_mscale     = hvx_vec_splat_f32(mscale);
+
+    for (uint32_t i0 = 0; i0 < n_cache; i0 += 2) {
+        const int sector = (i0 / 2) % sect_dims;
 
         if (indep_sects) {
             // Reset theta when crossing into a new section.
@@ -245,28 +343,34 @@ static __attribute__((noinline)) void mrope_cache_init(const float    pos_t,
             else if (sector == sec_e)       { theta_e = pos_e; }
         }
 
-        float theta;
-        if (is_imrope) {
-            // Interleaved: sector mod 3 selects component
-            if      (sector % 3 == 0 && sector < 3 * sections[0]) { theta = theta_t; }
-            else if (sector % 3 == 1 && sector < 3 * sections[1]) { theta = theta_h; }
-            else if (sector % 3 == 2 && sector < 3 * sections[2]) { theta = theta_w; }
-            else                                                   { theta = theta_e; }
-        } else {
-            // Contiguous sections
-            if      (sector < sections[0]) { theta = theta_t; }
-            else if (sector < sec_w)       { theta = theta_h; }
-            else if (sector < sec_e)       { theta = theta_w; }
-            else                           { theta = theta_e; }
-        }
+        const float theta = mrope_pick_theta(theta_t, theta_h, theta_w, theta_e,
+                                             sector, sections, sec_w, sec_e, is_imrope);
 
-        rope_yarn_one(theta / ff, freq_scale, corr_dims, i0, ext_factor, mscale, cache);
+        if (use_hvx) {
+            if (n_thetas == 0) {
+                block_i0 = i0;
+            }
+            thetas[n_thetas++] = theta;
+            if (n_thetas == 32) {
+                rope_cache_hvx_32(cache, block_i0, hvx_vmemu(thetas), freq_factors, v_freq_scale, v_mscale);
+                n_thetas = 0;
+            }
+        } else {
+            const float ff = freq_factors ? freq_factors[i0 / 2] : 1.0f;
+            rope_yarn_one(theta / ff, freq_scale, corr_dims, i0, ext_factor, mscale, cache);
+        }
 
         theta_t *= theta_scale;
         theta_h *= theta_scale;
         theta_w *= theta_scale;
         theta_e *= theta_scale;
     }
+
+    for (uint32_t k = 0; k < n_thetas; k++) {
+        const uint32_t i0 = block_i0 + 2 * k;
+        const float ff = freq_factors ? freq_factors[i0 / 2] : 1.0f;
+        rope_yarn_one(thetas[k] / ff, freq_scale, corr_dims, i0, ext_factor, mscale, cache);
+    }
 }
 
 #define M_PI 3.1415926535897932384626433
@@ -283,52 +387,54 @@ static void rope_corr_dims(int     n_dims,
     dims[1]     = MIN(n_dims - 1, end);
 }
 
+static inline void hvx_rope_neox_mul(HVX_Vector v0, HVX_Vector v1, HVX_Vector vcos, HVX_Vector vsin,
+                                     HVX_Vector * o0, HVX_Vector * o1) {
+    HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(v0, vcos);
+    HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(v0, vsin);
+    HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(v1, vcos);
+    HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(v1, vsin);
+    *o0 = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vsub_Vqf32Vqf32(vx0_c, vx1_s));
+    *o1 = Q6_Vsf_equals_Vqf32(Q6_Vqf32_vadd_Vqf32Vqf32(vx0_s, vx1_c));
+}
+
+// theta_cache full 32-pair blocks are deinterleaved [cos | sin].
 static inline void hvx_rope_neox_f32_aa(float * restrict dst, const float * restrict src0, uint32_t ne, const float * restrict theta_cache) {
     const uint32_t he = ne / 2;
     const uint32_t nvec = he / 32;
     const uint32_t nloe = he % 32;
 
+    if (nloe == 0) {
+        const HVX_Vector * vs = (const HVX_Vector *) src0;
+        const HVX_Vector * vt = (const HVX_Vector *) theta_cache;
+        HVX_Vector * vd = (HVX_Vector *) dst;
+        for (uint32_t i = 0; i < nvec; i++) {
+            HVX_Vector o0, o1;
+            hvx_rope_neox_mul(vs[i], vs[nvec + i], vt[i * 2 + 0], vt[i * 2 + 1], &o0, &o1);
+            vd[i] = o0;
+            vd[nvec + i] = o1;
+        }
+        return;
+    }
+
     for (uint32_t i = 0; i < nvec; i++) {
-        HVX_Vector v0 = ((const HVX_Vector *) src0)[i];
-        HVX_Vector v1 = hvx_vmemu(src0 + he + i * 32);
-
-        HVX_Vector v2 = ((const HVX_Vector *) theta_cache)[i * 2 + 0];
-        HVX_Vector v3 = ((const HVX_Vector *) theta_cache)[i * 2 + 1];
-
-        HVX_VectorPair vcos_sin = Q6_W_vdeal_VVR(v3, v2, -4);
-
-        HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(v0, Q6_V_lo_W(vcos_sin));
-        HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(v0, Q6_V_hi_W(vcos_sin));
-        HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(v1, Q6_V_lo_W(vcos_sin));
-        HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(v1, Q6_V_hi_W(vcos_sin));
-
-        HVX_Vector v4 = Q6_Vqf32_vsub_Vqf32Vqf32(vx0_c, vx1_s);
-        HVX_Vector v5 = Q6_Vqf32_vadd_Vqf32Vqf32(vx0_s, vx1_c);
-
-        ((HVX_Vector *) dst)[i] = Q6_Vsf_equals_Vqf32(v4);
-        hvx_vmemu(dst + he + i * 32) = Q6_Vsf_equals_Vqf32(v5);
+        HVX_Vector o0, o1;
+        hvx_rope_neox_mul(((const HVX_Vector *) src0)[i],
+                          hvx_vmemu(src0 + he + i * 32),
+                          ((const HVX_Vector *) theta_cache)[i * 2 + 0],
+                          ((const HVX_Vector *) theta_cache)[i * 2 + 1],
+                          &o0, &o1);
+        ((HVX_Vector *) dst)[i] = o0;
+        hvx_vmemu(dst + he + i * 32) = o1;
     }
 
-    if (nloe > 0) {
-        HVX_Vector v0 = hvx_vmemu(src0 + nvec * 32);
-        HVX_Vector v1 = hvx_vmemu(src0 + he + nvec * 32);
-
-        HVX_Vector v2 = ((const HVX_Vector *) theta_cache)[nvec * 2 + 0];
-        HVX_Vector v3 = ((const HVX_Vector *) theta_cache)[nvec * 2 + 1];
-
-        HVX_VectorPair vcos_sin = Q6_W_vdeal_VVR(v3, v2, -4);
-
-        HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(v0, Q6_V_lo_W(vcos_sin));
-        HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(v0, Q6_V_hi_W(vcos_sin));
-        HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(v1, Q6_V_lo_W(vcos_sin));
-        HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(v1, Q6_V_hi_W(vcos_sin));
-
-        HVX_Vector v4 = Q6_Vqf32_vsub_Vqf32Vqf32(vx0_c, vx1_s);
-        HVX_Vector v5 = Q6_Vqf32_vadd_Vqf32Vqf32(vx0_s, vx1_c);
-
-        hvx_vec_store_u(dst + nvec * 32, nloe * sizeof(float), Q6_Vsf_equals_Vqf32(v4));
-        hvx_vec_store_u(dst + he + nvec * 32, nloe * sizeof(float), Q6_Vsf_equals_Vqf32(v5));
-    }
+    HVX_Vector v0 = hvx_vmemu(src0 + nvec * 32);
+    HVX_Vector v1 = hvx_vmemu(src0 + he + nvec * 32);
+    HVX_Vector vcos = hvx_vmemu(theta_cache + nvec * 64);
+    HVX_Vector vsin = hvx_vmemu(theta_cache + nvec * 64 + 32);
+    HVX_Vector o0, o1;
+    hvx_rope_neox_mul(v0, v1, vcos, vsin, &o0, &o1);
+    hvx_vec_store_u(dst + nvec * 32, nloe * sizeof(float), o0);
+    hvx_vec_store_u(dst + he + nvec * 32, nloe * sizeof(float), o1);
 }
 
 static inline void hvx_rope_f32_aa(float * restrict dst, const float * restrict src0, uint32_t ne, const float * restrict theta_cache) {
@@ -339,16 +445,15 @@ static inline void hvx_rope_f32_aa(float * restrict dst, const float * restrict
         HVX_Vector v0 = ((const HVX_Vector *) src0)[i * 2 + 0];
         HVX_Vector v1 = ((const HVX_Vector *) src0)[i * 2 + 1];
 
-        HVX_Vector v2 = ((const HVX_Vector *) theta_cache)[i * 2 + 0];
-        HVX_Vector v3 = ((const HVX_Vector *) theta_cache)[i * 2 + 1];
+        HVX_Vector vcos = ((const HVX_Vector *) theta_cache)[i * 2 + 0];
+        HVX_Vector vsin = ((const HVX_Vector *) theta_cache)[i * 2 + 1];
 
-        HVX_VectorPair vx0_x1   = Q6_W_vdeal_VVR(v1, v0, -4);
-        HVX_VectorPair vcos_sin = Q6_W_vdeal_VVR(v3, v2, -4);
+        HVX_VectorPair vx0_x1 = Q6_W_vdeal_VVR(v1, v0, -4);
 
-        HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), Q6_V_lo_W(vcos_sin));
-        HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), Q6_V_hi_W(vcos_sin));
-        HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), Q6_V_lo_W(vcos_sin));
-        HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), Q6_V_hi_W(vcos_sin));
+        HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), vcos);
+        HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), vsin);
+        HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), vcos);
+        HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), vsin);
 
         HVX_Vector v4 = Q6_Vqf32_vsub_Vqf32Vqf32(vx0_c, vx1_s);
         HVX_Vector v5 = Q6_Vqf32_vadd_Vqf32Vqf32(vx0_s, vx1_c);
@@ -362,15 +467,15 @@ static inline void hvx_rope_f32_aa(float * restrict dst, const float * restrict
     if (nloe > 0) {
         if (nloe <= 32) {
             HVX_Vector v0 = hvx_vmemu(src0 + nvec * 64);
-            HVX_Vector v2 = hvx_vmemu(theta_cache + nvec * 64);
+            HVX_Vector vcos = hvx_vmemu(theta_cache + nvec * 64);
+            HVX_Vector vsin = hvx_vmemu(theta_cache + nvec * 64 + 32);
 
-            HVX_VectorPair vx0_x1   = Q6_W_vdeal_VVR(Q6_V_vzero(), v0, -4);
-            HVX_VectorPair vcos_sin = Q6_W_vdeal_VVR(Q6_V_vzero(), v2, -4);
+            HVX_VectorPair vx0_x1 = Q6_W_vdeal_VVR(Q6_V_vzero(), v0, -4);
 
-            HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), Q6_V_lo_W(vcos_sin));
-            HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), Q6_V_hi_W(vcos_sin));
-            HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), Q6_V_lo_W(vcos_sin));
-            HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), Q6_V_hi_W(vcos_sin));
+            HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), vcos);
+            HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), vsin);
+            HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), vcos);
+            HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), vsin);
 
             HVX_Vector v4 = Q6_Vqf32_vsub_Vqf32Vqf32(vx0_c, vx1_s);
             HVX_Vector v5 = Q6_Vqf32_vadd_Vqf32Vqf32(vx0_s, vx1_c);
@@ -382,16 +487,15 @@ static inline void hvx_rope_f32_aa(float * restrict dst, const float * restrict
             HVX_Vector v0 = hvx_vmemu(src0 + nvec * 64);
             HVX_Vector v1 = hvx_vmemu(src0 + nvec * 64 + 32);
 
-            HVX_Vector v2 = hvx_vmemu(theta_cache + nvec * 64);
-            HVX_Vector v3 = hvx_vmemu(theta_cache + nvec * 64 + 32);
+            HVX_Vector vcos = hvx_vmemu(theta_cache + nvec * 64);
+            HVX_Vector vsin = hvx_vmemu(theta_cache + nvec * 64 + 32);
 
-            HVX_VectorPair vx0_x1   = Q6_W_vdeal_VVR(v1, v0, -4);
-            HVX_VectorPair vcos_sin = Q6_W_vdeal_VVR(v3, v2, -4);
+            HVX_VectorPair vx0_x1 = Q6_W_vdeal_VVR(v1, v0, -4);
 
-            HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), Q6_V_lo_W(vcos_sin));
-            HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), Q6_V_hi_W(vcos_sin));
-            HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), Q6_V_lo_W(vcos_sin));
-            HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), Q6_V_hi_W(vcos_sin));
+            HVX_Vector vx0_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), vcos);
+            HVX_Vector vx0_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_lo_W(vx0_x1), vsin);
+            HVX_Vector vx1_c = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), vcos);
+            HVX_Vector vx1_s = Q6_Vqf32_vmpy_VsfVsf(Q6_V_hi_W(vx0_x1), vsin);
 
             HVX_Vector v4 = Q6_Vqf32_vsub_Vqf32Vqf32(vx0_c, vx1_s);
             HVX_Vector v5 = Q6_Vqf32_vadd_Vqf32Vqf32(vx0_s, vx1_c);
@@ -404,54 +508,23 @@ static inline void hvx_rope_f32_aa(float * restrict dst, const float * restrict
     }
 }
 
-static void inline rope_basic_f32(struct htp_rope_context * rctx, uint8_t * restrict dst, uint8_t * restrict src,
-                   uint32_t nr, uint32_t ne0, const float * restrict theta_cache) {
-    const uint32_t n_offs = rctx->n_offs; // VLEN-aligned (enforced by supports_op)
+static void inline rope_basic_f32_inplace(struct htp_rope_context * rctx, uint8_t * src,
+                   uint32_t nr, const float * restrict theta_cache) {
+    const uint32_t n_offs = rctx->n_offs;
     #pragma unroll(4)
     for (uint32_t i = 0; i < nr; i++) {
-        float * d = (float *) (dst + i * rctx->dst_row_size_aligned);
         float * s = (float *) (src + i * rctx->src0_row_size_aligned);
-
-        hvx_rope_f32_aa(d + n_offs, s + n_offs, rctx->n_dims, theta_cache);
-
-        // fill the remain channels with data from src tensor
-        if (n_offs > 0) {
-            hvx_copy_f32_uu((uint8_t *) d, (uint8_t *) s, n_offs);
-        }
-        if (n_offs + rctx->n_dims < ne0) {
-            hvx_copy_f32_uu((uint8_t *)(d + n_offs + rctx->n_dims), (uint8_t *)(s + n_offs + rctx->n_dims), ne0 - n_offs - rctx->n_dims);
-        }
+        hvx_rope_f32_aa(s + n_offs, s + n_offs, rctx->n_dims, theta_cache);
     }
 }
 
-static void inline rope_neox_f32(struct htp_rope_context * rctx, uint8_t * restrict dst, uint8_t * restrict src,
-                   uint32_t nr, uint32_t ne0, const float * restrict theta_cache) {
-    const uint32_t n_offs = rctx->n_offs; // VLEN-aligned (enforced by supports_op)
+static void inline rope_neox_f32_inplace(struct htp_rope_context * rctx, uint8_t * src,
+                   uint32_t nr, uint32_t ne, const float * restrict theta_cache) {
+    const uint32_t n_offs = rctx->n_offs;
     #pragma unroll(4)
     for (uint32_t i = 0; i < nr; i++) {
-        float * d = (float *) (dst + i * rctx->dst_row_size_aligned);
         float * s = (float *) (src + i * rctx->src0_row_size_aligned);
-
-        hvx_rope_neox_f32_aa(d + n_offs, s + n_offs, rctx->n_dims, theta_cache);
-
-        // fill the remain channels with data from src tensor
-        if (n_offs > 0) {
-            hvx_copy_f32_uu((uint8_t *) d, (uint8_t *) s, n_offs);
-        }
-        if (n_offs + rctx->n_dims < ne0) {
-            hvx_copy_f32_uu((uint8_t *)(d + n_offs + rctx->n_dims), (uint8_t *)(s + n_offs + rctx->n_dims), ne0 - n_offs - rctx->n_dims);
-        }
-    }
-}
-
-static void inline rope_vision_f32(struct htp_rope_context * rctx, uint8_t * restrict dst, uint8_t * restrict src,
-                   uint32_t nr, uint32_t ne0, const float * restrict theta_cache) {
-    #pragma unroll(4)
-    for (uint32_t i = 0; i < nr; i++) {
-        float * d = (float *) (dst + i * rctx->dst_row_size_aligned);
-        float * s = (float *) (src + i * rctx->src0_row_size_aligned);
-
-        hvx_rope_neox_f32_aa(d, s, ne0, theta_cache);
+        hvx_rope_neox_f32_aa(s + n_offs, s + n_offs, ne, theta_cache);
     }
 }
 
@@ -477,20 +550,18 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
         return;
     }
 
-    uint64_t tt = HAP_perf_get_qtimer_count();
-
     const int32_t mode    = rctx->mode;
     // MROPE, IMROPE and VISION use NEOX-style pairing for the rotation
     const bool    is_neox = (mode & HTP_ROPE_TYPE_NEOX) || (mode & HTP_ROPE_TYPE_MROPE);
     const bool    is_vision = (mode == HTP_ROPE_TYPE_VISION);
 
     // VTCM setup
-    uint8_t * src0_spad_base = octx->src0_spad.data + (ith * octx->src0_spad.size_per_thread);
+    uint8_t * src0_spad_base = rctx->vtcm_base + (ith * rctx->spad_per_thread);
     float *   theta_cache    = (float *) (src0_spad_base);
               src0_spad_base = src0_spad_base + rctx->theta_cache_offset;
-    uint8_t * dst_spad_base  = octx->dst_spad.data + (ith * octx->dst_spad.size_per_thread);
 
     dma_queue * dma_queue = octx->ctx->dma[ith];
+    struct htp_thread_trace * tr = &octx->ctx->trace[ith];
     const int32_t * pos = (const int32_t *) src1->data;
     const float * freq_factors = src2 ? (const float *) src2->data : NULL;
 
@@ -501,6 +572,7 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
 
     uint32_t ir = src0_start_row;
     uint32_t prev_i2 = (uint32_t) -1;
+    uint32_t cur_slot = 0;
 
     for (uint32_t i3 = i3_start; i3 < ne3; i3++) { // batch
         const uint32_t i2_init = (i3 == i3_start) ? i2_start : 0;
@@ -513,35 +585,30 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
                 const uint32_t nrows = MIN(src0_end_row - ir, ne1 - i1);
 
                 // Depth before prefetch
-                uint32_t dma_depth = dma_queue_depth(dma_queue);
+                const uint32_t dma_depth = dma_queue_depth(dma_queue);
 
-                // FARF(HIGH, "rope-block %u: ir %u n-rows %u dma-depth %u : usec %u", ith, ir, nrows, dma_depth,
-                //             (unsigned) HAP_perf_qtimer_count_to_us(HAP_perf_get_qtimer_count() - rctx->t_start));
+                // Prefetch up to 2 blocks
+                const uint32_t p_nrows = MIN(nrows, 2 * HTP_ROPE_SPAD_BLOCK);
+                for (uint32_t pr = 0; pr < p_nrows; pr += HTP_ROPE_SPAD_BLOCK) {
+                    const uint32_t pnr = MIN(nrows - pr, HTP_ROPE_SPAD_BLOCK);
+                    const uint32_t slot = (cur_slot + pr / HTP_ROPE_SPAD_BLOCK) % HTP_ROPE_SPAD_NSLOTS;
+                    uint8_t * spad_slot = rope_spad_slot(src0_spad_base, slot, rctx->src0_row_size_aligned);
+                    const uint8_t * src_addr = (const uint8_t *) src0->data + i3 * nb03 + i2 * nb02 + (i1 + pr) * nb01;
 
-                // Prefetch loop
-                for (uint32_t pnr = 0, pr = 0; pr < nrows && pr < HTP_ROPE_SPAD_NROWS; pr += pnr) {
-                    pnr = MIN(nrows - pr, HTP_ROPE_SPAD_BLOCK);
+                    // Dummy DMA transaction for sequencing (interleaving wr, rd, wr, rd, ...)
+                    dma_queue_push(dma_queue, dma_make_ptr((void *) dst->data, spad_slot), 0, 0, 0, 0);
 
-                    uint32_t pi1 = i1 + pr;
-                    uint32_t pir = ir + pr;
-
-                    // Dummy DMA transaction for sequencing (interleaving dst,src,dst,...)
-                    dma_queue_push_vtcm_to_ddr(dma_queue, dma_make_ptr((void *) dst->data, dst_spad_base + pr * rctx->dst_row_size_aligned), 0, 0, 0);
-
-                    const uint8_t * src_addr = (const uint8_t *) src0->data + i3 * nb03 + i2 * nb02 + pi1 * nb01;
-                          uint8_t * src_spad = src0_spad_base + pr * rctx->src0_row_size_aligned;
-
-                    // Copy only the row payload while striding the DDR source
-                    dma_queue_push(dma_queue, dma_make_ptr(src_spad, src_addr),
+                    dma_queue_push(dma_queue, dma_make_ptr(spad_slot, src_addr),
                         rctx->src0_row_size_aligned, rctx->src0_row_stride, rctx->src0_row_size, pnr);
-
-                    // FARF(HIGH, "rope-prefetch %u: pr %u i1 %u i2 %u i3 %u src-spad %p src-addr %p pnr %u", ith, pir, pi1, i2, i3, src_spad, src_addr, pnr);
                 }
 
                 // Update theta cache
                 if (i2 != prev_i2) {
                     prev_i2 = i2;
 
+                    htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_A_PREP, i2);
+                    // VISION rotates the full row; other modes only rotate n_dims.
+                    const uint32_t n_cache = is_vision ? ne0 : (uint32_t) rctx->n_dims;
                     const bool is_mrope = (rctx->mode & HTP_ROPE_TYPE_MROPE) != 0;
                     if (is_mrope) {
                         // src1 holds four position arrays stacked along ne0:
@@ -554,66 +621,71 @@ static void rope_job_f32(unsigned int nth, unsigned int ith, void * data) {
                             (float) pos[i2 + ne2 * 3],
                             rctx->sections, is_imrope, is_vision,
                             rctx->freq_scale, freq_factors, rctx->corr_dims,
-                            ne0, rctx->ext_factor, rctx->attn_factor,
-                            theta_cache, rctx->theta_scale);
+                            n_cache, rctx->ext_factor, rctx->attn_factor,
+                            theta_cache, rctx->theta_scale, rctx->theta_powers, rctx->theta_scale_32);
                     } else {
                        rope_cache_init(pos[i2], rctx->freq_scale, freq_factors, rctx->corr_dims,
-                                        ne0, rctx->ext_factor, rctx->attn_factor,
-                                        theta_cache, rctx->theta_scale);
+                                        n_cache, rctx->ext_factor, rctx->attn_factor,
+                                        theta_cache, rctx->theta_scale, rctx->theta_powers, rctx->theta_scale_32);
                     }
+                    htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_A_PREP, i2);
                 }
 
                 // Skip output DMA transactions from prev block (if any)
-                // No need to wait for those here since we're explicitly waiting for the latest prefecthes below.
-                for (uint32_t d=0; d < dma_depth; d++) { dma_queue_pop_nowait(dma_queue); }
+                for (uint32_t d = 0; d < dma_depth; d++) { dma_queue_pop_nowait(dma_queue); }
 
                 // Compute loop
-                for (uint32_t cnr = 0, cr = 0; cr < nrows; cr += cnr, ir += cnr, i1 += cnr) {
-                    // Number of rows to compute
+                const uint32_t ne = is_vision ? ne0 : rctx->n_dims;
+                const uint32_t base_i1 = i1;
+                const uint32_t base_ir = ir;
+
+                for (uint32_t cnr = 0, cr = 0; cr < nrows; cr += cnr) {
                     cnr = MIN(nrows - cr, HTP_ROPE_SPAD_BLOCK);
+                    const uint32_t slot = (cur_slot + cr / HTP_ROPE_SPAD_BLOCK) % HTP_ROPE_SPAD_NSLOTS;
+                    const uint32_t cur_ir   = base_ir + cr;
+                    const uint32_t cur_i1   = base_i1 + cr;
 
-                    uint8_t * dst_spad = (uint8_t *) dma_queue_pop(dma_queue).src;
-                    uint8_t * src_spad = (uint8_t *) dma_queue_pop(dma_queue).dst;
+                    dma_queue_pop(dma_queue);
+                    uint8_t * cur_spad = (uint8_t *) dma_queue_pop(dma_queue).dst;
 
-                    // FARF(HIGH, "rope-compute %u: ir %u i1 %u i2 %u i3 %u src-spad %p cnr %u : usec %u", ith, ir, i1, i2, i3, src_spad, cnr,
-                    //         (unsigned) HAP_perf_qtimer_count_to_us(HAP_perf_get_qtimer_count() - rctx->t_start));
-
-                    if (is_vision) {
-                        rope_vision_f32(rctx, dst_spad, src_spad, cnr, ne0, theta_cache);
-                    } else if (is_neox) {
-                        rope_neox_f32(rctx, dst_spad, src_spad, cnr, ne0, theta_cache);
+                    htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_COMP, cur_ir);
+                    if (is_neox || is_vision) {
+                        rope_neox_f32_inplace(rctx, cur_spad, cnr, ne, theta_cache);
                     } else {
-                        rope_basic_f32(rctx, dst_spad, src_spad, cnr, ne0, theta_cache);
+                        rope_basic_f32_inplace(rctx, cur_spad, cnr, theta_cache);
                     }
+                    htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_COMP, cur_ir);
 
-                    uint8_t * dst_addr = (uint8_t *) dst->data + i3 * nb3 + i2 * nb2 + i1 * nb1;
+                    uint8_t * dst_addr = (uint8_t *) dst->data + i3 * nb3 + i2 * nb2 + cur_i1 * nb1;
+                    dma_queue_push(dma_queue, dma_make_ptr(dst_addr, cur_spad),
+                        rctx->dst_row_stride, rctx->src0_row_size_aligned, rctx->dst_row_size, cnr);
 
-                    // Write only the row payload while striding the DDR dst
-                    dma_queue_push(dma_queue, dma_make_ptr(dst_addr, dst_spad),
-                        rctx->dst_row_stride, rctx->dst_row_size_aligned, rctx->dst_row_size, cnr);
+                    // Prefetch 2 blocks ahead into the slot just freed
+                    if ((cr + 2 * HTP_ROPE_SPAD_BLOCK) < nrows) {
+                        const uint32_t p_cr   = cr + 2 * HTP_ROPE_SPAD_BLOCK;
+                        const uint32_t pnr    = MIN(nrows - p_cr, HTP_ROPE_SPAD_BLOCK);
+                        const uint32_t p_slot = (cur_slot + p_cr / HTP_ROPE_SPAD_BLOCK) % HTP_ROPE_SPAD_NSLOTS;
+                        uint8_t * p_spad      = rope_spad_slot(src0_spad_base, p_slot, rctx->src0_row_size_aligned);
+                        const uint8_t * src_addr = (const uint8_t *) src0->data + i3 * nb03 + i2 * nb02 + (base_i1 + p_cr) * nb01;
 
-                    // Prefetch more rows (if any)
-                    if ((cr + HTP_ROPE_SPAD_NROWS) < nrows) {
-                        uint32_t pnr = MIN(nrows - (cr + HTP_ROPE_SPAD_NROWS), HTP_ROPE_SPAD_BLOCK);
-                        uint32_t pi1 = i1 + HTP_ROPE_SPAD_NROWS;
-                        uint32_t pir = ir + HTP_ROPE_SPAD_NROWS;
-
-                        const uint8_t * src_addr = (const uint8_t *) src0->data + i3 * nb03 + i2 * nb02 + pi1 * nb01;
-                        dma_queue_push(dma_queue, dma_make_ptr(src_spad, src_addr),
+                        dma_queue_push(dma_queue, dma_make_ptr(p_spad, src_addr),
                             rctx->src0_row_size_aligned, rctx->src0_row_stride, rctx->src0_row_size, pnr);
-
-                        // FARF(HIGH, "rope-prefetch %u: pr %u i1 %u i2 %u i3 %u src-spad %p src-addr %p pnr %u", ith, pir, pi1, i2, i3, src_spad, src_addr, pnr);
                     }
                 }
+
+                const uint32_t n_chunks = (nrows + HTP_ROPE_SPAD_BLOCK - 1) / HTP_ROPE_SPAD_BLOCK;
+                cur_slot = (cur_slot + n_chunks) % HTP_ROPE_SPAD_NSLOTS;
+
+                ir += nrows;
+                i1 += nrows;
             }
         }
     }
 
 done:
     dma_queue_flush(dma_queue);
-    tt = HAP_perf_get_qtimer_count() - tt;
 
-    FARF(HIGH, "rope-f32: %d/%d: (%u:%u) usec %u\n", ith, nth, src0_start_row, src0_end_row, (unsigned) HAP_perf_qtimer_count_to_us(tt));
+    FARF(HIGH, "rope-f32: %d/%d: (%u:%u)\n", ith, nth, src0_start_row, src0_end_row);
 }
 
 static int execute_op_rope_f32(struct htp_ops_context * octx) {
@@ -624,8 +696,6 @@ static int execute_op_rope_f32(struct htp_ops_context * octx) {
     const struct htp_tensor * src2 = octx->src[2];
     const struct htp_tensor * dst  = octx->dst;
 
-    const char * op_type = "rope-f32";
-
     switch (octx->op) {
         case HTP_OP_ROPE:
             break;
@@ -635,48 +705,23 @@ static int execute_op_rope_f32(struct htp_ops_context * octx) {
             return HTP_STATUS_NO_SUPPORT;
     }
 
-    const uint32_t ne0 = dst->ne[0];
-    const uint32_t src0_nrows = src0->ne[1] * src0->ne[2] * src0->ne[3];
-    const uint32_t n_threads = MIN(octx->n_threads, src0_nrows);
+    const struct htp_rope_kernel_params * kparams = (const struct htp_rope_kernel_params *) octx->kernel_params;
+    assert(kparams->n_threads > 0);
+    assert(octx->ctx->vtcm_size >= kparams->vtcm_size);
 
+    const uint32_t ne0 = dst->ne[0];
     const size_t src0_row_size   = src0->ne[0] * sizeof(float);
     const size_t src0_row_stride = src0->nb[1];
     const size_t dst_row_size    = dst->ne[0] * sizeof(float);
     const size_t dst_row_stride  = dst->nb[1];
 
-    // Aligned row sizes for VTCM
-    const size_t src0_row_size_aligned    = hex_round_up(src0_row_size, VLEN);
-    const size_t dst_row_size_aligned     = hex_round_up(dst_row_stride, VLEN);
-    const size_t theta_cache_size_aligned = hex_round_up(src0->ne[0] * sizeof(float), 256);
-
-    // Calculate spad sizes per thread
-    size_t src0_spad_per_thread = theta_cache_size_aligned + HTP_ROPE_SPAD_NROWS * src0_row_size_aligned;
-    size_t dst_spad_per_thread  = HTP_ROPE_SPAD_NROWS * dst_row_size_aligned;
-    size_t spad_per_thread = src0_spad_per_thread + dst_spad_per_thread;
-
-    // Check if we fit in VTCM
-    size_t total_vtcm_needed = spad_per_thread * n_threads;
-    if (octx->ctx->vtcm_size < total_vtcm_needed) {
-        FARF(ERROR, "%s : current VTCM reservation %zu is too small, needed %zu\n", op_type, octx->ctx->vtcm_size, total_vtcm_needed);
-        return HTP_STATUS_VTCM_TOO_SMALL;
-    }
-
-    octx->src0_spad.size_per_thread = src0_spad_per_thread;
-    octx->dst_spad.size_per_thread  = dst_spad_per_thread;
-    octx->src0_spad.size = n_threads * src0_spad_per_thread;
-    octx->dst_spad.size  = n_threads * dst_spad_per_thread;
-    octx->src1_spad.size = 0;
-
-    octx->src0_spad.data = octx->ctx->vtcm_base;                        octx->src0_spad.src = NULL;
-    octx->src1_spad.data = NULL;                                        octx->src1_spad.src = NULL;
-    octx->dst_spad.data  = octx->src0_spad.data + octx->src0_spad.size; octx->dst_spad.src  = NULL;
-
     struct htp_rope_context rctx;
     memset(&rctx, 0, sizeof(struct htp_rope_context));
 
-    rctx.t_start = HAP_perf_get_qtimer_count();
-
-    rctx.octx = octx;
+    rctx.octx                  = octx;
+    rctx.vtcm_base             = (uint8_t *) octx->ctx->vtcm_base;
+    rctx.spad_per_thread       = kparams->spad_per_thread;
+    rctx.theta_cache_offset    = kparams->theta_cache_offset;
 
     const int32_t * op_params = &octx->op_params[0];
     rctx.n_dims     = ((const int32_t *) op_params)[1];
@@ -693,31 +738,29 @@ static int execute_op_rope_f32(struct htp_ops_context * octx) {
     memcpy(&rctx.sections,    (int32_t *) op_params + 11, sizeof(int) * 4);
 
     rctx.theta_scale = powf(rctx.freq_base, -2.0f / rctx.n_dims);
+    rctx.theta_powers[0] = 1.0f;
+    for (int j = 1; j < 32; j++) {
+        rctx.theta_powers[j] = rctx.theta_powers[j - 1] * rctx.theta_scale;
+    }
+    rctx.theta_scale_32 = rctx.theta_powers[31] * rctx.theta_scale;
 
     rope_corr_dims(rctx.n_dims, rctx.n_ctx_orig, rctx.freq_base, rctx.beta_fast, rctx.beta_slow, rctx.corr_dims);
 
-    rctx.src0_row_size   = src0_row_size;
-    rctx.src0_row_stride = src0_row_stride;
-    rctx.dst_row_size    = dst_row_size;
-    rctx.dst_row_stride  = dst_row_stride;
-    rctx.src0_row_size_aligned = src0_row_size_aligned;
-    rctx.dst_row_size_aligned  = dst_row_size_aligned;
-    rctx.theta_cache_offset    = theta_cache_size_aligned;
+    rctx.src0_row_size         = src0_row_size;
+    rctx.src0_row_stride       = src0_row_stride;
+    rctx.dst_row_size          = dst_row_size;
+    rctx.dst_row_stride        = dst_row_stride;
+    rctx.src0_row_size_aligned = kparams->src0_row_size_aligned;
 
-    rctx.src0_nrows = src0_nrows;
-    rctx.src0_nrows_per_thread = (src0_nrows + n_threads - 1) / n_threads;
-
-    if (src0_nrows > 0) {
-        rctx.div_ne2_ne1 = init_fastdiv_values(dst->ne[2] * dst->ne[1]);
-        rctx.div_ne1     = init_fastdiv_values(dst->ne[1]);
-    }
+    rctx.src0_nrows            = kparams->src0_nrows;
+    rctx.src0_nrows_per_thread = kparams->src0_nrows_per_thread;
+    rctx.div_ne2_ne1           = kparams->div_ne2_ne1;
+    rctx.div_ne1               = kparams->div_ne1;
 
     FARF(HIGH, "rope-f32 n-rows %u n-dims %d ne0 %u ext-factor %.6f theta-scale %.6f attn-factor %.6f\n", rctx.src0_nrows, rctx.n_dims, ne0,
          rctx.ext_factor, rctx.theta_scale, rctx.attn_factor);
 
-    if (!(octx->flags & HTP_OPFLAGS_SKIP_COMPUTE)) {
-        worker_pool_run_func(octx->ctx->worker_pool, rope_job_f32, &rctx, n_threads);
-    }
+    work_queue_run(octx->ctx->work_queue, rope_job_f32, &rctx, kparams->n_threads);
 
     return err;
 }

ggml/src/ggml-hexagon/htp/rope-ops.h

@@ -0,0 +1,56 @@
+#ifndef HTP_ROPE_OPS_H
+#define HTP_ROPE_OPS_H
+
+#include "hex-common.h"
+#include "hex-fastdiv.h"
+
+#define HTP_ROPE_SPAD_BLOCK  8
+#define HTP_ROPE_SPAD_NSLOTS 4
+#define HTP_ROPE_SPAD_NROWS  (HTP_ROPE_SPAD_BLOCK * HTP_ROPE_SPAD_NSLOTS)
+
+struct htp_rope_kernel_params {
+    uint32_t n_threads;
+    uint32_t src0_nrows;
+    uint32_t src0_nrows_per_thread;
+    uint32_t vtcm_size;
+    uint32_t spad_per_thread;
+    uint32_t theta_cache_offset;
+    uint32_t src0_row_size_aligned;
+
+    struct fastdiv_values div_ne2_ne1;
+    struct fastdiv_values div_ne1;
+};
+
+#if defined(__cplusplus)
+static_assert(sizeof(struct htp_rope_kernel_params) <= 128, "htp_rope_kernel_params is too large for kernel_params blob");
+#else
+_Static_assert(sizeof(struct htp_rope_kernel_params) <= 128, "htp_rope_kernel_params is too large for kernel_params blob");
+#endif
+
+struct htp_rope_vtcm_layout {
+    size_t total_bytes;
+    size_t bytes_per_thread;
+    size_t theta_cache_size_aligned;
+    size_t src0_row_size_aligned;
+};
+
+static inline void htp_rope_vtcm_layout_build(
+    struct htp_rope_vtcm_layout * layout,
+    uint32_t ne00,
+    uint32_t n_threads
+) {
+    const size_t src0_row_size            = ne00 * sizeof(float);
+    const size_t src0_row_size_aligned    = hex_round_up((uint32_t) src0_row_size, 128);
+    const size_t theta_cache_size_aligned = hex_round_up((uint32_t) src0_row_size, 256);
+
+    layout->src0_row_size_aligned    = src0_row_size_aligned;
+    layout->theta_cache_size_aligned = theta_cache_size_aligned;
+    layout->bytes_per_thread         = theta_cache_size_aligned + HTP_ROPE_SPAD_NROWS * src0_row_size_aligned;
+    layout->total_bytes              = layout->bytes_per_thread * n_threads;
+}
+
+static inline uint8_t * rope_spad_slot(uint8_t * base, uint32_t slot, size_t row_size_aligned) {
+    return base + (slot * HTP_ROPE_SPAD_BLOCK) * row_size_aligned;
+}
+
+#endif // HTP_ROPE_OPS_H

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

@@ -260,6 +260,20 @@ float16_t dequantFuncTQ1_0(const in decodeBufTQ1_0 bl, const in uint blockCoords
     return bl.block.d * (float16_t(int(xi)) - float16_t(1.0));
 }
 
+f16vec4 dequantFuncTQ1_0_v(const in decodeBufTQ1_0 bl, const in uint blockCoords[2], const in uint coordInBlock[2])
+{
+    const uint e = coordInBlock[1];
+    f16vec4 v;
+    [[unroll]] for (uint k = 0u; k < 4u; ++k) {
+        const uint ee = e + k;
+        const uint bidx = tq1_0_byte_of(ee);
+        const uint qbyte = uint(bidx < 48u ? bl.block.qs[bidx] : bl.block.qh[bidx - 48u]);
+        const uint xi = tq1_0_trit(qbyte, tq1_0_digit_of(ee));
+        v[k] = bl.block.d * (float16_t(int(xi)) - float16_t(1.0));
+    }
+    return v;
+}
+
 layout(buffer_reference, std430, buffer_reference_align = 2) buffer decodeBufTQ2_0 {
    block_tq2_0 block;
 };
@@ -1054,7 +1068,7 @@ float16_t dequantFuncIQ2_S(const in decodeBufIQ2_S bl, const in uint blockCoords
     const uint scale = (bl.block.scales[ib32] >> ((idx & 0x10) >> 2)) & 0xf;
     const uint qs = bl.block.qs[ib8];
     const uint qh = bl.block.qh[ib32];
-    const uint sign = bl.block.qs[QUANT_K / 8 + ib8] >> (idx & 0x6);
+    const uint sign = bl.block.qs[QUANT_K_IQ2_S / 8 + ib8] >> (idx & 0x6);
 
     const float d = float(bl.block.d);
     const float db = d * 0.25 * (0.5 + scale);
@@ -1076,7 +1090,7 @@ f16vec4 dequantFuncIQ2_S_v(const in decodeBufIQ2_S bl, const in uint blockCoords
     const uint scale = (bl.block.scales[ib32] >> ((idx & 0x10) >> 2)) & 0xf;
     const uint qs    = bl.block.qs[ib8];
     const uint qh    = bl.block.qh[ib32];
-    const uint sb    = uint(bl.block.qs[QUANT_K / 8 + ib8]) >> (idx & 0x6u);
+    const uint sb    = uint(bl.block.qs[QUANT_K_IQ2_S / 8 + ib8]) >> (idx & 0x6u);
 
     const float d  = float(bl.block.d);
     const float db = d * 0.25 * (0.5 + scale);
@@ -1107,7 +1121,7 @@ float16_t dequantFuncIQ3_XXS(const in decodeBufIQ3_XXS bl, const in uint blockCo
     uint idx = coordInBlock[1];
 
     const uint iqs = (idx & 0xFC) >> 2;             // 0..63
-    const uint is = QUANT_K / 4 + ((idx & 0xE0) >> 3);// 8 values
+    const uint is = QUANT_K_IQ3_XXS / 4 + ((idx & 0xE0) >> 3);// 8 values
 
     const float d = float(bl.block.d);
     const uint qs = bl.block.qs[iqs];
@@ -1130,7 +1144,7 @@ f16vec4 dequantFuncIQ3_XXS_v(const in decodeBufIQ3_XXS bl, const in uint blockCo
     const uint idx = coordInBlock[1];
 
     const uint iqs = idx >> 2;
-    const uint is  = QUANT_K / 4 + ((idx & 0xE0) >> 3);
+    const uint is  = QUANT_K_IQ3_XXS / 4 + ((idx & 0xE0) >> 3);
 
     const float d     = float(bl.block.d);
     const uint  qs    = bl.block.qs[iqs];

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

@@ -1,32 +1,22 @@
 #if !defined(GGML_FA_TYPES_COMP)
 #define GGML_FA_TYPES_COMP
 
-// FaTypeK / FaTypeV spec constant values. These mirror enum ggml_type so the
-// host can pass the type directly. Keep in sync with ggml.h.
-#define FA_TYPE_F32   0u
-#define FA_TYPE_F16   1u
-#define FA_TYPE_Q4_0  2u
-#define FA_TYPE_Q4_1  3u
-#define FA_TYPE_Q5_0  6u
-#define FA_TYPE_Q5_1  7u
-#define FA_TYPE_Q8_0  8u
-#define FA_TYPE_IQ4_NL 20u
-#define FA_TYPE_BF16 30u
+#include "ggml_type_ids.glsl"
 
 // Number of matrix elements per buffer block, derived from the K/V type spec
 // constant. F32 is treated as a vec4 "block" of 4 floats. F16 uses block size 1
 // and bypasses the dequant path entirely. Quants follow their ggml block sizes.
 uint fa_block_elems(uint ty) {
     switch (ty) {
-        case FA_TYPE_F32:  return 4u;
-        case FA_TYPE_F16:  return 1u;
-        case FA_TYPE_Q4_0: return uint(QUANT_K_Q4_0);
-        case FA_TYPE_Q4_1: return uint(QUANT_K_Q4_1);
-        case FA_TYPE_Q5_0: return uint(QUANT_K_Q5_0);
-        case FA_TYPE_Q5_1: return uint(QUANT_K_Q5_1);
-        case FA_TYPE_Q8_0: return uint(QUANT_K_Q8_0);
-        case FA_TYPE_IQ4_NL: return uint(QUANT_K_IQ4_NL);
-        case FA_TYPE_BF16: return 1u;
+        case GGML_TYPE_F32:  return 4u;
+        case GGML_TYPE_F16:  return 1u;
+        case GGML_TYPE_Q4_0: return uint(QUANT_K_Q4_0);
+        case GGML_TYPE_Q4_1: return uint(QUANT_K_Q4_1);
+        case GGML_TYPE_Q5_0: return uint(QUANT_K_Q5_0);
+        case GGML_TYPE_Q5_1: return uint(QUANT_K_Q5_1);
+        case GGML_TYPE_Q8_0: return uint(QUANT_K_Q8_0);
+        case GGML_TYPE_IQ4_NL: return uint(QUANT_K_IQ4_NL);
+        case GGML_TYPE_BF16: return 1u;
         default:           return 1u;
     }
 }
@@ -36,18 +26,18 @@ uint fa_block_elems(uint ty) {
 // of int32s per 32-element block on the MMQ K path: ints_per_block == 8 / R.
 uint fa_quant_r_mmq(uint ty) {
     switch (ty) {
-        case FA_TYPE_Q4_0: return uint(QUANT_R_Q4_0);
-        case FA_TYPE_Q4_1: return uint(QUANT_R_Q4_1);
-        case FA_TYPE_Q5_0: return uint(QUANT_R_Q5_0);
-        case FA_TYPE_Q5_1: return uint(QUANT_R_Q5_1);
-        case FA_TYPE_Q8_0: return uint(QUANT_R_Q8_0);
+        case GGML_TYPE_Q4_0: return uint(QUANT_R_Q4_0);
+        case GGML_TYPE_Q4_1: return uint(QUANT_R_Q4_1);
+        case GGML_TYPE_Q5_0: return uint(QUANT_R_Q5_0);
+        case GGML_TYPE_Q5_1: return uint(QUANT_R_Q5_1);
+        case GGML_TYPE_Q8_0: return uint(QUANT_R_Q8_0);
         default:           return 1u;
     }
 }
 
 bool fa_type_needs_shmem(uint ty) {
     switch (ty) {
-        case FA_TYPE_IQ4_NL: return true;
+        case GGML_TYPE_IQ4_NL: return true;
         default:             return false;
     }
 }

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

@@ -8,7 +8,9 @@ layout(local_size_x = 512, local_size_y = 1, local_size_z = 1) in;
 layout (binding = 0) writeonly buffer D {D_TYPE data_d[];};
 
 void main() {
-    const uint i = gl_GlobalInvocationID.x;
+    // 2D grid flattening: each x workgroup covers gl_WorkGroupSize.x elements,
+    // each y workgroup covers gl_NumWorkGroups.x * gl_WorkGroupSize.x elements.
+    const uint i = (gl_GlobalInvocationID.y * gl_NumWorkGroups.x * gl_WorkGroupSize.x) + gl_GlobalInvocationID.x;
 
     if (i >= p.KX) {
         return;

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

@@ -134,7 +134,7 @@ void main() {
 
         // Q8_0 K only needs (qd, _); the asymmetric Q4_*/Q5_* family also stores
         // the row-sum scaled by qd, used in k_dot_correction.
-        if (FaTypeK == FA_TYPE_Q8_0) {
+        if (FaTypeK == GGML_TYPE_Q8_0) {
             if (buf_iqs == 0) {
                 Qf[buf_ib].ds = FLOAT_TYPEV2(qd, 0.0f);
             }
@@ -367,7 +367,7 @@ void main() {
                 // Q4_*/Q5_* take the block-8 fast path when one step covers a full
                 // block; Q8_0 always goes through the per-int get_k_qs* helpers
                 // (its qs is byte-packed, not nibble-packed).
-                const bool block8_fast = (d_per_step == 8) && (FaTypeK != FA_TYPE_Q8_0);
+                const bool block8_fast = (d_per_step == 8) && (FaTypeK != GGML_TYPE_Q8_0);
 
                 if (SHMEM_STAGING != 0) {
                     const uint k_block_idx = (d_tid * (HSK_per_thread / 4) + d_block) / 8;
@@ -375,7 +375,7 @@ void main() {
                     k_dm = ACC_TYPEV2(kblocksh[buf_ib].dm);
 
                     if (block8_fast) {
-                        const bool has_qh = (FaTypeK == FA_TYPE_Q5_0) || (FaTypeK == FA_TYPE_Q5_1);
+                        const bool has_qh = (FaTypeK == GGML_TYPE_Q5_0) || (FaTypeK == GGML_TYPE_Q5_1);
                         [[unroll]] for (uint32_t d = 0; d < 4; d++) {
                             uint vui = kblocksh[buf_ib].qs[d];
                             k_quants[d    ] = int32_t( vui       & 0x0F0F0F0F);

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

@@ -105,8 +105,8 @@ layout (binding = 6) readonly buffer MO {uint32_t data_mask_opt[];};
 #define BLOCK_SIZE_V fa_block_elems(FaTypeV)
 // F16 reads f16 elements directly from the binding; everything else routes
 // through dequantize4 / the MMQ helpers to unpack from the packed block layout.
-#define USE_DECODE_K (FaTypeK != FA_TYPE_F16)
-#define USE_DECODE_V (FaTypeV != FA_TYPE_F16)
+#define USE_DECODE_K (FaTypeK != GGML_TYPE_F16)
+#define USE_DECODE_V (FaTypeV != GGML_TYPE_F16)
 
 #define CEIL_DIV(a, b) (((a) + (b) - 1) / (b))
 

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

@@ -40,26 +40,28 @@ layout(buffer_reference, std430, buffer_reference_align = 1) buffer decodeBufFA_
 #if !defined(BFLOAT16)
 float16_t faDecodeK(const decodeBufFA_K bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
     switch (FaTypeK) {
-        case FA_TYPE_F32:  return dequantFuncF32 (decodeBufF32 (bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q4_0: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q4_1: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q5_0: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q5_1: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q8_0: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_IQ4_NL: return dequantFuncIQ4_NL(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_F32:  return dequantFuncF32 (decodeBufF32 (bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q4_0: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q4_1: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q5_0: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q5_1: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q8_0: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_IQ4_NL: return dequantFuncIQ4_NL(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q1_0: return dequantFuncQ1_0(decodeBufQ1_0(bl_in), blockCoords, coordInBlock);
         default: return float16_t(0);
     }
 }
 
 float16_t faDecodeV(const decodeBufFA_V bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
     switch (FaTypeV) {
-        case FA_TYPE_F32:  return dequantFuncF32 (decodeBufF32 (bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q4_0: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q4_1: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q5_0: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q5_1: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q8_0: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_IQ4_NL: return dequantFuncIQ4_NL(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_F32:  return dequantFuncF32 (decodeBufF32 (bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q4_0: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q4_1: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q5_0: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q5_1: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q8_0: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_IQ4_NL: return dequantFuncIQ4_NL(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q1_0: return dequantFuncQ1_0(decodeBufQ1_0(bl_in), blockCoords, coordInBlock);
         default: return float16_t(0);
     }
 }
@@ -67,26 +69,26 @@ float16_t faDecodeV(const decodeBufFA_V bl_in, const uint blockCoords[2], const
 // V=4 vector decode for K/V; dispatches to per-format _v decoders.
 f16vec4 faDecodeKVector(const decodeBufFA_K bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
     switch (FaTypeK) {
-        case FA_TYPE_F32: return f16vec4(decodeBufF32(bl_in).block);
-        case FA_TYPE_Q4_0: return dequantFuncQ4_0_v(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q4_1: return dequantFuncQ4_1_v(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q5_0: return dequantFuncQ5_0_v(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q5_1: return dequantFuncQ5_1_v(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q8_0: return dequantFuncQ8_0_v(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_IQ4_NL: return dequantFuncIQ4_NL_v(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_F32: return f16vec4(decodeBufF32(bl_in).block);
+        case GGML_TYPE_Q4_0: return dequantFuncQ4_0_v(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q4_1: return dequantFuncQ4_1_v(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q5_0: return dequantFuncQ5_0_v(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q5_1: return dequantFuncQ5_1_v(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q8_0: return dequantFuncQ8_0_v(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_IQ4_NL: return dequantFuncIQ4_NL_v(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
         default: return f16vec4(0);
     }
 }
 
 f16vec4 faDecodeVVector(const decodeBufFA_V bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
     switch (FaTypeV) {
-        case FA_TYPE_F32: return f16vec4(decodeBufF32(bl_in).block);
-        case FA_TYPE_Q4_0: return dequantFuncQ4_0_v(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q4_1: return dequantFuncQ4_1_v(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q5_0: return dequantFuncQ5_0_v(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q5_1: return dequantFuncQ5_1_v(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_Q8_0: return dequantFuncQ8_0_v(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
-        case FA_TYPE_IQ4_NL: return dequantFuncIQ4_NL_v(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_F32: return f16vec4(decodeBufF32(bl_in).block);
+        case GGML_TYPE_Q4_0: return dequantFuncQ4_0_v(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q4_1: return dequantFuncQ4_1_v(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q5_0: return dequantFuncQ5_0_v(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q5_1: return dequantFuncQ5_1_v(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_Q8_0: return dequantFuncQ8_0_v(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
+        case GGML_TYPE_IQ4_NL: return dequantFuncIQ4_NL_v(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
         default: return f16vec4(0);
     }
 }

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

@@ -121,25 +121,25 @@ layout (binding = 1) readonly buffer K_PACKED_Q5_1_P32 { block_q5_1_packed32 dat
 FLOAT_TYPEV4 dequantize4(uint ib, uint iqs, uint a_offset, uint binding_idx) {
     if (binding_idx == BINDING_IDX_K) {
         switch (FaTypeK) {
-            case FA_TYPE_F32:  FA_DEQUANT4_F32 (k_packed_f32)
-            case FA_TYPE_Q4_0: FA_DEQUANT4_Q4_0(k_packed_q4_0)
-            case FA_TYPE_Q4_1: FA_DEQUANT4_Q4_1(k_packed_q4_1)
-            case FA_TYPE_Q5_0: FA_DEQUANT4_Q5_0(k_packed_q5_0)
-            case FA_TYPE_Q5_1: FA_DEQUANT4_Q5_1(k_packed_q5_1)
-            case FA_TYPE_Q8_0: FA_DEQUANT4_Q8_0(k_packed_q8_0)
-            case FA_TYPE_IQ4_NL: FA_DEQUANT4_IQ4_NL(k_packed_iq4_nl)
-            case FA_TYPE_BF16: FA_DEQUANT4_BF16(k_packed_bf16)
+            case GGML_TYPE_F32:  FA_DEQUANT4_F32 (k_packed_f32)
+            case GGML_TYPE_Q4_0: FA_DEQUANT4_Q4_0(k_packed_q4_0)
+            case GGML_TYPE_Q4_1: FA_DEQUANT4_Q4_1(k_packed_q4_1)
+            case GGML_TYPE_Q5_0: FA_DEQUANT4_Q5_0(k_packed_q5_0)
+            case GGML_TYPE_Q5_1: FA_DEQUANT4_Q5_1(k_packed_q5_1)
+            case GGML_TYPE_Q8_0: FA_DEQUANT4_Q8_0(k_packed_q8_0)
+            case GGML_TYPE_IQ4_NL: FA_DEQUANT4_IQ4_NL(k_packed_iq4_nl)
+            case GGML_TYPE_BF16: FA_DEQUANT4_BF16(k_packed_bf16)
         }
     } else {
         switch (FaTypeV) {
-            case FA_TYPE_F32:  FA_DEQUANT4_F32 (v_packed_f32)
-            case FA_TYPE_Q4_0: FA_DEQUANT4_Q4_0(v_packed_q4_0)
-            case FA_TYPE_Q4_1: FA_DEQUANT4_Q4_1(v_packed_q4_1)
-            case FA_TYPE_Q5_0: FA_DEQUANT4_Q5_0(v_packed_q5_0)
-            case FA_TYPE_Q5_1: FA_DEQUANT4_Q5_1(v_packed_q5_1)
-            case FA_TYPE_Q8_0: FA_DEQUANT4_Q8_0(v_packed_q8_0)
-            case FA_TYPE_IQ4_NL: FA_DEQUANT4_IQ4_NL(v_packed_iq4_nl)
-            case FA_TYPE_BF16: FA_DEQUANT4_BF16(v_packed_bf16)
+            case GGML_TYPE_F32:  FA_DEQUANT4_F32 (v_packed_f32)
+            case GGML_TYPE_Q4_0: FA_DEQUANT4_Q4_0(v_packed_q4_0)
+            case GGML_TYPE_Q4_1: FA_DEQUANT4_Q4_1(v_packed_q4_1)
+            case GGML_TYPE_Q5_0: FA_DEQUANT4_Q5_0(v_packed_q5_0)
+            case GGML_TYPE_Q5_1: FA_DEQUANT4_Q5_1(v_packed_q5_1)
+            case GGML_TYPE_Q8_0: FA_DEQUANT4_Q8_0(v_packed_q8_0)
+            case GGML_TYPE_IQ4_NL: FA_DEQUANT4_IQ4_NL(v_packed_iq4_nl)
+            case GGML_TYPE_BF16: FA_DEQUANT4_BF16(v_packed_bf16)
         }
     }
     return FLOAT_TYPEV4(0);

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

@@ -4,20 +4,20 @@
 
 int32_t get_k_qs(uint ib, uint iqs, uint a_offset) {
     switch (FaTypeK) {
-        case FA_TYPE_Q4_0: {
+        case GGML_TYPE_Q4_0: {
             uint vui = pack32(u16vec2(k_packed_q4_0.data[a_offset + ib].qs[(iqs & 0xF) / 2 + 0],
                                       k_packed_q4_0.data[a_offset + ib].qs[(iqs & 0xF) / 2 + 1]));
             uint shift = (iqs & 0x10) >> 2;
             vui >>= shift;
             return int32_t(vui & 0x0F0F0F0F);
         }
-        case FA_TYPE_Q4_1: { // uses packed32 alias
+        case GGML_TYPE_Q4_1: { // uses packed32 alias
             uint vui = k_packed_q4_1_p32.data[a_offset + ib].qs[(iqs & 0xF) / 4];
             uint shift = (iqs & 0x10) >> 2;
             vui >>= shift;
             return int32_t(vui & 0x0F0F0F0F);
         }
-        case FA_TYPE_Q5_0: {
+        case GGML_TYPE_Q5_0: {
             uint vui = pack32(u16vec2(k_packed_q5_0.data[a_offset + ib].qs[(iqs & 0xF) / 2 + 0],
                                       k_packed_q5_0.data[a_offset + ib].qs[(iqs & 0xF) / 2 + 1]));
             uint qh = pack32(u16vec2(k_packed_q5_0.data[a_offset + ib].qh[0],
@@ -27,7 +27,7 @@ int32_t get_k_qs(uint ib, uint iqs, uint a_offset) {
             uint qh_bits = (qh >> iqs) & 0xF;
             return int32_t(vui & 0x0F0F0F0F) | int32_t((qh_bits * 0x02040810u) & 0x10101010u);
         }
-        case FA_TYPE_Q5_1: { // qs via packed32, qh via packed16
+        case GGML_TYPE_Q5_1: { // qs via packed32, qh via packed16
             uint vui = k_packed_q5_1_p32.data[a_offset + ib].qs[(iqs & 0xF) / 4];
             uint qh  = k_packed_q5_1.data[a_offset + ib].qh;
             uint shift = (iqs & 0x10) >> 2;
@@ -35,7 +35,7 @@ int32_t get_k_qs(uint ib, uint iqs, uint a_offset) {
             uint qh_bits = (qh >> iqs) & 0xF;
             return int32_t(vui & 0x0F0F0F0F) | int32_t((qh_bits * 0x02040810u) & 0x10101010u);
         }
-        case FA_TYPE_Q8_0: {
+        case GGML_TYPE_Q8_0: {
             return pack32(i16vec2(k_packed_q8_0.data[a_offset + ib].qs[iqs / 2],
                                   k_packed_q8_0.data[a_offset + ib].qs[iqs / 2 + 1]));
         }
@@ -47,11 +47,11 @@ int32_t get_k_qs(uint ib, uint iqs, uint a_offset) {
 // return (d, 0) so call sites always see the same shape.
 FLOAT_TYPEV2 get_k_scale(uint ib, uint a_offset) {
     switch (FaTypeK) {
-        case FA_TYPE_Q4_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q4_0.data[a_offset + ib].d), 0.0);
-        case FA_TYPE_Q4_1: return FLOAT_TYPEV2(k_packed_q4_1_p32.data[a_offset + ib].dm);
-        case FA_TYPE_Q5_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q5_0.data[a_offset + ib].d), 0.0);
-        case FA_TYPE_Q5_1: return FLOAT_TYPEV2(k_packed_q5_1_p32.data[a_offset + ib].dm);
-        case FA_TYPE_Q8_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q8_0.data[a_offset + ib].d), 0.0);
+        case GGML_TYPE_Q4_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q4_0.data[a_offset + ib].d), 0.0);
+        case GGML_TYPE_Q4_1: return FLOAT_TYPEV2(k_packed_q4_1_p32.data[a_offset + ib].dm);
+        case GGML_TYPE_Q5_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q5_0.data[a_offset + ib].d), 0.0);
+        case GGML_TYPE_Q5_1: return FLOAT_TYPEV2(k_packed_q5_1_p32.data[a_offset + ib].dm);
+        case GGML_TYPE_Q8_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q8_0.data[a_offset + ib].d), 0.0);
         default: return FLOAT_TYPEV2(0);
     }
 }
@@ -61,16 +61,16 @@ void k_block_to_shmem(const uint buf_ib, const uint global_ib, const uint iqs, c
     // explicit casts. The bit pattern is what we care about here -- the actual
     // signed/unsigned interpretation happens downstream in the dot product.
     switch (FaTypeK) {
-        case FA_TYPE_Q4_0: {
+        case GGML_TYPE_Q4_0: {
             kblocksh[buf_ib].qs[iqs] = int32_t(pack32(u16vec2(k_packed_q4_0.data[a_offset + global_ib].qs[iqs * 2],
                                                               k_packed_q4_0.data[a_offset + global_ib].qs[iqs * 2 + 1])));
             break;
         }
-        case FA_TYPE_Q4_1: {
+        case GGML_TYPE_Q4_1: {
             kblocksh[buf_ib].qs[iqs] = int32_t(k_packed_q4_1_p32.data[a_offset + global_ib].qs[iqs]);
             break;
         }
-        case FA_TYPE_Q5_0: {
+        case GGML_TYPE_Q5_0: {
             kblocksh[buf_ib].qs[iqs] = int32_t(pack32(u16vec2(k_packed_q5_0.data[a_offset + global_ib].qs[iqs * 2],
                                                               k_packed_q5_0.data[a_offset + global_ib].qs[iqs * 2 + 1])));
             if (iqs == 0) {
@@ -79,14 +79,14 @@ void k_block_to_shmem(const uint buf_ib, const uint global_ib, const uint iqs, c
             }
             break;
         }
-        case FA_TYPE_Q5_1: {
+        case GGML_TYPE_Q5_1: {
             kblocksh[buf_ib].qs[iqs] = int32_t(k_packed_q5_1_p32.data[a_offset + global_ib].qs[iqs]);
             if (iqs == 0) {
                 kblocksh[buf_ib].qh = k_packed_q5_1.data[a_offset + global_ib].qh;
             }
             break;
         }
-        case FA_TYPE_Q8_0: {
+        case GGML_TYPE_Q8_0: {
             kblocksh[buf_ib].qs[iqs] = pack32(i16vec2(k_packed_q8_0.data[a_offset + global_ib].qs[iqs * 2],
                                                       k_packed_q8_0.data[a_offset + global_ib].qs[iqs * 2 + 1]));
             break;
@@ -96,11 +96,11 @@ void k_block_to_shmem(const uint buf_ib, const uint global_ib, const uint iqs, c
     if (iqs == 0) {
         // Q4_0/Q5_0/Q8_0 store dm.x = d; Q4_1/Q5_1 store dm = (d, m) pair.
         switch (FaTypeK) {
-            case FA_TYPE_Q4_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q4_0.data[a_offset + global_ib].d), 0.0); break;
-            case FA_TYPE_Q4_1: kblocksh[buf_ib].dm = FLOAT_TYPEV2(k_packed_q4_1_p32.data[a_offset + global_ib].dm); break;
-            case FA_TYPE_Q5_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q5_0.data[a_offset + global_ib].d), 0.0); break;
-            case FA_TYPE_Q5_1: kblocksh[buf_ib].dm = FLOAT_TYPEV2(k_packed_q5_1_p32.data[a_offset + global_ib].dm); break;
-            case FA_TYPE_Q8_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q8_0.data[a_offset + global_ib].d), 0.0); break;
+            case GGML_TYPE_Q4_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q4_0.data[a_offset + global_ib].d), 0.0); break;
+            case GGML_TYPE_Q4_1: kblocksh[buf_ib].dm = FLOAT_TYPEV2(k_packed_q4_1_p32.data[a_offset + global_ib].dm); break;
+            case GGML_TYPE_Q5_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q5_0.data[a_offset + global_ib].d), 0.0); break;
+            case GGML_TYPE_Q5_1: kblocksh[buf_ib].dm = FLOAT_TYPEV2(k_packed_q5_1_p32.data[a_offset + global_ib].dm); break;
+            case GGML_TYPE_Q8_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q8_0.data[a_offset + global_ib].d), 0.0); break;
         }
     }
 }
@@ -121,31 +121,31 @@ struct fa_k_qs_block8 {
 fa_k_qs_block8 get_k_qs_block8(uint ib, uint a_offset) {
     fa_k_qs_block8 r;
     uint qh = 0;
-    if (FaTypeK == FA_TYPE_Q5_0) {
+    if (FaTypeK == GGML_TYPE_Q5_0) {
         qh = pack32(u16vec2(k_packed_q5_0.data[a_offset + ib].qh[0],
                             k_packed_q5_0.data[a_offset + ib].qh[1]));
-    } else if (FaTypeK == FA_TYPE_Q5_1) {
+    } else if (FaTypeK == GGML_TYPE_Q5_1) {
         qh = k_packed_q5_1.data[a_offset + ib].qh;
     }
-    const bool has_qh = (FaTypeK == FA_TYPE_Q5_0) || (FaTypeK == FA_TYPE_Q5_1);
+    const bool has_qh = (FaTypeK == GGML_TYPE_Q5_0) || (FaTypeK == GGML_TYPE_Q5_1);
     [[unroll]] for (uint32_t d = 0; d < 4; d++) {
         uint vui = 0;
         switch (FaTypeK) {
-            case FA_TYPE_Q4_0: { // packed16
+            case GGML_TYPE_Q4_0: { // packed16
                 vui = pack32(u16vec2(k_packed_q4_0.data[a_offset + ib].qs[d * 2 + 0],
                                      k_packed_q4_0.data[a_offset + ib].qs[d * 2 + 1]));
                 break;
             }
-            case FA_TYPE_Q4_1: { // packed32 alias
+            case GGML_TYPE_Q4_1: { // packed32 alias
                 vui = k_packed_q4_1_p32.data[a_offset + ib].qs[d];
                 break;
             }
-            case FA_TYPE_Q5_0: { // packed16
+            case GGML_TYPE_Q5_0: { // packed16
                 vui = pack32(u16vec2(k_packed_q5_0.data[a_offset + ib].qs[d * 2 + 0],
                                      k_packed_q5_0.data[a_offset + ib].qs[d * 2 + 1]));
                 break;
             }
-            case FA_TYPE_Q5_1: { // packed32 alias
+            case GGML_TYPE_Q5_1: { // packed32 alias
                 vui = k_packed_q5_1_p32.data[a_offset + ib].qs[d];
                 break;
             }
@@ -164,21 +164,21 @@ fa_k_qs_block8 get_k_qs_block8(uint ib, uint a_offset) {
 
 int32_t get_k_qs_shmem(const uint buf_ib, const uint pos) {
     switch (FaTypeK) {
-        case FA_TYPE_Q4_0:
-        case FA_TYPE_Q4_1: {
+        case GGML_TYPE_Q4_0:
+        case GGML_TYPE_Q4_1: {
             uint sub = pos % 4;
             uint shift = ((pos % 8) >= 4) ? 4u : 0u;
             return int32_t((uint(kblocksh[buf_ib].qs[sub]) >> shift) & 0x0F0F0F0Fu);
         }
-        case FA_TYPE_Q5_0:
-        case FA_TYPE_Q5_1: {
+        case GGML_TYPE_Q5_0:
+        case GGML_TYPE_Q5_1: {
             uint sub = pos % 4;
             uint shift = ((pos % 8) >= 4) ? 4u : 0u;
             int32_t result = int32_t((uint(kblocksh[buf_ib].qs[sub]) >> shift) & 0x0F0F0F0Fu);
             uint qh_bits = (kblocksh[buf_ib].qh >> (pos * 4u)) & 0xFu;
             return result | int32_t((qh_bits * 0x02040810u) & 0x10101010u);
         }
-        case FA_TYPE_Q8_0: {
+        case GGML_TYPE_Q8_0: {
             return kblocksh[buf_ib].qs[pos];
         }
         default: return 0;
@@ -187,10 +187,10 @@ int32_t get_k_qs_shmem(const uint buf_ib, const uint pos) {
 
 ACC_TYPE k_dot_correction(const uint qib, const ACC_TYPEV2 k_dm) {
     switch (FaTypeK) {
-        case FA_TYPE_Q4_0: return -ACC_TYPE(8.0)  * ACC_TYPE(Qf[qib].ds.y) * k_dm.x;
-        case FA_TYPE_Q5_0: return -ACC_TYPE(16.0) * ACC_TYPE(Qf[qib].ds.y) * k_dm.x;
-        case FA_TYPE_Q4_1:
-        case FA_TYPE_Q5_1: return ACC_TYPE(Qf[qib].ds.y) * k_dm.y;
+        case GGML_TYPE_Q4_0: return -ACC_TYPE(8.0)  * ACC_TYPE(Qf[qib].ds.y) * k_dm.x;
+        case GGML_TYPE_Q5_0: return -ACC_TYPE(16.0) * ACC_TYPE(Qf[qib].ds.y) * k_dm.x;
+        case GGML_TYPE_Q4_1:
+        case GGML_TYPE_Q5_1: return ACC_TYPE(Qf[qib].ds.y) * k_dm.y;
         default: return ACC_TYPE(0.0);
     }
 }

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

@@ -0,0 +1,34 @@
+#if !defined(GGML_TYPE_IDS_COMP)
+#define GGML_TYPE_IDS_COMP
+
+// ggml_type enum values — must match ggml.h
+#define GGML_TYPE_F32      0u
+#define GGML_TYPE_F16      1u
+#define GGML_TYPE_Q4_0     2u
+#define GGML_TYPE_Q4_1     3u
+#define GGML_TYPE_Q5_0     6u
+#define GGML_TYPE_Q5_1     7u
+#define GGML_TYPE_Q8_0     8u
+#define GGML_TYPE_Q2_K    10u
+#define GGML_TYPE_Q3_K    11u
+#define GGML_TYPE_Q4_K    12u
+#define GGML_TYPE_Q5_K    13u
+#define GGML_TYPE_Q6_K    14u
+#define GGML_TYPE_IQ2_XXS 16u
+#define GGML_TYPE_IQ2_XS  17u
+#define GGML_TYPE_IQ3_XXS 18u
+#define GGML_TYPE_IQ1_S   19u
+#define GGML_TYPE_IQ4_NL  20u
+#define GGML_TYPE_IQ3_S   21u
+#define GGML_TYPE_IQ2_S   22u
+#define GGML_TYPE_IQ4_XS  23u
+#define GGML_TYPE_IQ1_M   29u
+#define GGML_TYPE_BF16    30u
+#define GGML_TYPE_TQ1_0   34u
+#define GGML_TYPE_TQ2_0   35u
+#define GGML_TYPE_MXFP4   39u
+#define GGML_TYPE_NVFP4   40u
+#define GGML_TYPE_Q1_0    41u
+#define GGML_TYPE_Q2_0    42u
+
+#endif // !defined(GGML_TYPE_IDS_COMP)

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

@@ -0,0 +1,2 @@
+void init_iq_shmem(uvec3 wgsize) {
+}

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

@@ -12,9 +12,9 @@
 
 #include "types.glsl"
 #include "fa_types.glsl"
-#define FaTypeV FA_TYPE_F32
+#define FaTypeV GGML_TYPE_F32
 
-layout(constant_id = 0) const uint FaTypeK = FA_TYPE_F32;
+layout(constant_id = 0) const uint FaTypeK = GGML_TYPE_F32;
 layout(constant_id = 1) const uint FaBlockBytesK = 4;
 layout(constant_id = 2) const uint SUBGROUP_SIZE = 32;
 
@@ -84,11 +84,11 @@ void main() {
     const uint k_block_elems = fa_block_elems(FaTypeK);
     const uint k_elem_bytes = FaBlockBytesK / k_block_elems;
 
-    if (FaTypeK == FA_TYPE_F16) {
+    if (FaTypeK == GGML_TYPE_F16) {
         k_row[tid] = float(k_f16[k_offset / k_elem_bytes + tid]);
-    } else if (FaTypeK == FA_TYPE_F32) {
+    } else if (FaTypeK == GGML_TYPE_F32) {
         k_row[tid] = k_f32[k_offset / k_elem_bytes + tid];
-    } else if (FaTypeK == FA_TYPE_BF16) {
+    } else if (FaTypeK == GGML_TYPE_BF16) {
         k_row[tid] = bf16_to_fp32(uint(k_bf16[k_offset / k_elem_bytes + tid]));
     } else if (4 * tid < HEAD_SIZE) {
         const uint coord = 4 * tid;

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

@@ -9,6 +9,9 @@
 #if defined(DATA_A_IQ1_M)
 #extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
 #endif
+#if !defined(DATA_A_F32) && !defined(DATA_A_F16) && !defined(DATA_A_BF16)
+#extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
+#endif
 
 #if defined(DATA_A_BF16) && defined(COOPMAT)
 #extension GL_EXT_bfloat16 : enable
@@ -28,24 +31,54 @@
 #extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
 #endif
 
+#ifdef MULMAT_QUANT
+#include "ggml_type_ids.glsl"
+layout (constant_id = 12) const uint MmTypeA = 0;
+#endif
+
 #include "types.glsl"
 #include "dot_product_funcs.glsl"
 
+#ifndef MULMAT_QUANT
 #ifndef LOAD_VEC_A
 #define LOAD_VEC_A 1
 #endif
+#endif
 #ifndef LOAD_VEC_B
 #define LOAD_VEC_B 1
 #endif
 
 layout (constant_id = 11) const uint ALIGNED = 0;
 
+#ifdef MULMAT_QUANT
+
+uint mm_load_vec_a() {
+    switch (MmTypeA) {
+    case GGML_TYPE_Q1_0:
+    case GGML_TYPE_Q4_0:
+    case GGML_TYPE_Q4_1:
+    case GGML_TYPE_Q5_1:
+        return 8u;
+    case GGML_TYPE_Q2_0:
+    case GGML_TYPE_Q5_0:
+    case GGML_TYPE_Q8_0:
+    case GGML_TYPE_Q2_K:
+    case GGML_TYPE_Q4_K:
+    case GGML_TYPE_Q5_K:
+        return 4u;
+    default:
+        return 2u;
+    }
+}
+#endif
+
 #if !defined(TO_FLOAT_TYPE)
 #define TO_FLOAT_TYPE FLOAT_TYPE
 #endif
 
 layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
 
+#ifndef MULMAT_QUANT
 layout (binding = 0) readonly buffer A {A_TYPE data_a[];};
 #if defined(DATA_A_F32)
 layout (binding = 0) readonly buffer A_SCALAR {float data_a_scalar[];};
@@ -60,6 +93,30 @@ layout (binding = 0) readonly buffer A_PACKED16 {A_TYPE_PACKED16 data_a_packed16
 #if defined(A_TYPE_PACKED32)
 layout (binding = 0) readonly buffer A_PACKED32 {A_TYPE_PACKED32 data_a_packed32[];};
 #endif
+#else
+// Unpacked struct aliases
+layout (binding = 0) readonly buffer BUF_Q1_0     { block_q1_0 data[];     } a_q1_0;
+layout (binding = 0) readonly buffer BUF_Q2_0     { block_q2_0 data[];     } a_q2_0;
+layout (binding = 0) readonly buffer BUF_Q2_K     { block_q2_K data[];     } a_q2_k;
+layout (binding = 0) readonly buffer BUF_Q3_K     { block_q3_K data[];     } a_q3_k;
+layout (binding = 0) readonly buffer BUF_Q4_K     { block_q4_K data[];     } a_q4_k;
+layout (binding = 0) readonly buffer BUF_Q5_K     { block_q5_K data[];     } a_q5_k;
+layout (binding = 0) readonly buffer BUF_Q6_K     { block_q6_K data[];     } a_q6_k;
+layout (binding = 0) readonly buffer BUF_TQ1_0    { block_tq1_0 data[];    } a_tq1_0;
+layout (binding = 0) readonly buffer BUF_TQ2_0    { block_tq2_0 data[];    } a_tq2_0;
+// Packed16 aliases
+layout (binding = 0) readonly buffer BUF_Q4_0_P16   { block_q4_0_packed16 data[];   } a_q4_0_p16;
+layout (binding = 0) readonly buffer BUF_Q5_0_P16   { block_q5_0_packed16 data[];   } a_q5_0_p16;
+layout (binding = 0) readonly buffer BUF_Q8_0_P16   { block_q8_0_packed16 data[];   } a_q8_0_p16;
+layout (binding = 0) readonly buffer BUF_Q3_K_P16   { block_q3_K_packed16 data[];   } a_q3_k_p16;
+layout (binding = 0) readonly buffer BUF_Q6_K_P16   { block_q6_K_packed16 data[];   } a_q6_k_p16;
+// Packed32 aliases
+layout (binding = 0) readonly buffer BUF_Q4_1_P32   { block_q4_1_packed32 data[];   } a_q4_1_p32;
+layout (binding = 0) readonly buffer BUF_Q5_1_P32   { block_q5_1_packed32 data[];   } a_q5_1_p32;
+layout (binding = 0) readonly buffer BUF_Q2_K_P32   { block_q2_K_packed32 data[];   } a_q2_k_p32;
+layout (binding = 0) readonly buffer BUF_Q4_K_P32   { block_q4_K_packed32 data[];   } a_q4_k_p32;
+layout (binding = 0) readonly buffer BUF_Q5_K_P32   { block_q5_K_packed32 data[];   } a_q5_k_p32;
+#endif
 
 layout (binding = 1) readonly buffer B {B_TYPE data_b[];};
 layout (binding = 1) readonly buffer B_SCALAR {B_TYPE_SCALAR data_b_scalar[];};
@@ -121,8 +178,13 @@ layout (constant_id = 3) const uint BK = 16;  // Assumed to be 32 if working wit
 #endif
 
 #ifdef COOPMAT
+#ifdef MULMAT_QUANT
+layout(constant_id = 13) const uint SHMEM_STRIDE_PAD = 4;
+layout(constant_id = 14) const bool APPLY_SLM_A_RESHAPE = false;
+#else
 layout(constant_id = 12) const uint SHMEM_STRIDE_PAD = 4;
 layout(constant_id = 13) const bool APPLY_SLM_A_RESHAPE = false;
+#endif
 #else
 const uint SHMEM_STRIDE_PAD = 1;
 const bool APPLY_SLM_A_RESHAPE = false;
@@ -141,6 +203,10 @@ shared ACC_TYPE coopmat_stage[TM * TN * NUM_WARPS];
 #include "mul_mm_id_funcs.glsl"
 #include "mul_mm_funcs.glsl"
 
+#ifdef MULMAT_QUANT
+#include "iq_shmem_init.glsl"
+#endif
+
 void main() {
     const uint ic = gl_WorkGroupID.y;
 
@@ -150,7 +216,7 @@ void main() {
         return;
     }
 #endif
-#ifdef NEEDS_INIT_IQ_SHMEM
+#if defined(NEEDS_INIT_IQ_SHMEM) || defined(MULMAT_QUANT)
     init_iq_shmem(gl_WorkGroupSize);
 #endif
 
@@ -200,9 +266,12 @@ void main() {
 #if defined(DATA_A_F32) || defined(DATA_A_F16) || defined(DATA_A_BF16)
     const uint LOAD_VEC_A_EFF = (ALIGNED != 0) ? LOAD_VEC_A : 1;
     const uint LOAD_VEC_BATCH_A = (ALIGNED != 0) ? 1 : 2;
-#else
+#elif !defined(MULMAT_QUANT)
     const uint LOAD_VEC_A_EFF = LOAD_VEC_A;
     const uint LOAD_VEC_BATCH_A = 1;
+#else
+    const uint LOAD_VEC_A_EFF = mm_load_vec_a();
+    const uint LOAD_VEC_BATCH_A = 1;
 #endif
     const uint LOAD_VEC_B_EFF = (ALIGNED != 0) ? LOAD_VEC_B : 1;
     const uint LOAD_VEC_BATCH_B = (ALIGNED != 0) ? 1 : 2;

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

@@ -21,6 +21,13 @@
 #extension GL_EXT_bfloat16 : enable
 #endif
 
+#include "ggml_type_ids.glsl"
+
+#ifdef MULMAT_QUANT
+layout (constant_id = 7) const uint MmTypeA = 0;
+layout (constant_id = 8) const uint MmABlockBytes = 2;
+#endif
+
 #include "types.glsl"
 #include "utils.glsl"
 
@@ -37,6 +44,24 @@ layout (constant_id = 4) const bool enable_smaller_matrices = false;
 const uint BNover2 = enable_smaller_matrices ? (BN / 2) : BN;
 const uint BNover4 = enable_smaller_matrices ? (BN / 4) : BN;
 layout (constant_id = 5) const uint ALIGNED = 0;
+layout (constant_id = 6) const uint subgroup_size = 32;
+
+#ifdef MULMAT_QUANT
+
+uint mm_quant_k() {
+    switch (MmTypeA) {
+    case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1:
+    case GGML_TYPE_Q8_0:
+        return 32u;
+    case GGML_TYPE_Q1_0:
+        return 128u;
+    case GGML_TYPE_Q2_0:
+        return 64u;
+    default:
+        return 256u;
+    }
+}
+#endif
 
 layout (push_constant) uniform parameter
 {
@@ -72,21 +97,73 @@ layout (push_constant) uniform parameter
 } p;
 
 
+#ifndef MULMAT_QUANT
 layout (binding = 0) readonly buffer A {A_TYPE data_a[];};
+#else
+layout (binding = 0) readonly buffer A {uint8_t data_a[];};
+#endif
 layout (binding = 1) readonly buffer B {B_TYPE data_b[];};
 layout (binding = 2) writeonly buffer D {D_TYPE data_d[];};
 #if defined(MUL_MAT_ID) && defined(GGML_VULKAN_COOPMAT2_DECODE_VECTOR)
 layout (binding = 1) readonly buffer B4 {B_TYPEV4 data_b_v4[];};
 #endif
 
-#if QUANT_K > 1
+#if defined(MULMAT_QUANT) || QUANT_K > 1
 #include "dequant_funcs_cm2.glsl"
+#ifndef MULMAT_QUANT
+// Per-type path: use the alias set by dequant_funcs_cm2.glsl
 #if defined(dequantFuncA_v) && defined(GGML_VULKAN_COOPMAT2_DECODE_VECTOR)
 #define DECODEFUNCA , dequantFuncA, dequantFuncA_v
 #else
 #define DECODEFUNCA , dequantFuncA
 #endif
 #else
+layout(buffer_reference, std430, buffer_reference_align = 1) buffer decodeBufA {
+    uint8_t raw[MmABlockBytes];
+};
+
+float16_t mmDecodeA(const in decodeBufA bl_in, const in uint blockCoords[2], const in uint coordInBlock[2]) {
+    switch (MmTypeA) {
+    case GGML_TYPE_Q1_0:  return dequantFuncQ1_0 (decodeBufQ1_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q2_0:  return dequantFuncQ2_0 (decodeBufQ2_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q4_0:  return dequantFuncQ4_0 (decodeBufQ4_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q4_1:  return dequantFuncQ4_1 (decodeBufQ4_1 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q5_0:  return dequantFuncQ5_0 (decodeBufQ5_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q5_1:  return dequantFuncQ5_1 (decodeBufQ5_1 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q8_0:  return dequantFuncQ8_0 (decodeBufQ8_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q2_K:  return dequantFuncQ2_K (decodeBufQ2_K (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q3_K:  return dequantFuncQ3_K (decodeBufQ3_K (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q6_K:  return dequantFuncQ6_K (decodeBufQ6_K (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_TQ1_0: return dequantFuncTQ1_0(decodeBufTQ1_0(bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_TQ2_0: return dequantFuncTQ2_0(decodeBufTQ2_0(bl_in), blockCoords, coordInBlock);
+    default:              return float16_t(0);
+    }
+}
+
+#ifdef GGML_VULKAN_COOPMAT2_DECODE_VECTOR
+f16vec4 mmDecodeA_v(const in decodeBufA bl_in, const in uint blockCoords[2], const in uint coordInBlock[2]) {
+    switch (MmTypeA) {
+    case GGML_TYPE_Q1_0:  return dequantFuncQ1_0_v (decodeBufQ1_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q2_0:  return dequantFuncQ2_0_v (decodeBufQ2_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q4_0:  return dequantFuncQ4_0_v (decodeBufQ4_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q4_1:  return dequantFuncQ4_1_v (decodeBufQ4_1 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q5_0:  return dequantFuncQ5_0_v (decodeBufQ5_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q5_1:  return dequantFuncQ5_1_v (decodeBufQ5_1 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q8_0:  return dequantFuncQ8_0_v (decodeBufQ8_0 (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q2_K:  return dequantFuncQ2_K_v (decodeBufQ2_K (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q3_K:  return dequantFuncQ3_K_v (decodeBufQ3_K (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_Q6_K:  return dequantFuncQ6_K_v (decodeBufQ6_K (bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_TQ1_0: return dequantFuncTQ1_0_v(decodeBufTQ1_0(bl_in), blockCoords, coordInBlock);
+    case GGML_TYPE_TQ2_0: return dequantFuncTQ2_0_v(decodeBufTQ2_0(bl_in), blockCoords, coordInBlock);
+    default:              return f16vec4(0);
+    }
+}
+#define DECODEFUNCA , mmDecodeA, mmDecodeA_v
+#else
+#define DECODEFUNCA , mmDecodeA
+#endif
+#endif
+#else
 #define DECODEFUNCA
 #endif
 
@@ -114,7 +191,6 @@ layout(buffer_reference, std430, buffer_reference_align = 2) buffer decodeBufB {
 };
 
 uint _ne1;
-layout (constant_id = 6) const uint subgroup_size = 32;
 shared uvec4 ballots_sh[BLOCK_SIZE / subgroup_size];
 
 B_TYPE decodeFuncB(const in decodeBufB bl, const in uint blockCoords[2], const in uint coordInBlock[2])
@@ -246,6 +322,10 @@ void load_row_ids_hoisted(uint expert_idx, uint ic) {
 }
 #endif
 
+#ifdef MULMAT_QUANT
+#include "iq_shmem_init.glsl"
+#endif
+
 void main() {
     const uint tid = gl_LocalInvocationIndex;
     const uint ic = gl_WorkGroupID.y;
@@ -264,7 +344,7 @@ void main() {
 #endif
 #endif
 
-#ifdef NEEDS_INIT_IQ_SHMEM
+#if defined(NEEDS_INIT_IQ_SHMEM) || defined(MULMAT_QUANT)
     init_iq_shmem(gl_WorkGroupSize);
 #endif
 
@@ -305,22 +385,33 @@ void main() {
     const uint end_k = min(p.K, (ik + 1) * p.k_split);
 #endif
 
+#ifdef MULMAT_QUANT
+    const uint qk = mm_quant_k();
+#else
+    const uint qk = QUANT_K;
+#endif
+
 #ifdef MUL_MAT_ID
-    uint pos_a = expert_idx * (p.batch_stride_a / QUANT_K);
+    uint pos_a = expert_idx * (p.batch_stride_a / qk);
     uint pos_b = 0;
 #else
-    uint pos_a = batch_idx_a * (p.batch_stride_a / QUANT_K);
+    uint pos_a = batch_idx_a * (p.batch_stride_a / qk);
     uint pos_b = batch_idx * p.batch_stride_b;
     uint pos_d = batch_idx * p.batch_stride_d + ik * p.batch_stride_d * p.num_batches;
 #endif
 
-    uint stride_a = p.stride_a / QUANT_K;
+#ifdef MULMAT_QUANT
+    // pos_a is a byte offset into the raw buffer; strides stay in block units
+    pos_a *= MmABlockBytes;
+#endif
+
+    uint stride_a = p.stride_a / qk;
     uint stride_b = p.stride_b;
 
     // Hint to the compiler that values are aligned (want 16B alignment).
     // Quants are always block-aligned, no alignment needed.
     if (ALIGNED != 0) {
-#if QUANT_K == 1
+#if !defined(MULMAT_QUANT) && QUANT_K == 1
         stride_a &= ~7;
 #endif
         stride_b &= ~7;
@@ -335,10 +426,8 @@ void main() {
 #endif
     tensorLayoutNV<2, gl_CooperativeMatrixClampModeConstantNV> tensorLayoutD = createTensorLayoutNV(2, gl_CooperativeMatrixClampModeConstantNV);
 
-#if QUANT_K > 1
-    tensorLayoutA = setTensorLayoutBlockSizeNV(tensorLayoutA, 1, QUANT_K);
-    tensorLayoutAClamp = setTensorLayoutBlockSizeNV(tensorLayoutAClamp, 1, QUANT_K);
-#endif
+    tensorLayoutA = setTensorLayoutBlockSizeNV(tensorLayoutA, 1, qk);
+    tensorLayoutAClamp = setTensorLayoutBlockSizeNV(tensorLayoutAClamp, 1, qk);
 #if defined(MUL_MAT_ID) && defined(GGML_VULKAN_COOPMAT2_DECODE_VECTOR)
     tensorLayoutB = setTensorLayoutBlockSizeNV(tensorLayoutB, 1, BK);
 #endif
@@ -368,19 +457,19 @@ void main() {
     const uint START_ALIGN_K = 256;
     // For Qi_K (block size 256), unroll whole 256 element tiles.
     // For legacy quants (block size 32), unroll 8x.
-    const uint UNROLL_K = (QUANT_K == 256) ? 256 : (BK * 8);
+    const uint UNROLL_K = (qk == 256) ? 256 : (BK * 8);
     const uint unroll_count = UNROLL_K / BK;
 
     // Detect a fast path where all loads are entirely in bounds and no clamping is required
     if ((ir + 1) * BM <= p.M && (ic + 1) * BN <= p.padded_N && (start_k % START_ALIGN_K) == 0 && (end_k % BK) == 0 &&
-#if QUANT_K == 1
+#if !defined(MULMAT_QUANT) && QUANT_K == 1
         (stride_a % 8) == 0 &&
 #endif
         (stride_b % 8) == 0) {
         // Hint to the compiler that values are aligned (want 16B alignment)
         start_k &= ~(START_ALIGN_K-1);
         stride_b &= ~7;
-#if QUANT_K == 1
+#if !defined(MULMAT_QUANT) && QUANT_K == 1
         stride_a &= ~7;
 #endif
 
@@ -551,10 +640,10 @@ void main() {
             [[dont_unroll]]
             for (uint block_k = start_k, i = 0; i < k_iters; block_k += BK, ++i) {
 
-                if ((block_k % QUANT_K) == 0) {
+                if ((block_k % qk) == 0) {
                     store_scales(tid);
                 }
-                if (block_k + BK < end_k && ((block_k + BK) % QUANT_K) == 0) {
+                if (block_k + BK < end_k && ((block_k + BK) % qk) == 0) {
                     fetch_scales(ir * BM, pos_a, stride_a, block_k + BK, tid, false);
                 }
 
@@ -595,10 +684,10 @@ void main() {
             [[dont_unroll]]
             for (uint block_k = start_k, i = 0; i < k_iters; block_k += BK, ++i) {
 
-                if ((block_k % QUANT_K) == 0) {
+                if ((block_k % qk) == 0) {
                     store_scales(tid);
                 }
-                if (block_k + BK < end_k && ((block_k + BK) % QUANT_K) == 0) {
+                if (block_k + BK < end_k && ((block_k + BK) % qk) == 0) {
                     fetch_scales(ir * BM, pos_a, stride_a, block_k + BK, tid, false);
                 }
 
@@ -639,10 +728,10 @@ void main() {
         [[dont_unroll]]
         for (uint block_k = start_k, i = 0; i < k_iters; block_k += BK, ++i) {
 
-            if ((block_k % QUANT_K) == 0) {
+            if ((block_k % qk) == 0) {
                 store_scales(tid);
             }
-            if (block_k + BK < end_k && ((block_k + BK) % QUANT_K) == 0) {
+            if (block_k + BK < end_k && ((block_k + BK) % qk) == 0) {
                 fetch_scales(ir * BM, pos_a, stride_a, block_k + BK, tid, false);
             }
 

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

@@ -955,6 +955,7 @@ shared uint16_t iq1s_grid[2048];
 shared uint32_t iq1s_grid_gpu[2048];
 #endif
 
+#if defined(DATA_A_IQ1_S) || defined(DATA_A_IQ1_M)
 #define NEEDS_INIT_IQ_SHMEM
 void init_iq_shmem(uvec3 wgsize)
 {
@@ -978,6 +979,17 @@ void init_iq_shmem(uvec3 wgsize)
     barrier();
 }
 #endif
+#endif
+
+#if defined(DATA_A_IQ2_XXS) || defined(DATA_A_IQ2_XS) || defined(DATA_A_IQ2_S)
+#if defined(DATA_A_IQ2_S)
+shared uvec2 iq2s_grid[1024];
+#elif defined(DATA_A_IQ2_XS)
+shared uvec2 iq2xs_grid[512];
+#else
+shared uvec2 iq2xxs_grid[256];
+#endif
+#endif
 
 #define QUANT_K_IQ2_XXS 256
 #define QUANT_R_IQ2_XXS 1
@@ -1063,8 +1075,7 @@ const uvec2[256] iq2xxs_grid_const = {
     uvec2(0x08080808, 0x2b2b082b), uvec2(0x08192b08, 0x2b2b1908), uvec2(0x19190808, 0x2b2b2b08), uvec2(0x08081908, 0x2b2b2b19)
 };
 
-shared uvec2 iq2xxs_grid[256];
-
+#if defined(DATA_A_IQ2_XXS)
 #define NEEDS_INIT_IQ_SHMEM
 void init_iq_shmem(uvec3 wgsize)
 {
@@ -1076,12 +1087,15 @@ void init_iq_shmem(uvec3 wgsize)
     }
     barrier();
 }
+#endif
 
+#if defined(DATA_A_IQ2_XXS)
 #define QUANT_K QUANT_K_IQ2_XXS
 #define QUANT_R QUANT_R_IQ2_XXS
 #define A_TYPE block_iq2_xxs
 #define A_TYPE_PACKED16 block_iq2_xxs_packed16
 #endif
+#endif
 
 #define QUANT_K_IQ2_XS 256
 #define QUANT_R_IQ2_XS 1
@@ -1233,8 +1247,7 @@ const uvec2 iq2xs_grid_const[512] = {
     uvec2(0x082b2b08, 0x2b2b2b2b), uvec2(0x082b2b2b, 0x2b2b2b2b), uvec2(0x2b190819, 0x2b2b2b2b), uvec2(0x2b2b2b2b, 0x2b2b2b2b),
 };
 
-shared uvec2 iq2xs_grid[512];
-
+#if defined(DATA_A_IQ2_XS)
 #define NEEDS_INIT_IQ_SHMEM
 void init_iq_shmem(uvec3 wgsize)
 {
@@ -1246,12 +1259,15 @@ void init_iq_shmem(uvec3 wgsize)
     }
     barrier();
 }
+#endif
 
+#if defined(DATA_A_IQ2_XS)
 #define QUANT_K QUANT_K_IQ2_XS
 #define QUANT_R QUANT_R_IQ2_XS
 #define A_TYPE block_iq2_xs
 #define A_TYPE_PACKED16 block_iq2_xs_packed16
 #endif
+#endif
 
 #define QUANT_K_IQ2_S 256
 #define QUANT_R_IQ2_S 1
@@ -1533,8 +1549,7 @@ const uvec2 iq2s_grid_const[1024] = {
     uvec2(0x082b082b, 0x2b2b2b2b), uvec2(0x082b2b08, 0x2b2b2b2b), uvec2(0x2b082b08, 0x2b2b2b2b), uvec2(0x2b2b2b2b, 0x2b2b2b2b)
 };
 
-shared uvec2 iq2s_grid[1024];
-
+#if defined(DATA_A_IQ2_S)
 #define NEEDS_INIT_IQ_SHMEM
 void init_iq_shmem(uvec3 wgsize)
 {
@@ -1546,12 +1561,23 @@ void init_iq_shmem(uvec3 wgsize)
     }
     barrier();
 }
+#endif
 
+#if defined(DATA_A_IQ2_S)
 #define QUANT_K QUANT_K_IQ2_S
 #define QUANT_R QUANT_R_IQ2_S
 #define A_TYPE block_iq2_s
 #define A_TYPE_PACKED16 block_iq2_s_packed16
 #endif
+#endif
+
+#if defined(DATA_A_IQ3_XXS) || defined(DATA_A_IQ3_S)
+#if defined(DATA_A_IQ3_S)
+shared uint32_t iq3s_grid[512];
+#else
+shared uint32_t iq3xxs_grid[256];
+#endif
+#endif
 
 #define QUANT_K_IQ3_XXS 256
 #define QUANT_R_IQ3_XXS 1
@@ -1605,8 +1631,7 @@ const uint32_t iq3xxs_grid_const[256] = {
     0x3e1c1c1c, 0x3e1c3404, 0x3e24140c, 0x3e24240c, 0x3e2c0404, 0x3e2c0414, 0x3e2c1424, 0x3e341c04,
 };
 
-shared uint32_t iq3xxs_grid[256];
-
+#if defined(DATA_A_IQ3_XXS)
 #define NEEDS_INIT_IQ_SHMEM
 void init_iq_shmem(uvec3 wgsize)
 {
@@ -1618,12 +1643,15 @@ void init_iq_shmem(uvec3 wgsize)
     }
     barrier();
 }
+#endif
 
+#if defined(DATA_A_IQ3_XXS)
 #define QUANT_K QUANT_K_IQ3_XXS
 #define QUANT_R QUANT_R_IQ3_XXS
 #define A_TYPE block_iq3_xxs
 #define A_TYPE_PACKED16 block_iq3_xxs_packed16
 #endif
+#endif
 
 #define QUANT_K_IQ3_S 256
 #define QUANT_R_IQ3_S 1
@@ -1715,8 +1743,7 @@ const uint32_t iq3s_grid_const[512] = {
     0x0f090307, 0x0f090501, 0x0f090b01, 0x0f0b0505, 0x0f0b0905, 0x0f0d0105, 0x0f0d0703, 0x0f0f0101,
 };
 
-shared uint32_t iq3s_grid[512];
-
+#if defined(DATA_A_IQ3_S)
 #define NEEDS_INIT_IQ_SHMEM
 void init_iq_shmem(uvec3 wgsize)
 {
@@ -1728,12 +1755,15 @@ void init_iq_shmem(uvec3 wgsize)
     }
     barrier();
 }
+#endif
 
+#if defined(DATA_A_IQ3_S)
 #define QUANT_K QUANT_K_IQ3_S
 #define QUANT_R QUANT_R_IQ3_S
 #define A_TYPE block_iq3_s
 #define A_TYPE_PACKED16 block_iq3_s_packed16
 #endif
+#endif
 
 #define QUANT_K_IQ4_XS 256
 #define QUANT_R_IQ4_XS 1
@@ -1847,6 +1877,7 @@ const int8_t kvalues_iq4nl_const[16] = {
 
 shared FLOAT_TYPE kvalues_iq4nl[16];
 
+#if defined(DATA_A_IQ4_NL) || defined(DATA_A_IQ4_XS)
 #define NEEDS_INIT_IQ_SHMEM
 void init_iq_shmem(uvec3 wgsize)
 {
@@ -1857,6 +1888,7 @@ void init_iq_shmem(uvec3 wgsize)
     barrier();
 }
 #endif
+#endif
 
 #if defined(DATA_A_MXFP4) || defined(DATA_A_NVFP4)
 #if !defined(USE_OCP_FP4)
@@ -1886,7 +1918,7 @@ float ue4m3_to_fp32_build(uint u) {
 }
 #endif
 
-#if !defined(USE_OCP_FP4)
+#if (defined(DATA_A_MXFP4) || defined(DATA_A_NVFP4)) && !defined(USE_OCP_FP4)
 #define NEEDS_INIT_IQ_SHMEM
 void init_iq_shmem(uvec3 wgsize)
 {

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

@@ -246,6 +246,17 @@ bool is_iq_quant(const std::string& type_name) {
     return string_starts_with(type_name, "iq");
 }
 
+bool is_lut_quant(const std::string& type_name) {
+    return is_iq_quant(type_name) || type_name == "mxfp4" || type_name == "nvfp4";
+}
+
+std::string lut_load_vec_a(const std::string& type_name) {
+    if (type_name == "iq1_s" || type_name == "iq1_m" || type_name == "iq2_xxs" || type_name == "iq2_xs" || type_name == "iq2_s") {
+        return "8";
+    }
+    return "4";
+}
+
 static const char path_separator = '/';
 
 std::string join_paths(const std::string& path1, const std::string& path2) {
@@ -583,20 +594,28 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c
     }
 
     for (const auto& tname : type_names) {
-        std::string load_vec_quant = "2";
-        if ((tname == "q1_0") || (tname == "q4_0") || (tname == "q4_1") || (tname == "q5_1") || (tname == "iq1_s") || (tname == "iq1_m") || (tname == "iq2_xxs") || (tname == "iq2_xs") || (tname == "iq2_s"))
-            load_vec_quant = "8";
-        else if ((tname == "q2_0") || (tname == "q5_0") || (tname == "q8_0") || (tname == "q2_k") || (tname == "q4_k") || (tname == "q5_k") || (tname == "iq3_xxs") || (tname == "iq3_s") || (tname == "iq4_xs") || (tname == "iq4_nl") || (tname == "mxfp4") || (tname == "nvfp4"))
-            load_vec_quant = "4";
-
         if (tname == "bf16") {
             continue;
         }
 
-        std::string data_a_key = "DATA_A_" + to_uppercase(tname);
-        // For aligned matmul loads
-        std::string load_vec_a = (coopmat2 || tname == "f32" || tname == "f16" || tname == "bf16") ? load_vec : load_vec_quant;
+        // Float types keep per-type compilation (different accumulation loop structure)
+        if (tname == "f32" || tname == "f16") {
+            std::string data_a_key = "DATA_A_" + to_uppercase(tname);
 
+            const std::map<std::string, std::string> float_type_dict = {
+                {"FLOAT_TYPE",   FLOAT_TYPE(1, tname)},
+                {"FLOAT_TYPEV2", FLOAT_TYPE(2, tname)},
+                {"FLOAT_TYPEV4", FLOAT_TYPE(4, tname)},
+                {"FLOAT_TYPEV8", FLOAT_TYPE(8, tname)},
+            };
+
+            if (!coopmat2) {
+                string_to_spv(shader_name + "_" + tname + "_f32" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", load_vec}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
+            }
+            continue;
+        }
+
+        std::string data_a_key = "DATA_A_" + to_uppercase(tname);
         const std::map<std::string, std::string> float_type_dict = {
             {"FLOAT_TYPE",   FLOAT_TYPE(1, tname)},
             {"FLOAT_TYPEV2", FLOAT_TYPE(2, tname)},
@@ -604,30 +623,52 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c
             {"FLOAT_TYPEV8", FLOAT_TYPE(8, tname)},
         };
 
-        // don't generate f32 variants for coopmat2
-        if (!coopmat2) {
-            string_to_spv(shader_name + "_" + tname + "_f32" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", load_vec_a}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
-        }
-
-        if (tname != "f16" && tname != "f32") {
-            string_to_spv(shader_name + "_" + tname + "_f16" + dot2_sfx, source_name,  merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", load_vec_a}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
-        }
-
-#if defined(GGML_VULKAN_FLOAT_E2M1_GLSLC_SUPPORT) && defined(GGML_VULKAN_FLOAT_E4M3_GLSLC_SUPPORT)
-        if ((coopmat || coopmat2) && (tname == "mxfp4" || tname == "nvfp4")) {
-            if (!coopmat2) {
-                string_to_spv(shader_name + "_" + tname + "_f32_ocp" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"USE_OCP_FP4", "1"}, {"LOAD_VEC_A", load_vec_a}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
-            }
-            string_to_spv(shader_name + "_" + tname + "_f16_ocp" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"USE_OCP_FP4", "1"}, {"LOAD_VEC_A", load_vec_a}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
-        }
-#endif
-
 #if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT)
-        // Integer dot mmq performs better with f32 accumulators (different shader, skip for dot2)
         if (!f16acc && !coopmat && !coopmat2 && !dot2 && (is_legacy_quant(tname) || is_k_quant(tname) || tname == "mxfp4")) {
             string_to_spv(shader_name + "_" + tname + "_q8_1", "mul_mmq.comp", merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"D_TYPE", "float"},}), fp16, coopmat, coopmat2, f16acc);
         }
 #endif
+
+        if (is_lut_quant(tname)) {
+            std::string lva = lut_load_vec_a(tname);
+
+            string_to_spv(shader_name + "_" + tname + "_f16" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", lva}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
+
+            if (!coopmat2) {
+                string_to_spv(shader_name + "_" + tname + "_f32" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", lva}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
+            }
+
+#if defined(GGML_VULKAN_FLOAT_E2M1_GLSLC_SUPPORT) && defined(GGML_VULKAN_FLOAT_E4M3_GLSLC_SUPPORT)
+            if ((tname == "mxfp4" || tname == "nvfp4") && (coopmat || coopmat2)) {
+                string_to_spv(shader_name + "_" + tname + "_f16_ocp" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"USE_OCP_FP4", "1"}, {"LOAD_VEC_A", lva}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
+                if (!coopmat2) {
+                    string_to_spv(shader_name + "_" + tname + "_f32_ocp" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"USE_OCP_FP4", "1"}, {"LOAD_VEC_A", lva}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
+                }
+            }
+#endif
+            continue;
+        }
+
+        // dedicated shader needed due to regression on Ampere
+        if (coopmat2 && (tname == "q4_k" || tname == "q5_k")) {
+            string_to_spv(shader_name + "_" + tname + "_f16" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", load_vec}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
+        }
+    }
+
+    // Quant shader: one SPIR-V for all quant types, selected via MmTypeA spec constant
+    {
+        const std::map<std::string, std::string> quant_float_type_dict = {
+            {"FLOAT_TYPE",   FLOAT_TYPE(1, "q4_0")},
+            {"FLOAT_TYPEV2", FLOAT_TYPE(2, "q4_0")},
+            {"FLOAT_TYPEV4", FLOAT_TYPE(4, "q4_0")},
+            {"FLOAT_TYPEV8", FLOAT_TYPE(8, "q4_0")},
+        };
+
+        string_to_spv(shader_name + "_quant_f16" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, quant_float_type_dict), {{"MULMAT_QUANT", "1"}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
+
+        if (!coopmat2) {
+            string_to_spv(shader_name + "_quant_f32" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, quant_float_type_dict), {{"MULMAT_QUANT", "1"}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
+        }
     }
 }
 

gguf-py/pyproject.toml

@@ -7,7 +7,7 @@ dynamic = ["classifiers"]
 readme = "README.md"
 authors = [{name = "GGML", email = "ggml@ggml.ai"}]
 requires-python = '>=3.10'
-dependencies = ['numpy (>=1.17)', 'tqdm (>=4.27)', 'pyyaml (>=5.1)', 'requests (>=2.25)']
+dependencies = ['numpy (>=2.2.6)', 'tqdm (>=4.27)', 'pyyaml (>=5.1)', 'requests (>=2.25)']
 classifiers = [
     "Programming Language :: Python :: 3",
     "License :: OSI Approved :: MIT License",

include/llama.h

@@ -1357,7 +1357,7 @@ extern "C" {
     LLAMA_API struct llama_sampler * llama_sampler_chain_get(      struct llama_sampler * chain, int32_t i);
 
     // the total number of samplers in the chain
-    LLAMA_API int                    llama_sampler_chain_n  (const struct llama_sampler * chain);
+    LLAMA_API int32_t                llama_sampler_chain_n  (const struct llama_sampler * chain);
 
     // after removing a sampler, the chain will no longer own it, and it will not be freed when the chain is freed
     LLAMA_API struct llama_sampler * llama_sampler_chain_remove(   struct llama_sampler * chain, int32_t i);

pyproject.toml

@@ -8,7 +8,7 @@ readme = "README.md"
 authors = [{name = "GGML", email = "ggml@ggml.ai"}]
 requires-python = '>=3.10,<3.15'
 dependencies = [
-    'numpy (>=1.26.4,<3.0.0)',
+    'numpy (>=2.2.6,<3.0.0)',
     'sentencepiece (>=0.1.98,<0.3.0)',
     'transformers (==4.57.6)',
     'protobuf (>=4.21.0,<5.0.0)',

requirements/requirements-convert_legacy_llama.txt

@@ -1,4 +1,4 @@
-numpy~=1.26.4
+numpy~=2.2.6
 sentencepiece>=0.1.98,<0.3.0
 
 transformers==4.57.6

requirements/requirements-gguf_editor_gui.txt

@@ -1,3 +1,3 @@
-numpy~=1.26.4
+numpy~=2.2.6
 PySide6~=6.9.0
 gguf>=0.17.0

requirements/requirements-server-bench.txt

@@ -1,5 +1,5 @@
 datasets~=4.8.0
 matplotlib~=3.10.0
-numpy~=1.26.4
+numpy~=2.2.6
 requests~=2.32.3
 tqdm~=4.67.1

requirements/requirements-tool_bench.txt

@@ -1,7 +1,7 @@
 aiohttp~=3.9.3
 pytest~=8.3.3
 matplotlib~=3.10.0
-numpy~=1.26.4
+numpy~=2.2.6
 openai~=2.14.0
 pandas~=2.2.3
 prometheus-client~=0.20.0

scripts/make-release-checks.sh

@@ -22,9 +22,9 @@ for arg in "$@"; do
     esac
 done
 
-MAJOR=$(grep "set(LLAMA_VERSION_MAJOR" "$REPO_ROOT/CMakeLists.txt" | grep -oP '\d+')
-MINOR=$(grep "set(LLAMA_VERSION_MINOR" "$REPO_ROOT/CMakeLists.txt" | grep -oP '\d+')
-PATCH=$(grep "set(LLAMA_VERSION_PATCH" "$REPO_ROOT/CMakeLists.txt" | grep -oP '\d+')
+MAJOR=$(grep "set(LLAMA_VERSION_MAJOR" "$REPO_ROOT/CMakeLists.txt" | sed 's/.*MAJOR \([0-9]*\).*/\1/')
+MINOR=$(grep "set(LLAMA_VERSION_MINOR" "$REPO_ROOT/CMakeLists.txt" | sed 's/.*MINOR \([0-9]*\).*/\1/')
+PATCH=$(grep "set(LLAMA_VERSION_PATCH" "$REPO_ROOT/CMakeLists.txt" | sed 's/.*PATCH \([0-9]*\).*/\1/')
 VERSION="v${MAJOR}.${MINOR}.${PATCH}"
 echo "Determined version: ${VERSION}"
 if [[ -n "${GITHUB_OUTPUT:-}" ]]; then
@@ -91,9 +91,9 @@ else
     fi
 fi
 
-MAJOR=$(grep "set(GGML_VERSION_MAJOR" "$REPO_ROOT/ggml/CMakeLists.txt" | grep -oP '\d+')
-MINOR=$(grep "set(GGML_VERSION_MINOR" "$REPO_ROOT/ggml/CMakeLists.txt" | grep -oP '\d+')
-PATCH=$(grep "set(GGML_VERSION_PATCH" "$REPO_ROOT/ggml/CMakeLists.txt" | grep -oP '\d+')
+MAJOR=$(grep "set(GGML_VERSION_MAJOR" "$REPO_ROOT/ggml/CMakeLists.txt" | sed 's/.*MAJOR \([0-9]*\).*/\1/')
+MINOR=$(grep "set(GGML_VERSION_MINOR" "$REPO_ROOT/ggml/CMakeLists.txt" | sed 's/.*MINOR \([0-9]*\).*/\1/')
+PATCH=$(grep "set(GGML_VERSION_PATCH" "$REPO_ROOT/ggml/CMakeLists.txt" | sed 's/.*PATCH \([0-9]*\).*/\1/')
 GGML_VERSION="v${MAJOR}.${MINOR}.${PATCH}"
 echo "Local ggml version: ${GGML_VERSION}"
 

src/llama-memory-hybrid-idx.cpp

@@ -55,6 +55,10 @@ llama_memory_hybrid_idx::llama_memory_hybrid_idx(
         // K-shift must not rotate them while the stream copies in the same update still apply
         hparams_idx.rope_type = LLAMA_ROPE_TYPE_NONE;
 
+        // fool llama_kv_cache into thinking this is a MLA cache, so it won't cache V tensors
+        hparams_idx.n_embd_head_k_mla_impl = model.hparams.indexer_head_size;
+        hparams_idx.n_embd_head_v_mla_impl = model.hparams.indexer_head_size;
+
         LLAMA_LOG_INFO("%s: creating indexer KV cache, size = %u cells\n", __func__, kv_size);
 
         return new llama_kv_cache(

src/llama-model.cpp

@@ -936,6 +936,7 @@ const char * llm_type_name(llm_type type) {
         case LLM_TYPE_17B_128E:      return "17Bx128E (Maverick)";
         case LLM_TYPE_A13B:          return "A13B";
         case LLM_TYPE_1B_A400M:      return "1B.A400M";
+        case LLM_TYPE_3B_A800M:      return "3B.A800M";
         case LLM_TYPE_7B_A1B:        return "7B.A1B";
         case LLM_TYPE_8B_A1B:        return "8B.A1B";
         case LLM_TYPE_7_9B_A1_3B:    return "7.9B.A1.3B";
@@ -946,6 +947,7 @@ const char * llm_type_name(llm_type type) {
         case LLM_TYPE_26B_A4B:       return "26B.A4B";
         case LLM_TYPE_30B_A3B:       return "30B.A3B";
         case LLM_TYPE_31B_A3_5B:     return "31B.A3.5B";
+        case LLM_TYPE_32B_A9B:       return "32B.A9B";
         case LLM_TYPE_35B_A3B:       return "35B.A3B";
         case LLM_TYPE_48B_A3B:       return "48B.A3B";
         case LLM_TYPE_75B_A9B:       return "75B.A9B";

src/llama-model.h

@@ -117,6 +117,7 @@ enum llm_type {
     LLM_TYPE_17B_128E, // llama4 Maverick
     LLM_TYPE_A13B,
     LLM_TYPE_1B_A400M, // Granite3 MoE
+    LLM_TYPE_3B_A800M, // Granite3 MoE
     LLM_TYPE_7B_A1B,
     LLM_TYPE_8B_A1B, // lfm2moe
     LLM_TYPE_7_9B_A1_3B, // Ling-3.0-tiny
@@ -127,6 +128,7 @@ enum llm_type {
     LLM_TYPE_26B_A4B, // Gemma4
     LLM_TYPE_30B_A3B,
     LLM_TYPE_31B_A3_5B,
+    LLM_TYPE_32B_A9B, // Granite4 Hybrid
     LLM_TYPE_35B_A3B, // Qwen3.5
     LLM_TYPE_48B_A3B, // Kimi Linear
     LLM_TYPE_75B_A9B, // Nemotron 3 Puzzle

src/llama-sampler.cpp

@@ -1006,7 +1006,7 @@ struct llama_sampler * llama_sampler_chain_remove(struct llama_sampler * chain,
     return result;
 }
 
-int llama_sampler_chain_n(const struct llama_sampler * chain) {
+int32_t llama_sampler_chain_n(const struct llama_sampler * chain) {
     const auto * p = (const llama_sampler_chain *) chain->ctx;
 
     return p->samplers.size();

src/models/bert.cpp

@@ -29,15 +29,13 @@ void llama_model_bert::load_arch_tensors(llama_model_loader &) {
     tok_embd     = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD,  "weight"), {n_embd, n_vocab}, 0);
     type_embd    = create_tensor(tn(LLM_TENSOR_TOKEN_TYPES, "weight"), {n_embd, n_token_types}, TENSOR_NOT_REQUIRED);
 
-    if (arch == LLM_ARCH_BERT) {
-        pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD,    "weight"), {n_embd, n_ctx_train}, 0);
+    pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD,    "weight"), {n_embd, n_ctx_train}, 0);
 
-        cls   = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
-        cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"),   {n_embd},         TENSOR_NOT_REQUIRED);
+    cls   = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
+    cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"),   {n_embd},         TENSOR_NOT_REQUIRED);
 
-        cls_out   = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
-        cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"),   {hparams.n_cls_out},         TENSOR_NOT_REQUIRED);
-    }
+    cls_out   = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
+    cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"),   {hparams.n_cls_out},         TENSOR_NOT_REQUIRED);
 
     tok_norm   = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "weight", 0), {n_embd}, 0);
     tok_norm_b = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "bias",   0), {n_embd}, 0);
@@ -53,20 +51,10 @@ void llama_model_bert::load_arch_tensors(llama_model_loader &) {
         layer.attn_out_norm   = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "weight", i), {n_embd}, 0);
         layer.attn_out_norm_b = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "bias", i),   {n_embd}, 0);
 
-        if (hparams.moe_every_n_layers > 0 && i % hparams.moe_every_n_layers == 1) {
-            layer.ffn_up_exps   = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS,   "weight", i), {  n_embd, n_ff,   n_expert}, 0);
-            layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {  n_ff,   n_embd, n_expert}, 0);
-            layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP,   "weight", i), {n_embd, n_expert}, 0);
-        } else {
-            layer.ffn_up     = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", i), {n_embd, n_ff}, 0);
-            layer.ffn_up_b   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "bias", i),   {n_ff}, TENSOR_NOT_REQUIRED);
-            layer.ffn_down   = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
-            layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i),   {n_embd}, TENSOR_NOT_REQUIRED);
-
-            if (arch == LLM_ARCH_NOMIC_BERT) {
-                layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
-            }
-        }
+        layer.ffn_up     = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", i), {n_embd, n_ff}, 0);
+        layer.ffn_up_b   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "bias", i),   {n_ff}, TENSOR_NOT_REQUIRED);
+        layer.ffn_down   = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
+        layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i),   {n_embd}, TENSOR_NOT_REQUIRED);
 
         layer.layer_out_norm   = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, 0);
         layer.layer_out_norm_b = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "bias", i),   {n_embd}, 0);

src/models/granite-hybrid.cpp

@@ -30,7 +30,7 @@ void llama_model_granite_hybrid::load_arch_hparams(llama_model_loader & ml) {
         case 768: type = LLM_TYPE_350M; break;
         case 1536: type = (hparams.n_ff() == 512 ? LLM_TYPE_7B_A1B : LLM_TYPE_1B); break;
         case 2048: case 2560: type = LLM_TYPE_3B; break;
-        case 4096: type = LLM_TYPE_32B; break;
+        case 4096: type = LLM_TYPE_32B_A9B; break;
         default: type = LLM_TYPE_UNKNOWN;
     }
 

src/models/granite-moe.cpp

@@ -9,8 +9,7 @@ void llama_model_granite_moe::load_arch_hparams(llama_model_loader & ml) {
 
     switch (hparams.n_layer()) {
         case 24: type = LLM_TYPE_1B_A400M; break;
-        case 32: type = LLM_TYPE_3B; break;
-        case 40: type = LLM_TYPE_3B; break;
+        case 32: type = LLM_TYPE_3B_A800M; break;
         // Add additional layer/vocab/etc checks here for other model sizes
         default: type = LLM_TYPE_UNKNOWN;
     }

src/models/granite.cpp

@@ -38,7 +38,16 @@ void llama_model_granite::load_arch_hparams(llama_model_loader & ml) {
 
     switch (hparams.n_layer()) {
         case 32: type = LLM_TYPE_3B; break;
-        case 40: type = LLM_TYPE_3B; break;
+        case 40: {
+            switch (hparams.n_embd) {
+                case 2048: type = LLM_TYPE_2B; break;
+                case 2560: type = LLM_TYPE_3B; break;
+                case 4096: type = LLM_TYPE_8B; break;
+                default: type = LLM_TYPE_UNKNOWN;
+            }
+            break;
+        }
+        case 64: type = LLM_TYPE_30B; break;
         // Add additional layer/vocab/etc checks here for other model sizes
         default: type = LLM_TYPE_UNKNOWN;
     }

src/models/jina-bert-v3.cpp

@@ -19,16 +19,6 @@ void llama_model_jina_bert_v3::load_arch_tensors(llama_model_loader &) {
     tok_embd     = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD,  "weight"), {n_embd, n_vocab}, 0);
     type_embd    = create_tensor(tn(LLM_TENSOR_TOKEN_TYPES, "weight"), {n_embd, n_token_types}, TENSOR_NOT_REQUIRED);
 
-    if (arch == LLM_ARCH_BERT) {
-        pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD,    "weight"), {n_embd, n_ctx_train}, 0);
-
-        cls   = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
-        cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"),   {n_embd},         TENSOR_NOT_REQUIRED);
-
-        cls_out   = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
-        cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"),   {hparams.n_cls_out},         TENSOR_NOT_REQUIRED);
-    }
-
     tok_norm   = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "weight", 0), {n_embd}, 0);
     tok_norm_b = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "bias",   0), {n_embd}, 0);
 
@@ -43,20 +33,10 @@ void llama_model_jina_bert_v3::load_arch_tensors(llama_model_loader &) {
         layer.attn_out_norm   = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "weight", i), {n_embd}, 0);
         layer.attn_out_norm_b = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "bias", i),   {n_embd}, 0);
 
-        if (hparams.moe_every_n_layers > 0 && i % hparams.moe_every_n_layers == 1) {
-            layer.ffn_up_exps   = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS,   "weight", i), {  n_embd, n_ff,   n_expert}, 0);
-            layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {  n_ff,   n_embd, n_expert}, 0);
-            layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP,   "weight", i), {n_embd, n_expert}, 0);
-        } else {
-            layer.ffn_up     = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", i), {n_embd, n_ff}, 0);
-            layer.ffn_up_b   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "bias", i),   {n_ff}, TENSOR_NOT_REQUIRED);
-            layer.ffn_down   = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
-            layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i),   {n_embd}, TENSOR_NOT_REQUIRED);
-
-            if (arch == LLM_ARCH_NOMIC_BERT) {
-                layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
-            }
-        }
+        layer.ffn_up     = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", i), {n_embd, n_ff}, 0);
+        layer.ffn_up_b   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "bias", i),   {n_ff}, TENSOR_NOT_REQUIRED);
+        layer.ffn_down   = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
+        layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i),   {n_embd}, TENSOR_NOT_REQUIRED);
 
         layer.layer_out_norm   = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, 0);
         layer.layer_out_norm_b = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "bias", i),   {n_embd}, 0);

src/models/nomic-bert-moe.cpp

@@ -4,12 +4,10 @@ void llama_model_nomic_bert_moe::load_arch_hparams(llama_model_loader & ml) {
     ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS,    hparams.f_norm_eps);
     ml.get_key(LLM_KV_MOE_EVERY_N_LAYERS,         hparams.moe_every_n_layers, 0);
 
-    if (hparams.n_layer() == 12 && hparams.n_embd == 768) {
-        if (arch == LLM_ARCH_NOMIC_BERT) {
-            type = LLM_TYPE_137M;
-        } else if (arch == LLM_ARCH_NOMIC_BERT_MOE && hparams.moe_every_n_layers == 2) {
-            type = LLM_TYPE_475M;
-        }
+    switch (hparams.n_layer()) {
+        case 12:
+            type = LLM_TYPE_475M; break;
+        default: type = LLM_TYPE_UNKNOWN;
     }
 }
 
@@ -22,16 +20,6 @@ void llama_model_nomic_bert_moe::load_arch_tensors(llama_model_loader &) {
     tok_embd     = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD,  "weight"), {n_embd, n_vocab}, 0);
     type_embd    = create_tensor(tn(LLM_TENSOR_TOKEN_TYPES, "weight"), {n_embd, n_token_types}, TENSOR_NOT_REQUIRED);
 
-    if (arch == LLM_ARCH_BERT) {
-        pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD,    "weight"), {n_embd, n_ctx_train}, 0);
-
-        cls   = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
-        cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"),   {n_embd},         TENSOR_NOT_REQUIRED);
-
-        cls_out   = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
-        cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"),   {hparams.n_cls_out},         TENSOR_NOT_REQUIRED);
-    }
-
     tok_norm   = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "weight", 0), {n_embd}, 0);
     tok_norm_b = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "bias",   0), {n_embd}, 0);
 
@@ -55,10 +43,6 @@ void llama_model_nomic_bert_moe::load_arch_tensors(llama_model_loader &) {
             layer.ffn_up_b   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "bias", i),   {n_ff}, TENSOR_NOT_REQUIRED);
             layer.ffn_down   = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
             layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i),   {n_embd}, TENSOR_NOT_REQUIRED);
-
-            if (arch == LLM_ARCH_NOMIC_BERT) {
-                layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
-            }
         }
 
         layer.layer_out_norm   = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, 0);

src/models/nomic-bert.cpp

@@ -1,15 +1,12 @@
 #include "models.h"
 
 void llama_model_nomic_bert::load_arch_hparams(llama_model_loader & ml) {
-    ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS,    hparams.f_norm_eps);
-    ml.get_key(LLM_KV_MOE_EVERY_N_LAYERS,         hparams.moe_every_n_layers, 0);
+    ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS, hparams.f_norm_eps);
 
-    if (hparams.n_layer() == 12 && hparams.n_embd == 768) {
-        if (arch == LLM_ARCH_NOMIC_BERT) {
-            type = LLM_TYPE_137M;
-        } else if (arch == LLM_ARCH_NOMIC_BERT_MOE && hparams.moe_every_n_layers == 2) {
-            type = LLM_TYPE_475M;
-        }
+    switch (hparams.n_layer()) {
+        case 12:
+            type = LLM_TYPE_137M; break;
+        default: type = LLM_TYPE_UNKNOWN;
     }
 }
 
@@ -22,16 +19,6 @@ void llama_model_nomic_bert::load_arch_tensors(llama_model_loader &) {
     tok_embd     = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD,  "weight"), {n_embd, n_vocab}, 0);
     type_embd    = create_tensor(tn(LLM_TENSOR_TOKEN_TYPES, "weight"), {n_embd, n_token_types}, TENSOR_NOT_REQUIRED);
 
-    if (arch == LLM_ARCH_BERT) {
-        pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD,    "weight"), {n_embd, n_ctx_train}, 0);
-
-        cls   = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
-        cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"),   {n_embd},         TENSOR_NOT_REQUIRED);
-
-        cls_out   = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
-        cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"),   {hparams.n_cls_out},         TENSOR_NOT_REQUIRED);
-    }
-
     tok_norm   = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "weight", 0), {n_embd}, 0);
     tok_norm_b = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "bias",   0), {n_embd}, 0);
 
@@ -46,20 +33,12 @@ void llama_model_nomic_bert::load_arch_tensors(llama_model_loader &) {
         layer.attn_out_norm   = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "weight", i), {n_embd}, 0);
         layer.attn_out_norm_b = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "bias", i),   {n_embd}, 0);
 
-        if (hparams.moe_every_n_layers > 0 && i % hparams.moe_every_n_layers == 1) {
-            layer.ffn_up_exps   = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS,   "weight", i), {  n_embd, n_ff,   n_expert}, 0);
-            layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), {  n_ff,   n_embd, n_expert}, 0);
-            layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP,   "weight", i), {n_embd, n_expert}, 0);
-        } else {
-            layer.ffn_up     = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", i), {n_embd, n_ff}, 0);
-            layer.ffn_up_b   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "bias", i),   {n_ff}, TENSOR_NOT_REQUIRED);
-            layer.ffn_down   = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
-            layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i),   {n_embd}, TENSOR_NOT_REQUIRED);
+        layer.ffn_up     = create_tensor(tn(LLM_TENSOR_FFN_UP,   "weight", i), {n_embd, n_ff}, 0);
+        layer.ffn_up_b   = create_tensor(tn(LLM_TENSOR_FFN_UP,   "bias", i),   {n_ff}, TENSOR_NOT_REQUIRED);
+        layer.ffn_down   = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
+        layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i),   {n_embd}, TENSOR_NOT_REQUIRED);
 
-            if (arch == LLM_ARCH_NOMIC_BERT) {
-                layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
-            }
-        }
+        layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
 
         layer.layer_out_norm   = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, 0);
         layer.layer_out_norm_b = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "bias", i),   {n_embd}, 0);

tests/test-backend-ops.cpp

@@ -462,17 +462,9 @@ static std::string var_to_str(ggml_scale_mode mode) {
 #define VARS_TO_STR16(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p) VAR_TO_STR(a) + "," + VARS_TO_STR15(b, c, d, e, f, g, h, i, j, k, l, m, n, o, p)
 #define VARS_TO_STR17(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q) VAR_TO_STR(a) + "," + VARS_TO_STR16(b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q)
 
-#ifdef GGML_USE_SYCL
-static bool inline _isinf(float f) {
-    return (*(uint32_t *)&f & 0x7fffffff) == 0x7f800000;
-}
-#else
-static bool inline _isinf(float f) { return std::isinf(f); }
-#endif
-
 // accept FLT_MAX as infinity
 static bool isinf_or_max(float f) {
-    return _isinf(f) || f == FLT_MAX || f == -FLT_MAX;
+    return std::isinf(f) || f == FLT_MAX || f == -FLT_MAX;
 }
 
 static bool ggml_is_view_op(enum ggml_op op) {
@@ -3336,6 +3328,14 @@ struct test_bin_bcast : public test_case {
         return op == ggml_div;
     }
 
+    double max_nmse_err() override {
+        if (op == ggml_add && type == GGML_TYPE_F16 && nf > 1) {
+            // Fused ADDs can keep FP32 intermediates while the CPU rounds each ADD to FP16.
+            return 1e-6;
+        }
+        return test_case::max_nmse_err();
+    }
+
     double max_maa_err() override {
         return op == ggml_add ? 1e-4 : 1e-3;
     }
@@ -4831,51 +4831,6 @@ struct test_mul_mat : public test_case {
     }
 };
 
-#define P 1.0f
-#define N -1.0f
-
-// constant Hadamard matrix via Paley I construction
-static constexpr float H12[12][12] = {
-    { P, P, P, P, P, P, P, P, P, P, P, P },
-    { P, N, P, N, P, P, P, N, N, N, P, N },
-    { P, N, N, P, N, P, P, P, N, N, N, P },
-    { P, P, N, N, P, N, P, P, P, N, N, N },
-    { P, N, P, N, N, P, N, P, P, P, N, N },
-    { P, N, N, P, N, N, P, N, P, P, P, N },
-    { P, N, N, N, P, N, N, P, N, P, P, P },
-    { P, P, N, N, N, P, N, N, P, N, P, P },
-    { P, P, P, N, N, N, P, N, N, P, N, P },
-    { P, P, P, P, N, N, N, P, N, N, P, N },
-    { P, N, P, P, P, N, N, N, P, N, N, P },
-    { P, P, N, P, P, P, N, N, N, P, N, N }
-};
-
-static constexpr float H20[20][20] = {
-    { P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P },
-    { P, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N },
-    { P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P },
-    { P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P },
-    { P, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N },
-    { P, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N },
-    { P, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N },
-    { P, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N },
-    { P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P },
-    { P, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N },
-    { P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P },
-    { P, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N },
-    { P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P },
-    { P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P },
-    { P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P },
-    { P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P },
-    { P, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N },
-    { P, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N },
-    { P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P },
-    { P, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N }
-};
-
-#undef P
-#undef N
-
 // GGML_HINT_SRC0_IS_HADAMARD
 struct test_mul_mat_hadamard : public test_mul_mat {
     test_mul_mat_hadamard(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32,
@@ -4900,58 +4855,20 @@ struct test_mul_mat_hadamard : public test_mul_mat {
     void initialize_tensors(ggml_context * ctx) override {
         for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
             if (strcmp(t->name, "a") == 0) {
-                const int64_t      n_cols = t->ne[0];
-                const int64_t      n_rows = ggml_nrows(t);
+                const int64_t n_cols = t->ne[0];
+                const int64_t n_rows = ggml_nrows(t);
                 std::vector<float> data(n_cols * n_rows);
-                float              scale = 1.0f / sqrtf((float) n_cols);
-
-                auto is_pow2 = [](const int64_t a) {
-                    return (a > 0) && ((a & (a - 1)) == 0);
-                };
-#ifdef GGML_USE_SYCL
-                const bool is_kronecker =
-                    ((n_cols % 12 == 0) && is_pow2(n_cols / 12)) || ((n_cols % 20 == 0) && is_pow2(n_cols / 20));
-#else
-                const bool is_kronecker = false;
-#endif
-                if (is_kronecker) {
-                    const int64_t B = (n_cols % 12 == 0 && is_pow2(n_cols / 12)) ? 12 : 20;
-                    for (int64_t r = 0; r < n_rows; r++) {
-                        float *       row_data = data.data() + r * n_cols;
-                        const int64_t r_mod    = r % n_cols;
-                        const int64_t r_b      = r_mod / B;
-                        const int64_t r_m      = r_mod % B;
-
-                        for (int64_t i = 0; i < n_cols; i++) {
-                            const int64_t c_b = i / B;
-                            const int64_t c_m = i % B;
-
-                            int     pop = 0;
-                            int64_t val = r_b & c_b;
-                            while (val) {
-                                pop += (val & 1);
-                                val >>= 1;
-                            }
-                            const float sign_m = (pop % 2 == 0) ? 1.0f : -1.0f;
-                            const float sign_b = (B == 12) ? H12[c_m][r_m] : H20[c_m][r_m];
-
-                            row_data[i] = scale * sign_b * sign_m;
-                        }
-                    }
-                }
-
-                else if (is_pow2(n_cols)) {
-                    for (int64_t r = 0; r < n_rows; r++) {
-                        float * row_data = data.data() + r * n_cols;
-                        for (int64_t i = 0; i < n_cols; i++) {
-                            int     pop_cnt = 0;
-                            int64_t val     = r & i;
-                            while (val) {
-                                pop_cnt += (val & 1);
-                                val >>= 1;
-                            }
-                            row_data[i] = (pop_cnt % 2 == 0) ? scale : -scale;
+                float scale = 1.0f / sqrtf((float)n_cols);
+                for (int64_t r = 0; r < n_rows; r++) {
+                    float * row_data = data.data() + r * n_cols;
+                    for (int64_t i = 0; i < n_cols; i++) {
+                        int pop = 0;
+                        int64_t val = r & i;
+                        while (val) {
+                            pop += (val & 1);
+                            val >>= 1;
                         }
+                        row_data[i] = (pop % 2 == 0) ? scale : -scale;
                     }
                 }
                 ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(float));
@@ -9716,16 +9633,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 512, 256)); // many rows
     test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 32, 1, 32)); // too small (N<64)
     test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1024, 1, 1024)); // too big (N>512)
-#ifdef GGML_USE_SYCL
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 1, 384));    // m=12 (N=384)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 32, 384));   // m=12 (batch)
-    test_cases.emplace_back(
-        new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 4, 384, { 2, 3 }));             // m=12 (multi-dim)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 768, 1, 768));   // m=12 (N=768)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 1, 640));   // m=20 (N=640)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 32, 640));  // m=20 (batch)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1280, 1, 1280));  // m=20 (N=1280)
-#endif
+
 #if 0
     // > 4GB A matrix. Too slow to be enabled by default.
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16,  900000,  3, 2592, {1, 1}, {1, 1}));
@@ -10282,6 +10190,12 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
         test_cases.emplace_back(new test_rope(type, {128, 32, 2, 1}, 32, GGML_ROPE_TYPE_NEOX, 512, 1.4245f, 0.7465f, 1.4245f, false, 0, true, true, 32));
     }
 
+    // Real-model RoPE: F32 forward, packed Q, 512-token prefill.
+    test_cases.emplace_back(new test_rope(GGML_TYPE_F32, {256,  8, 512, 1},  64, GGML_ROPE_TYPE_IMROPE, 512, 1.0f, 0.0f, 1.0f, false, 0, true)); // qwen3.5 0.8B
+    test_cases.emplace_back(new test_rope(GGML_TYPE_F32, {256, 16, 512, 1},  64, GGML_ROPE_TYPE_IMROPE, 512, 1.0f, 0.0f, 1.0f, false, 0, true)); // qwen3.5 4B
+    test_cases.emplace_back(new test_rope(GGML_TYPE_F32, {256,  8, 512, 1}, 256, GGML_ROPE_TYPE_NEOX,   512, 1.0f, 0.0f, 1.0f, false, 0, true)); // gemma4 E2B sliding
+    test_cases.emplace_back(new test_rope(GGML_TYPE_F32, {512,  8, 512, 1}, 128, GGML_ROPE_TYPE_NEOX,   512, 1.0f, 0.0f, 1.0f, true,  0, true)); // gemma4 E4B global
+
     for (int v : { 0, 1, 2, 3 }) {
         for (int dim : { 0, 1, 2, 3, }) {
             test_cases.emplace_back(new test_concat(GGML_TYPE_F32, {11, 12, 13, 14}, 7, dim, v));
@@ -11005,16 +10919,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
     test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 128, 2048, 128));
     test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 2048, 256));
     test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 512, 2048, 512));
-#ifdef GGML_USE_SYCL
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 1, 384));    // m=12 (N=384)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 32, 384));   // m=12 (batch)
-    test_cases.emplace_back(
-        new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 4, 384, { 2, 3 }));             // m=12 (multi-dim)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 768, 1, 768));   // m=12 (N=768)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 1, 640));   // m=20 (N=640)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 32, 640));  // m=20 (batch)
-    test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1280, 1, 1280));  // m=20 (N=1280)
-#endif
+
     test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 64, 64, 4, 4 }, { 32, 64, 4, 4 }));
     test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 128, 128, 4, 2 }, { 32, 128, 4, 2 }));
     // qwen3next with CHUNK_SIZE 64
@@ -11178,6 +11083,12 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
         }
     }
 
+    // Real-model RoPE: F32 forward, packed Q, 512-token prefill.
+    test_cases.emplace_back(new test_rope(GGML_TYPE_F32, {256,  8, 512, 1},  64, GGML_ROPE_TYPE_IMROPE, 512, 1.0f, 0.0f, 1.0f, false, 0, true)); // qwen3.5 0.8B
+    test_cases.emplace_back(new test_rope(GGML_TYPE_F32, {256, 16, 512, 1},  64, GGML_ROPE_TYPE_IMROPE, 512, 1.0f, 0.0f, 1.0f, false, 0, true)); // qwen3.5 4B
+    test_cases.emplace_back(new test_rope(GGML_TYPE_F32, {256,  8, 512, 1}, 256, GGML_ROPE_TYPE_NEOX,   512, 1.0f, 0.0f, 1.0f, false, 0, true)); // gemma4 E2B sliding
+    test_cases.emplace_back(new test_rope(GGML_TYPE_F32, {512,  8, 512, 1}, 128, GGML_ROPE_TYPE_NEOX,   512, 1.0f, 0.0f, 1.0f, true,  0, true)); // gemma4 E4B global
+
     std::vector<std::array<int64_t, 4>> reduce_rows_cases = {
         { 8192, 1,    1, 1 },
         { 8192, 8192, 1, 1 },

tools/server/tests/requirements.txt

@@ -2,7 +2,7 @@ aiohttp~=3.9.3
 pytest~=8.3.3
 pytest-xdist~=3.6
 filelock~=3.16
-numpy~=1.26.4
+numpy~=2.2.6
 openai~=2.14.0
 prometheus-client~=0.20.0
 requests~=2.32.3