hexagon: tensor repacking rework and support for tiled and flat quantized and fp weights (#30635)

* hex-repack: Remove need_repack hack and streamline the rules -- breaks GET_ROWS tests for K-Quants, will fix in the next round

Add support for GET_ROWS with tiled Q5_K
Add support for raw k-quant tensors

* hex-get-rows: vectorize raw k-quant ukernels

* hex-binary: refactor dma-loop thread functions (no need for macros, etc)

* hex-quant: move/cleanup flat and tiled dequant ukernels

* hex-repack: add support for repacking/tiling F16/F32/BF16 2D weights

* hex-get-rows: remove redundant cast

* hex-repack: proper filtering for BF16 that were leaking through during the tests

* hex-repack: streamling tiling rules (broken MM for now)

* hex-repack: fix mmid after tiling updates

* hex-repack: rename repacked to tiled in all flags/funcs/kernels/vars

* hex-mm: fix valid_rows calc in various kernels

* hex-mm: remove unused hvx_mm_nx_2d_flat

* hex-mm: cleanup and optimize hvx-tiled vec_dot_f16

* hex-mm: another pass at hvx-tiled f16 matmul

* hex-get-rows: optimize small src1

* hex-get-rows: one more pass at recovering get-rows perf regression

* hex-repack: export GGML_HEXAGON_RT_REPACK via run.py

* hex-mm-tiled: fix accuracy regression with hmx-tiled-f16 matmul

* hex-mm: fold tiled_f16 checks into the general flow

* hex-get-rows: remove dead checks for buffer_is_hexagon and tiling during supports_op

* hex-quant: proper tail handling in new dequant ukernels

* hex-get-rows: remove unused code and cleanup support checks

* hex-mm: cleanup ne1 padding/tiling checks

* hex-mm: cleanup fusion checks for tiled tensors

* hex-repack: add a little can_repack helper to reduce duplication

* hex-repack: never repack host buffers

* hex-repack: factor out repack functions into separate .h and .cpp

* hex-tensor: cleanup tensor ifaces for tiled/padded sizes

* hex-repack: update k-quant tiling comments and cleanup #defines

* hex-dcache: factor out l2/dcache dirty tracking and flushing

* hex-tensor: update docs and remove #defs duplication

* hex-tensor: remove unused/incorrect get_data interface

* hex-get-rows: cleanup vtcm layout build

* tests: disable huge MM tests

* hex-dcache: add missing header

docs/backend/snapdragon/developer.md

@@ -134,6 +134,18 @@ Writing high-performance operators for Hexagon requires following specific guide
 - **Distinction**: `octx->src[N]` pointers *can* be NULL by design and must be checked when optional.
   Examples include attention masks, optional bias or weights in fused kernels, and frequency factors.
 
+### `supports_op` Independence from Buffer and Repack State
+
+- `supports_op` (and op-specific checkers such as `supported_mul_mat()`, `supported_get_rows()`) runs during GGML graph scheduling **before** buffers are allocated and before tensor data is populated.
+- **Never check buffer usage or repacking state in `supports_op`**:
+  - `t->buffer` may be NULL or uninitialized.
+  - `t->data` may be NULL.
+  - `buffer->usage` is unknown at scheduling time.
+  - `is_tiled` or `extra->flags & GGML_HEXAGON_TENSOR_TILED` is not set yet.
+- Operator support decisions must depend strictly on static tensor properties: operand types (`t->type`), shapes (`t->ne`), strides (`t->nb`), and operation parameters.
+- Repacking (converting weights to tiled layout) is an internal optimization applied during `set_tensor()` or rt-repack. Non-weight tensors or test harness tensors with the same type remain in flat layout.
+- Therefore, any supported type (such as BF16, F16, quantized types) must be supported both when repacked into tiles and when stored flat in memory.
+
 ### Multiline Macro Formatting
 
 - Keep trailing backslashes in multiline `#define` macros cleanly aligned to a consistent column.
@@ -190,7 +202,7 @@ sessions.
 Common partitioning logic is factored into reusable inline helpers in
 [`htp-tensor.h`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h):
 
-1. [`htp_tensor_mdev_rows_per_chunk`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L71):
+1. [`htp_tensor_mdev_rows_per_chunk`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L104):
    Determines the minimum number of rows per chunk so that the chunk byte size is a multiple of 128 bytes:
 
    ```
@@ -201,7 +213,7 @@ Common partitioning logic is factored into reusable inline helpers in
    Returns `false` if the tensor cannot be safely row-partitioned (such as unaligned base pointer, permuted layout,
    or non-128-byte aligned outer strides).
 
-2. [`htp_tensor_mdev_partition`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L98):
+2. [`htp_tensor_mdev_partition`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L131):
    Calculates the per-device work range `struct htp_tensor_mdev_range { uint32_t start; uint32_t count; }` given
    `total_units`, `units_per_chunk`, `mdev_idx`, `mdev_count`, and the precomputed `mdev_count_div`.
    Handles chunk distribution across devices, assigns remainder units to the last device, and automatically triggers
@@ -240,11 +252,11 @@ if (nrows == 0) {
 For flat element-wise operations that write directly to DDR:
 - Partition total linear elements N = ne0 * ne1 * ne2 * ne3 in 128-byte cache line chunks (`elems_per_line = (elem_size == 4) ? 32 : 64`).
 - Requires strict 1D contiguity:
-  [`htp_tensor_is_contiguous(dst, elem_size)`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L32)
+  [`htp_tensor_is_contiguous(dst, elem_size)`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L65)
   and 128-byte aligned destination pointer
-  [`htp_tensor_mdev_data_aligned(dst)`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L51).
+  [`htp_tensor_mdev_data_aligned(dst)`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L84).
 - If contiguous and aligned, pass `elems_per_line` to
-  [`htp_tensor_mdev_partition`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L98);
+  [`htp_tensor_mdev_partition`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h#L131);
   otherwise pass 0 to trigger Device 0 fallback.
 
 ### Single-Device Fallback (Device 0)
@@ -312,7 +324,7 @@ Multi-device execution synchronizes worker sessions through atomic fence slots a
   (Output Slice 0)                             (Output Slice 1)
        |                                            |
   Tensor Cache Flush                           Tensor Cache Flush
-  (htp_tensor_flush_all)                       (htp_tensor_flush_all)
+  (htp_dcache_flush_dirty_tensors)             (htp_dcache_flush_dirty_tensors)
        |                                            |
   Post-Op/Batch Barrier ---------------------- Post-Op/Batch Barrier
   (htp_mdev_group_barrier)                     (htp_mdev_group_barrier)
@@ -350,8 +362,8 @@ Multi-device execution synchronizes worker sessions through atomic fence slots a
 ### Tensor Cache Flush and Pipeline Completion
 
 - In the kernel, ensure all pushed DMA operations have been popped in strict FIFO order to drain the queue.
-- Use [`htp_tensor_flush_all()`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h) to flush specific dirty tensors back to DDR:
-  - [`htp_tensor_flush_all()`](../../../ggml/src/ggml-hexagon/htp/htp-tensor.h) flushes modified tensor address ranges, or the
+- Use [`htp_dcache_flush_dirty_tensors()`](../../../ggml/src/ggml-hexagon/htp/htp-dcache.h) to flush specific dirty tensors back to DDR:
+  - [`htp_dcache_flush_dirty_tensors()`](../../../ggml/src/ggml-hexagon/htp/htp-dcache.h) flushes modified tensor address ranges, or the
     full D-cache when their total size exceeds the flush threshold, ensuring peer devices and the host CPU observe consistent
     data in DDR.
 - Never signal completion before all DMA transfers are drained and dirty tensor flushes have completed.

ggml/src/ggml-hexagon/CMakeLists.txt

@@ -50,6 +50,8 @@ endif()
 set(TARGET_NAME ggml-hexagon)
 ggml_add_backend_library(${TARGET_NAME}
     ggml-hexagon.cpp
+    htp-repack.cpp
+    htp-repack.h
     htp-drv.cpp
     htp-drv.h
     libdl.h

ggml/src/ggml-hexagon/htp-opnode.h

@@ -326,10 +326,12 @@ struct htp_opformat {
             const auto * kparams = (const struct htp_mm_kernel_params *) node.kernel_params;
             const char * path = "unknown";
             int32_t type = kparams->kernel_type;
-            if (type == HTP_MM_KERNEL_HMX_2D || type == HTP_MM_KERNEL_HMX_F16_BATCHED) {
-                path = "hmx-tiled";
+            if (type == HTP_MM_KERNEL_HMX_2D || type == HTP_MM_KERNEL_HMX_BATCHED) {
+                path = (kparams->tile_size > 0) ? "hmx-tiled" : "hmx-flat";
             } else if (type == HTP_MM_KERNEL_HVX_F16_F16_VTCM || type == HTP_MM_KERNEL_HVX_F32_F32_VTCM ||
-                       type == HTP_MM_KERNEL_HVX_QUANT_ROW    || type == HTP_MM_KERNEL_HVX_QUANT_BLOCK) {
+                       type == HTP_MM_KERNEL_HVX_BF16_F32_VTCM) {
+                path = "hvx-flat";
+            } else if (type == HTP_MM_KERNEL_HVX_QUANT_ROW || type == HTP_MM_KERNEL_HVX_QUANT_BLOCK) {
                 path = "hvx-tiled";
             }
             snprintf(str, max_size, "%s vtcm %d", path, (int) kparams->vtcm_size);

ggml/src/ggml-hexagon/htp-repack.cpp

@@ -0,0 +1,1000 @@
+#include "htp-repack.h"
+
+#define GGML_COMMON_IMPL_CPP
+#include "ggml-common.h"
+#include "ggml-impl.h"
+#include "ggml-quants.h"
+
+#include <algorithm>
+#include <cmath>
+#include <cstring>
+#include <initializer_list>
+
+// --- 4-bit block-32 nibble helpers (shared by Q4_0, IQ4_NL, Q4_1, MXFP4) ---
+
+static inline void unpack_nibbles_32(uint8_t * dst, const uint8_t * qs) {
+    for (unsigned int i = 0; i < 16; ++i) {
+        dst[i + 0]  = qs[i] & 0x0F;
+        dst[i + 16] = qs[i] >> 4;
+    }
+}
+
+static inline void pack_nibbles_32(uint8_t * qs, const uint8_t * src) {
+    for (unsigned int i = 0; i < 16; ++i) {
+        qs[i] = src[i + 0] | (src[i + 16] << 4);
+    }
+}
+
+static inline void transpose_nibbles_to_tile(uint8_t * tile_dst, const uint8_t tile_quants[32][32]) {
+    for (int cp = 0; cp < 16; cp++) {
+        for (int row = 0; row < 32; row++) {
+            tile_dst[cp * 32 + row] = (tile_quants[row][2 * cp + 1] << 4) | tile_quants[row][2 * cp];
+        }
+    }
+}
+
+static inline void transpose_tile_to_nibbles(uint8_t tile_quants[32][32], const uint8_t * tile_src) {
+    for (int cp = 0; cp < 16; cp++) {
+        for (int row = 0; row < 32; row++) {
+            const uint8_t val = tile_src[cp * 32 + row];
+            tile_quants[row][2 * cp + 0] = val & 0x0F;
+            tile_quants[row][2 * cp + 1] = val >> 4;
+        }
+    }
+}
+
+// --- Generic 32x32 tile grid walkers for block-32 and float types ---
+
+template <typename SrcElemT, int ELEMS_PER_KTILE, typename PackTileFn>
+static void repack_to_tiled_grid_32(ggml_tensor * t, const void * data, size_t offset, size_t size,
+                                    size_t tile_size, PackTileFn && pack_tile) {
+    const SrcElemT * src_matrix = (const SrcElemT *) data;
+    const int64_t ne0 = t->ne[0];
+    const int64_t ne1 = t->ne[1];
+    const int64_t ne2 = t->ne[2];
+    const int64_t ne3 = t->ne[3];
+    const int n_col_tiles = hex_round_up(ne1, 32) / 32;
+    const int n_k_tiles   = hex_round_up(ne0, 32) / 32;
+    const size_t matrix_size = (size_t) n_col_tiles * n_k_tiles * tile_size;
+
+    const size_t slice_size = (size_t) ne1 * ggml_row_size(t->type, ne0);
+    const int64_t start_slice = offset / slice_size;
+    int64_t end_slice = (offset + size + slice_size - 1) / slice_size;
+    if (end_slice > ne2 * ne3) {
+        end_slice = ne2 * ne3;
+    }
+
+    const int64_t row_stride = (ne0 / (32 / ELEMS_PER_KTILE));
+
+    for (int64_t slice_idx = start_slice; slice_idx < end_slice; slice_idx++) {
+        const SrcElemT * src_slice = src_matrix + (slice_idx - start_slice) * (ne1 * row_stride);
+        uint8_t * matrix_dst = (uint8_t *) t->data + slice_idx * matrix_size;
+
+        for (int ct = 0; ct < n_col_tiles; ct++) {
+            for (int kt = 0; kt < n_k_tiles; kt++) {
+                uint8_t * tile_dst = matrix_dst + ((size_t) ct * n_k_tiles + kt) * tile_size;
+                pack_tile(tile_dst, src_slice, ne0, ne1, row_stride, ct, kt);
+            }
+        }
+    }
+}
+
+template <typename DstElemT, int ELEMS_PER_KTILE, typename UnpackTileFn>
+static void unpack_from_tiled_grid_32(void * data, const ggml_tensor * t, size_t offset, size_t size,
+                                      size_t tile_size, UnpackTileFn && unpack_tile) {
+    uint8_t * dst_bytes = (uint8_t *) data;
+    const int64_t ne0 = t->ne[0];
+    const int64_t ne1 = t->ne[1];
+    const int64_t ne2 = t->ne[2];
+    const int64_t ne3 = t->ne[3];
+    const int n_col_tiles = hex_round_up(ne1, 32) / 32;
+    const int n_k_tiles   = hex_round_up(ne0, 32) / 32;
+    const size_t matrix_size = (size_t) n_col_tiles * n_k_tiles * tile_size;
+
+    const size_t row_size_bytes = ggml_row_size(t->type, ne0);
+    const size_t slice_size     = (size_t) ne1 * row_size_bytes;
+    const int64_t start_slice   = offset / slice_size;
+    int64_t end_slice           = (offset + size + slice_size - 1) / slice_size;
+    if (end_slice > ne2 * ne3) {
+        end_slice = ne2 * ne3;
+    }
+
+    const int64_t row_stride = (ne0 / (32 / ELEMS_PER_KTILE));
+
+    for (int64_t slice_idx = start_slice; slice_idx < end_slice; slice_idx++) {
+        const size_t cur_start_byte     = (std::max)(offset, (size_t) slice_idx * slice_size);
+        const size_t cur_end_byte       = (std::min)(offset + size, (size_t) (slice_idx + 1) * slice_size);
+        const size_t slice_offset_start = cur_start_byte - (size_t) slice_idx * slice_size;
+        const size_t slice_offset_end   = cur_end_byte - (size_t) slice_idx * slice_size;
+
+        const int64_t start_row = slice_offset_start / row_size_bytes;
+        const int64_t end_row   = (std::min)((int64_t) ((slice_offset_end + row_size_bytes - 1) / row_size_bytes), ne1);
+
+        const int start_ct = (int) (start_row / 32);
+        const int end_ct   = (std::min)((int) ((end_row + 31) / 32), n_col_tiles);
+
+        DstElemT * dst_slice = (DstElemT *) (dst_bytes + (cur_start_byte - offset));
+        const uint8_t * matrix_src = (const uint8_t *) t->data + slice_idx * matrix_size;
+
+        for (int ct = start_ct; ct < end_ct; ct++) {
+            for (int kt = 0; kt < n_k_tiles; kt++) {
+                const uint8_t * tile_src = matrix_src + ((size_t) ct * n_k_tiles + kt) * tile_size;
+                unpack_tile(dst_slice, tile_src, ne0, start_row, end_row, row_stride, ct, kt);
+            }
+        }
+    }
+}
+
+// --- Block-32 Quantized Types (Q4_0, IQ4_NL, Q4_1, Q8_0, MXFP4) ---
+
+static void repack_q4_0_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_to_tiled_grid_32<block_q4_0, 1>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q4_0,
+        [](uint8_t * tile_dst, const block_q4_0 * src_slice, int64_t ne0, int64_t ne1, int64_t row_stride, int ct, int kt) {
+            uint8_t tile_quants[32][32];
+            ggml_half * scale_dst = (ggml_half *) (tile_dst + 512);
+            for (int row = 0; row < 32; row++) {
+                const int64_t r = ct * 32 + row;
+                if (r < ne1 && kt < ne0 / 32) {
+                    const block_q4_0 & b = src_slice[r * row_stride + kt];
+                    unpack_nibbles_32(tile_quants[row], b.qs);
+                    scale_dst[row] = b.d;
+                } else {
+                    memset(tile_quants[row], 8, 32);
+                    scale_dst[row] = 0;
+                }
+            }
+            transpose_nibbles_to_tile(tile_dst, tile_quants);
+        });
+}
+
+static void repack_tiled_q4_0(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_from_tiled_grid_32<block_q4_0, 1>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q4_0,
+        [](block_q4_0 * dst_slice, const uint8_t * tile_src, int64_t ne0, int64_t start_row, int64_t end_row, int64_t row_stride, int ct, int kt) {
+            uint8_t tile_quants[32][32];
+            transpose_tile_to_nibbles(tile_quants, tile_src);
+            const ggml_half * scale_src = (const ggml_half *) (tile_src + 512);
+            for (int row = 0; row < 32; row++) {
+                const int64_t r = ct * 32 + row;
+                if (r >= start_row && r < end_row && kt < ne0 / 32) {
+                    block_q4_0 & b = dst_slice[(r - start_row) * row_stride + kt];
+                    pack_nibbles_32(b.qs, tile_quants[row]);
+                    b.d = scale_src[row];
+                }
+            }
+        });
+}
+
+static void repack_q4_1_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_to_tiled_grid_32<block_q4_1, 1>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q4_1,
+        [](uint8_t * tile_dst, const block_q4_1 * src_slice, int64_t ne0, int64_t ne1, int64_t row_stride, int ct, int kt) {
+            uint8_t tile_quants[32][32];
+            ggml_half * scale_dst = (ggml_half *) (tile_dst + 512);
+            for (int row = 0; row < 32; row++) {
+                const int64_t r = ct * 32 + row;
+                if (r < ne1 && kt < ne0 / 32) {
+                    const block_q4_1 & b = src_slice[r * row_stride + kt];
+                    unpack_nibbles_32(tile_quants[row], b.qs);
+                    scale_dst[2 * row + 0] = b.d;
+                    scale_dst[2 * row + 1] = b.m;
+                } else {
+                    memset(tile_quants[row], 0, 32);
+                    scale_dst[2 * row + 0] = 0;
+                    scale_dst[2 * row + 1] = 0;
+                }
+            }
+            transpose_nibbles_to_tile(tile_dst, tile_quants);
+        });
+}
+
+static void repack_tiled_q4_1(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_from_tiled_grid_32<block_q4_1, 1>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q4_1,
+        [](block_q4_1 * dst_slice, const uint8_t * tile_src, int64_t ne0, int64_t start_row, int64_t end_row, int64_t row_stride, int ct, int kt) {
+            uint8_t tile_quants[32][32];
+            transpose_tile_to_nibbles(tile_quants, tile_src);
+            const ggml_half * scale_src = (const ggml_half *) (tile_src + 512);
+            for (int row = 0; row < 32; row++) {
+                const int64_t r = ct * 32 + row;
+                if (r >= start_row && r < end_row && kt < ne0 / 32) {
+                    block_q4_1 & b = dst_slice[(r - start_row) * row_stride + kt];
+                    pack_nibbles_32(b.qs, tile_quants[row]);
+                    b.d = scale_src[2 * row + 0];
+                    b.m = scale_src[2 * row + 1];
+                }
+            }
+        });
+}
+
+static void repack_mxfp4_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_to_tiled_grid_32<block_mxfp4, 1>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_MXFP4,
+        [](uint8_t * tile_dst, const block_mxfp4 * src_slice, int64_t ne0, int64_t ne1, int64_t row_stride, int ct, int kt) {
+            uint8_t tile_quants[32][32];
+            uint8_t * scale_dst = tile_dst + 512;
+            for (int row = 0; row < 32; row++) {
+                const int64_t r = ct * 32 + row;
+                if (r < ne1 && kt < ne0 / 32) {
+                    const block_mxfp4 & b = src_slice[r * row_stride + kt];
+                    unpack_nibbles_32(tile_quants[row], b.qs);
+                    scale_dst[row] = b.e;
+                } else {
+                    memset(tile_quants[row], 0, 32);
+                    scale_dst[row] = 0;
+                }
+            }
+            transpose_nibbles_to_tile(tile_dst, tile_quants);
+        });
+}
+
+static void repack_tiled_mxfp4(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_from_tiled_grid_32<block_mxfp4, 1>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_MXFP4,
+        [](block_mxfp4 * dst_slice, const uint8_t * tile_src, int64_t ne0, int64_t start_row, int64_t end_row, int64_t row_stride, int ct, int kt) {
+            uint8_t tile_quants[32][32];
+            transpose_tile_to_nibbles(tile_quants, tile_src);
+            const uint8_t * scale_src = tile_src + 512;
+            for (int row = 0; row < 32; row++) {
+                const int64_t r = ct * 32 + row;
+                if (r >= start_row && r < end_row && kt < ne0 / 32) {
+                    block_mxfp4 & b = dst_slice[(r - start_row) * row_stride + kt];
+                    pack_nibbles_32(b.qs, tile_quants[row]);
+                    b.e = scale_src[row];
+                }
+            }
+        });
+}
+
+static void repack_q8_0_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_to_tiled_grid_32<block_q8_0, 1>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q8_0,
+        [](uint8_t * tile_dst, const block_q8_0 * src_slice, int64_t ne0, int64_t ne1, int64_t row_stride, int ct, int kt) {
+            for (int cp = 0; cp < 16; cp++) {
+                const int col0 = cp * 2;
+                const int col1 = col0 + 1;
+                for (int row = 0; row < 32; row++) {
+                    const int64_t r = ct * 32 + row;
+                    const block_q8_0 * b = (r < ne1 && kt < ne0 / 32) ? &src_slice[r * row_stride + kt] : nullptr;
+                    tile_dst[cp * 64 + 2 * row + 0] = b ? b->qs[col0] : 0;
+                    tile_dst[cp * 64 + 2 * row + 1] = b ? b->qs[col1] : 0;
+                }
+            }
+            ggml_half * scale_dst = (ggml_half *) (tile_dst + 1024);
+            for (int row = 0; row < 32; row++) {
+                const int64_t r = ct * 32 + row;
+                scale_dst[row] = (r < ne1 && kt < ne0 / 32) ? src_slice[r * row_stride + kt].d : 0;
+            }
+        });
+}
+
+static void repack_tiled_q8_0(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_from_tiled_grid_32<block_q8_0, 1>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q8_0,
+        [](block_q8_0 * dst_slice, const uint8_t * tile_src, int64_t ne0, int64_t start_row, int64_t end_row, int64_t row_stride, int ct, int kt) {
+            for (int cp = 0; cp < 16; cp++) {
+                const int col0 = cp * 2;
+                const int col1 = col0 + 1;
+                for (int row = 0; row < 32; row++) {
+                    const int64_t r = ct * 32 + row;
+                    if (r >= start_row && r < end_row && kt < ne0 / 32) {
+                        block_q8_0 & b = dst_slice[(r - start_row) * row_stride + kt];
+                        b.qs[col0] = tile_src[cp * 64 + 2 * row + 0];
+                        b.qs[col1] = tile_src[cp * 64 + 2 * row + 1];
+                    }
+                }
+            }
+            const ggml_half * scale_src = (const ggml_half *) (tile_src + 1024);
+            for (int row = 0; row < 32; row++) {
+                const int64_t r = ct * 32 + row;
+                if (r >= start_row && r < end_row && kt < ne0 / 32) {
+                    dst_slice[(r - start_row) * row_stride + kt].d = scale_src[row];
+                }
+            }
+        });
+}
+
+// --- Float Types (F16, F32, BF16 -> f16_tiled HMX layout) ---
+
+template <typename SrcT, typename ToFp16Fn>
+static void repack_float_to_f16_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size, ToFp16Fn && to_fp16) {
+    repack_to_tiled_grid_32<SrcT, 32>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_F16,
+        [&](uint8_t * tile_bytes, const SrcT * src_slice, int64_t ne0, int64_t ne1, int64_t row_stride, int ct, int kt) {
+            uint16_t * tile_dst = (uint16_t *) tile_bytes;
+            for (int cp = 0; cp < 16; cp++) {
+                const int64_t c0 = kt * 32 + cp * 2;
+                const int64_t c1 = c0 + 1;
+                for (int row = 0; row < 32; row++) {
+                    const int64_t r = ct * 32 + row;
+                    const SrcT * r_ptr = (r < ne1) ? &src_slice[r * row_stride] : nullptr;
+                    tile_dst[cp * 64 + row * 2 + 0] = (r_ptr && c0 < ne0) ? to_fp16(r_ptr[c0]) : 0;
+                    tile_dst[cp * 64 + row * 2 + 1] = (r_ptr && c1 < ne0) ? to_fp16(r_ptr[c1]) : 0;
+                }
+            }
+        });
+}
+
+template <typename DstT, typename FromFp16Fn>
+static void unpack_f16_tiled_to_float(void * data, const ggml_tensor * t, size_t offset, size_t size, FromFp16Fn && from_fp16) {
+    unpack_from_tiled_grid_32<DstT, 32>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_F16,
+        [&](DstT * dst_slice, const uint8_t * tile_bytes, int64_t ne0, int64_t start_row, int64_t end_row, int64_t row_stride, int ct, int kt) {
+            const uint16_t * tile_src = (const uint16_t *) tile_bytes;
+            for (int cp = 0; cp < 16; cp++) {
+                const int64_t c0 = kt * 32 + cp * 2;
+                const int64_t c1 = c0 + 1;
+                for (int row = 0; row < 32; row++) {
+                    const int64_t r = ct * 32 + row;
+                    if (r >= start_row && r < end_row) {
+                        DstT * r_ptr = dst_slice + (r - start_row) * row_stride;
+                        if (c0 < ne0) r_ptr[c0] = from_fp16(tile_src[cp * 64 + row * 2 + 0]);
+                        if (c1 < ne0) r_ptr[c1] = from_fp16(tile_src[cp * 64 + row * 2 + 1]);
+                    }
+                }
+            }
+        });
+}
+
+// --- Generic K-Quant (super-block 256) grid walkers ---
+
+template <typename BlockT, typename PackKtFn>
+static void repack_kquant_to_tiled_grid(ggml_tensor * t, const void * data, size_t offset, size_t size,
+                                        size_t tile_size, PackKtFn && pack_kt) {
+    GGML_ASSERT(offset == 0);
+    GGML_UNUSED(size);
+
+    const BlockT * src_matrix = (const BlockT *) data;
+    const int64_t ne0 = t->ne[0];
+    const int64_t ne1 = t->ne[1];
+    const int64_t ne2 = t->ne[2];
+    const int64_t ne3 = t->ne[3];
+    GGML_ASSERT(ne0 % QK_K == 0);
+
+    const int n_col_tiles    = hex_round_up(ne1, 32) / 32;
+    const int n_k_tiles      = hex_round_up(ne0, 32) / 32;
+    const size_t matrix_size = (size_t) n_col_tiles * n_k_tiles * tile_size;
+    const int64_t sb_per_row = ne0 / QK_K;
+
+    for (int64_t slice_idx = 0; slice_idx < ne2 * ne3; slice_idx++) {
+        const BlockT * src_slice = src_matrix + slice_idx * (ne1 * sb_per_row);
+        uint8_t * matrix_dst     = (uint8_t *) t->data + slice_idx * matrix_size;
+
+        // Padding rows and OR-ed bit planes require zeroed tiles
+        memset(matrix_dst, 0, matrix_size);
+
+        for (int64_t r = 0; r < ne1; r++) {
+            const int ct  = (int) (r / 32);
+            const int row = (int) (r % 32);
+            const BlockT * src_row = src_slice + r * sb_per_row;
+
+            for (int kt = 0; kt < n_k_tiles; kt++) {
+                uint8_t * tile = matrix_dst + ((size_t) ct * n_k_tiles + kt) * tile_size;
+                pack_kt(tile, &src_row[kt / 8], row, kt % 8);
+            }
+        }
+    }
+}
+
+template <typename BlockT, typename UnpackSbFn>
+static void unpack_kquant_from_tiled_grid(void * data, const ggml_tensor * t, size_t offset, size_t size,
+                                          size_t tile_size, UnpackSbFn && unpack_sb) {
+    uint8_t * dst_bytes = (uint8_t *) data;
+    const int64_t ne0 = t->ne[0];
+    const int64_t ne1 = t->ne[1];
+    const int64_t ne2 = t->ne[2];
+    const int64_t ne3 = t->ne[3];
+    GGML_ASSERT(ne0 % QK_K == 0);
+
+    const int n_col_tiles    = hex_round_up(ne1, 32) / 32;
+    const int n_k_tiles      = hex_round_up(ne0, 32) / 32;
+    const size_t matrix_size = (size_t) n_col_tiles * n_k_tiles * tile_size;
+    const int64_t sb_per_row = ne0 / QK_K;
+
+    const size_t row_size_bytes = ggml_row_size(t->type, ne0);
+    const size_t slice_size     = (size_t) ne1 * row_size_bytes;
+    const int64_t start_slice   = offset / slice_size;
+    int64_t end_slice           = (offset + size + slice_size - 1) / slice_size;
+    if (end_slice > ne2 * ne3) {
+        end_slice = ne2 * ne3;
+    }
+
+    for (int64_t slice_idx = start_slice; slice_idx < end_slice; slice_idx++) {
+        const size_t cur_start_byte     = (std::max)(offset, (size_t) slice_idx * slice_size);
+        const size_t cur_end_byte       = (std::min)(offset + size, (size_t) (slice_idx + 1) * slice_size);
+        const size_t slice_offset_start = cur_start_byte - (size_t) slice_idx * slice_size;
+        const size_t slice_offset_end   = cur_end_byte - (size_t) slice_idx * slice_size;
+
+        const int64_t start_row = slice_offset_start / row_size_bytes;
+        const int64_t end_row   = (std::min)((int64_t) ((slice_offset_end + row_size_bytes - 1) / row_size_bytes), ne1);
+
+        BlockT * dst_slice         = (BlockT *) (dst_bytes + (cur_start_byte - offset));
+        const uint8_t * matrix_src = (const uint8_t *) t->data + slice_idx * matrix_size;
+
+        for (int64_t r = start_row; r < end_row; r++) {
+            const int ct  = (int) (r / 32);
+            const int row = (int) (r % 32);
+            BlockT * dst_row = dst_slice + (r - start_row) * sb_per_row;
+
+            for (int64_t sb = 0; sb < sb_per_row; sb++) {
+                BlockT * b = &dst_row[sb];
+                memset(b, 0, sizeof(BlockT));
+                const uint8_t * sb_tiles = matrix_src + ((size_t) ct * n_k_tiles + sb * 8) * tile_size;
+                unpack_sb(b, sb_tiles, tile_size, row);
+            }
+        }
+    }
+}
+
+// --- K-Quant Helpers & Per-Type Implementations (Q2_K, Q3_K, Q4_K, Q5_K, Q6_K) ---
+
+// unsigned 6-bit value (0..63) of element e of a Q6_K block, same bit layout as dequantize_row_q6_K
+static inline uint8_t q6_K_get_quant(const block_q6_K * b, int e) {
+    const int c = e / 128;
+    const int w = e % 128;
+    const int g = w / 32;
+    const int l = w % 32;
+    const uint8_t * ql = b->ql + c * 64;
+    const uint8_t * qh = b->qh + c * 32;
+    uint8_t lo, hi;
+    switch (g) {
+        case 0:  lo = ql[l]      & 0xF; hi = (qh[l] >> 0) & 3; break;
+        case 1:  lo = ql[l + 32] & 0xF; hi = (qh[l] >> 2) & 3; break;
+        case 2:  lo = ql[l]      >> 4;  hi = (qh[l] >> 4) & 3; break;
+        default: lo = ql[l + 32] >> 4;  hi = (qh[l] >> 6) & 3; break;
+    }
+    return (uint8_t) (lo | (hi << 4));
+}
+
+static void repack_q6_K_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_kquant_to_tiled_grid<block_q6_K>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q6_K,
+        [](uint8_t * tile, const block_q6_K * b, int row, int kt_local) {
+            const float d     = GGML_FP16_TO_FP32(b->d);
+            uint8_t * lo_pl   = tile;
+            uint8_t * hi_pl   = tile + 512;
+            ggml_half * sc_pl = (ggml_half *) (tile + 768);
+
+            for (int lk = 0; lk < 32; lk++) {
+                const uint8_t q6 = q6_K_get_quant(b, kt_local * 32 + lk);
+                const int g   = lk >> 2;
+                const int pos = row * 4 + (lk & 3);
+                lo_pl[(g >> 1) * 128 + pos] |= (uint8_t) ((q6 & 0xF) << ((g & 1) * 4));
+                hi_pl[(g >> 2) * 128 + pos] |= (uint8_t) ((q6 >> 4) << ((g & 3) * 2));
+            }
+            for (int sub = 0; sub < 2; sub++) {
+                sc_pl[sub * 32 + row] = GGML_FP32_TO_FP16(d * (float) b->scales[kt_local * 2 + sub]);
+            }
+        });
+}
+
+static void repack_tiled_q6_K(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_kquant_from_tiled_grid<block_q6_K>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q6_K,
+        [](block_q6_K * b, const uint8_t * sb_tiles, size_t tile_size, int row) {
+            float sub_scales[16];
+            for (int kt_local = 0; kt_local < 8; kt_local++) {
+                const uint8_t *   tile  = sb_tiles + kt_local * tile_size;
+                const uint8_t *   lo_pl = tile;
+                const uint8_t *   hi_pl = tile + 512;
+                const ggml_half * sc_pl = (const ggml_half *) (tile + 768);
+
+                const int c  = kt_local / 4;
+                const int g  = kt_local % 4;
+                uint8_t * ql = b->ql + c * 64;
+                uint8_t * qh = b->qh + c * 32;
+
+                for (int lk = 0; lk < 32; lk++) {
+                    const int g_tile = lk >> 2;
+                    const int pos    = row * 4 + (lk & 3);
+                    const uint8_t lo = (lo_pl[(g_tile >> 1) * 128 + pos] >> ((g_tile & 1) * 4)) & 0xF;
+                    const uint8_t hi = (hi_pl[(g_tile >> 2) * 128 + pos] >> ((g_tile & 3) * 2)) & 3;
+
+                    switch (g) {
+                        case 0:  ql[lk]      |= lo;        qh[lk] |= (hi << 0); break;
+                        case 1:  ql[lk + 32] |= lo;        qh[lk] |= (hi << 2); break;
+                        case 2:  ql[lk]      |= (lo << 4); qh[lk] |= (hi << 4); break;
+                        default: ql[lk + 32] |= (lo << 4); qh[lk] |= (hi << 6); break;
+                    }
+                }
+
+                for (int sub = 0; sub < 2; sub++) {
+                    sub_scales[kt_local * 2 + sub] = GGML_FP16_TO_FP32(sc_pl[sub * 32 + row]);
+                }
+            }
+
+            float max_abs_scale = 0.0f;
+            for (int s = 0; s < 16; s++) {
+                max_abs_scale = (std::max)(max_abs_scale, fabsf(sub_scales[s]));
+            }
+
+            if (max_abs_scale == 0.0f) {
+                b->d = GGML_FP32_TO_FP16(0.0f);
+            } else {
+                b->d = GGML_FP32_TO_FP16(max_abs_scale / 127.0f);
+                const float d_actual = GGML_FP16_TO_FP32(b->d);
+                const float inv_d    = (d_actual != 0.0f) ? (1.0f / d_actual) : 0.0f;
+                for (int s = 0; s < 16; s++) {
+                    const int sc = (int) roundf(sub_scales[s] * inv_d);
+                    b->scales[s] = (int8_t) (std::max)(-128, (std::min)(127, sc));
+                }
+            }
+        });
+}
+
+// low 2 bits (0..3) of element e of a Q2_K or Q3_K block, same bit layout as dequantize_row_q2_K / q3_K
+static inline uint8_t q2_3_K_get_low2(const uint8_t * qs, int e) {
+    const int c = e / 128;
+    const int j = (e % 128) / 32;
+    const int l = e % 32;
+    return (qs[c * 32 + l] >> (2 * j)) & 3;
+}
+
+// hmask bit of element e of a Q3_K block, same bit layout as dequantize_row_q3_K
+static inline bool q3_K_get_hbit(const block_q3_K * b, int e) {
+    const int c = e / 128;
+    const int j = (e % 128) / 32;
+    const int l = e % 32;
+    return (b->hmask[l] >> (c * 4 + j)) & 1;
+}
+
+// signed 6-bit scale j (-32..31) of a Q3_K block, same packing as quantize_row_q3_K_ref
+static inline int q3_K_get_scale(const uint8_t * scales, int j) {
+    const int lo = (j < 8) ? (scales[j] & 0xF) : (scales[j - 8] >> 4);
+    const int hi = (scales[8 + j % 4] >> (2 * (j / 4))) & 3;
+    return (lo | (hi << 4)) - 32;
+}
+
+// read-back: find fp16 d and l[j] in [lmin, lmax] with fp16(d * l[j]) == prod[j] for all j, false if none
+static bool hexagon_recover_k_scales(const ggml_half * prod, int n, int lmin, int lmax, ggml_half * d_out, int * l_out) {
+    int jmax = 0;
+    for (int j = 1; j < n; j++) {
+        if (fabsf(GGML_FP16_TO_FP32(prod[j])) > fabsf(GGML_FP16_TO_FP32(prod[jmax]))) {
+            jmax = j;
+        }
+    }
+    const float pmax = GGML_FP16_TO_FP32(prod[jmax]);
+    if (pmax == 0.0f) {
+        *d_out = GGML_FP32_TO_FP16(0.0f);
+        for (int j = 0; j < n; j++) {
+            l_out[j] = 0;
+        }
+        return true;
+    }
+
+    // the quantizers put the largest scale at or near the range end, so try large |l| first
+    const int lext = (std::max)(-lmin, lmax);
+    for (int a = lext; a >= 1; a--) {
+        for (int sign : { 1, -1 }) {
+            const int lj = sign * a;
+            if (lj < lmin || lj > lmax) {
+                continue;
+            }
+            const ggml_half d0 = GGML_FP32_TO_FP16(pmax / (float) lj);
+            for (int ulp : { 0, -1, 1 }) {
+                ggml_half d = d0;
+                uint16_t bits;
+                memcpy(&bits, &d, sizeof(bits));
+                bits = (uint16_t) (bits + ulp);
+                memcpy(&d, &bits, sizeof(bits));
+
+                const float df = GGML_FP16_TO_FP32(d);
+                if (!std::isfinite(df) || df == 0.0f) {
+                    continue;
+                }
+                bool ok = true;
+                for (int j = 0; j < n && ok; j++) {
+                    const int l = (int) roundf(GGML_FP16_TO_FP32(prod[j]) / df);
+                    const ggml_half p = GGML_FP32_TO_FP16(df * (float) l);
+                    ok = l >= lmin && l <= lmax && memcmp(&p, &prod[j], sizeof(p)) == 0;
+                    l_out[j] = l;
+                }
+                if (ok) {
+                    *d_out = d;
+                    return true;
+                }
+            }
+        }
+    }
+    return false;
+}
+
+static void repack_q3_K_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_kquant_to_tiled_grid<block_q3_K>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q3_K,
+        [](uint8_t * tile, const block_q3_K * b, int row, int kt_local) {
+            const float d     = GGML_FP16_TO_FP32(b->d);
+            uint8_t * lo_pl   = tile;
+            uint8_t * neg_pl  = tile + 256;
+            ggml_half * sc_pl = (ggml_half *) (tile + 384);
+
+            for (int lk = 0; lk < 32; lk++) {
+                const int e   = kt_local * 32 + lk;
+                const int g   = lk >> 2;
+                const int pos = row * 4 + (lk & 3);
+                lo_pl[(g >> 2) * 128 + pos] |= (uint8_t) (q2_3_K_get_low2(b->qs, e) << ((g & 3) * 2));
+                if (!q3_K_get_hbit(b, e)) {
+                    neg_pl[pos] |= (uint8_t) (1 << g);
+                }
+            }
+            for (int sub = 0; sub < 2; sub++) {
+                sc_pl[sub * 32 + row] = GGML_FP32_TO_FP16(d * (float) q3_K_get_scale(b->scales, kt_local * 2 + sub));
+            }
+        });
+}
+
+static void repack_tiled_q3_K(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_kquant_from_tiled_grid<block_q3_K>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q3_K,
+        [](block_q3_K * b, const uint8_t * sb_tiles, size_t tile_size, int row) {
+            ggml_half sub_scales[16];
+            for (int kt_local = 0; kt_local < 8; kt_local++) {
+                const uint8_t *   tile   = sb_tiles + kt_local * tile_size;
+                const uint8_t *   lo_pl  = tile;
+                const uint8_t *   neg_pl = tile + 256;
+                const ggml_half * sc_pl  = (const ggml_half *) (tile + 384);
+
+                for (int lk = 0; lk < 32; lk++) {
+                    const int e   = kt_local * 32 + lk;
+                    const int g   = lk >> 2;
+                    const int pos = row * 4 + (lk & 3);
+                    const uint8_t lo = (lo_pl[(g >> 2) * 128 + pos] >> ((g & 3) * 2)) & 3;
+
+                    const int c = e / 128;
+                    const int j = (e % 128) / 32;
+                    const int l = e % 32;
+                    b->qs[c * 32 + l] |= (uint8_t) (lo << (2 * j));
+                    if (!((neg_pl[pos] >> g) & 1)) {
+                        b->hmask[l] |= (uint8_t) (1 << (c * 4 + j));
+                    }
+                }
+
+                for (int sub = 0; sub < 2; sub++) {
+                    sub_scales[kt_local * 2 + sub] = sc_pl[sub * 32 + row];
+                }
+            }
+
+            int ls[16];
+            if (!hexagon_recover_k_scales(sub_scales, 16, -32, 31, &b->d, ls)) {
+                float max_scale = 0.0f;
+                for (int s = 0; s < 16; s++) {
+                    if (fabsf(GGML_FP16_TO_FP32(sub_scales[s])) > fabsf(max_scale)) {
+                        max_scale = GGML_FP16_TO_FP32(sub_scales[s]);
+                    }
+                }
+                b->d = GGML_FP32_TO_FP16(-max_scale / 32.0f);
+                const float d_actual = GGML_FP16_TO_FP32(b->d);
+                const float inv_d    = (d_actual != 0.0f) ? (1.0f / d_actual) : 0.0f;
+                for (int s = 0; s < 16; s++) {
+                    ls[s] = (std::max)(-32, (std::min)(31, (int) roundf(GGML_FP16_TO_FP32(sub_scales[s]) * inv_d)));
+                }
+            }
+
+            for (int s = 0; s < 16; s++) {
+                const int l = ls[s] + 32;
+                if (s < 8) {
+                    b->scales[s] = l & 0xF;
+                } else {
+                    b->scales[s - 8] |= (uint8_t) ((l & 0xF) << 4);
+                }
+                b->scales[s % 4 + 8] |= (uint8_t) ((l >> 4) << (2 * (s / 4)));
+            }
+        });
+}
+
+static void repack_q2_K_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_kquant_to_tiled_grid<block_q2_K>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q2_K,
+        [](uint8_t * tile, const block_q2_K * b, int row, int kt_local) {
+            uint8_t * lo_pl   = tile;
+            uint8_t * sc_pl   = tile + 256;
+            ggml_half * dd_pl = (ggml_half *) (tile + 320);
+
+            for (int lk = 0; lk < 32; lk++) {
+                const int g   = lk >> 2;
+                const int pos = row * 4 + (lk & 3);
+                lo_pl[(g >> 2) * 128 + pos] |= (uint8_t) (q2_3_K_get_low2(b->qs, kt_local * 32 + lk) << ((g & 3) * 2));
+            }
+            for (int sub = 0; sub < 2; sub++) {
+                sc_pl[sub * 32 + row] = b->scales[kt_local * 2 + sub];
+            }
+            if (row / 8 == kt_local % 4) {
+                dd_pl[row % 8] = kt_local < 4 ? b->d : (ggml_half) (b->dmin ^ 0x8000);
+            }
+        });
+}
+
+static void repack_tiled_q2_K(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_kquant_from_tiled_grid<block_q2_K>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q2_K,
+        [](block_q2_K * b, const uint8_t * sb_tiles, size_t tile_size, int row) {
+            for (int kt_local = 0; kt_local < 8; kt_local++) {
+                const uint8_t *   tile  = sb_tiles + kt_local * tile_size;
+                const uint8_t *   lo_pl = tile;
+                const uint8_t *   sc_pl = tile + 256;
+                const ggml_half * dd_pl = (const ggml_half *) (tile + 320);
+
+                for (int lk = 0; lk < 32; lk++) {
+                    const int e   = kt_local * 32 + lk;
+                    const int g   = lk >> 2;
+                    const int pos = row * 4 + (lk & 3);
+                    const uint8_t lo = (lo_pl[(g >> 2) * 128 + pos] >> ((g & 3) * 2)) & 3;
+                    b->qs[(e / 128) * 32 + e % 32] |= (uint8_t) (lo << (2 * ((e % 128) / 32)));
+                }
+
+                for (int sub = 0; sub < 2; sub++) {
+                    b->scales[kt_local * 2 + sub] = sc_pl[sub * 32 + row];
+                }
+                if (row / 8 == kt_local % 4) {
+                    if (kt_local < 4) {
+                        b->d = dd_pl[row % 8];
+                    } else {
+                        b->dmin = (ggml_half) (dd_pl[row % 8] ^ 0x8000);
+                    }
+                }
+            }
+        });
+}
+
+static inline void get_scale_min_k4(int j, const uint8_t * q, uint8_t * d, uint8_t * m) {
+    if (j < 4) {
+        *d = q[j] & 63;
+        *m = q[j + 4] & 63;
+    } else {
+        *d = (q[j + 4] & 0xF) | ((q[j - 4] >> 6) << 4);
+        *m = (q[j + 4] >>  4) | ((q[j - 0] >> 6) << 4);
+    }
+}
+
+static inline void pack_k4_scales_mins(uint8_t * scales, ggml_half * d, ggml_half * dmin,
+                                       const float * sub_scales, const float * sub_mins) {
+    float max_scale = 0.0f;
+    float max_min   = 0.0f;
+    for (int j = 0; j < 8; j++) {
+        if (sub_scales[j] > max_scale) max_scale = sub_scales[j];
+        if (sub_mins[j]   > max_min)   max_min   = sub_mins[j];
+    }
+
+    float inv_scale = 0.0f;
+    if (max_scale > 0.0f) {
+        *d = GGML_FP32_TO_FP16(max_scale / 63.0f);
+        const float d_actual = GGML_FP16_TO_FP32(*d);
+        inv_scale = (d_actual > 0.0f) ? (1.0f / d_actual) : 0.0f;
+    } else {
+        *d = GGML_FP32_TO_FP16(0.0f);
+    }
+
+    float inv_min = 0.0f;
+    if (max_min > 0.0f) {
+        *dmin = GGML_FP32_TO_FP16(max_min / 63.0f);
+        const float dmin_actual = GGML_FP16_TO_FP32(*dmin);
+        inv_min = (dmin_actual > 0.0f) ? (1.0f / dmin_actual) : 0.0f;
+    } else {
+        *dmin = GGML_FP32_TO_FP16(0.0f);
+    }
+
+    for (int j = 0; j < 8; j++) {
+        uint8_t ls = (std::min)((uint8_t) 63, (uint8_t) roundf(inv_scale * sub_scales[j]));
+        uint8_t lm = (std::min)((uint8_t) 63, (uint8_t) roundf(inv_min * sub_mins[j]));
+        if (j < 4) {
+            scales[j]     = ls;
+            scales[j + 4] = lm;
+        } else {
+            scales[j + 4] = (ls & 0xF) | ((lm & 0xF) << 4);
+            scales[j - 4] |= ((ls >> 4) << 6);
+            scales[j - 0] |= ((lm >> 4) << 6);
+        }
+    }
+}
+
+static void repack_q4_K_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_kquant_to_tiled_grid<block_q4_K>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q4_K,
+        [](uint8_t * tile_dst, const block_q4_K * b, int row, int kt_local) {
+            const float d    = GGML_FP16_TO_FP32(b->d);
+            const float dmin = GGML_FP16_TO_FP32(b->dmin);
+
+            uint8_t sc, m;
+            get_scale_min_k4(kt_local, b->scales, &sc, &m);
+
+            const uint8_t * qs_sub = b->qs + (kt_local / 2) * 32;
+            const int shift = (kt_local & 1) ? 4 : 0;
+
+            for (int cp = 0; cp < 16; cp++) {
+                const uint8_t q0 = (qs_sub[2 * cp + 0] >> shift) & 0x0F;
+                const uint8_t q1 = (qs_sub[2 * cp + 1] >> shift) & 0x0F;
+                tile_dst[cp * 32 + row] = (uint8_t) ((q1 << 4) | q0);
+            }
+
+            ggml_half * scale_dst = (ggml_half *) (tile_dst + 512);
+            scale_dst[2 * row + 0] = GGML_FP32_TO_FP16(d * (float) sc);
+            scale_dst[2 * row + 1] = GGML_FP32_TO_FP16(-dmin * (float) m);
+        });
+}
+
+static void repack_tiled_q4_K(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_kquant_from_tiled_grid<block_q4_K>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q4_K,
+        [](block_q4_K * b, const uint8_t * sb_tiles, size_t tile_size, int row) {
+            float sub_scales[8];
+            float sub_mins[8];
+
+            for (int kt_local = 0; kt_local < 8; kt_local++) {
+                const uint8_t *   tile_src  = sb_tiles + kt_local * tile_size;
+                const ggml_half * scale_src = (const ggml_half *) (tile_src + 512);
+
+                uint8_t * qs_sub = b->qs + (kt_local / 2) * 32;
+                const int shift  = (kt_local & 1) ? 4 : 0;
+
+                for (int cp = 0; cp < 16; cp++) {
+                    const uint8_t val = tile_src[cp * 32 + row];
+                    qs_sub[2 * cp + 0] |= (uint8_t) ((val & 0x0F) << shift);
+                    qs_sub[2 * cp + 1] |= (uint8_t) ((val >> 4)   << shift);
+                }
+
+                const float D = GGML_FP16_TO_FP32(scale_src[2 * row + 0]);
+                const float M = GGML_FP16_TO_FP32(scale_src[2 * row + 1]);
+                sub_scales[kt_local] = (D > 0.0f) ? D : 0.0f;
+                sub_mins[kt_local]   = (-M > 0.0f) ? -M : 0.0f;
+            }
+
+            pack_k4_scales_mins(b->scales, &b->d, &b->dmin, sub_scales, sub_mins);
+        });
+}
+
+static void repack_q5_K_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+    repack_kquant_to_tiled_grid<block_q5_K>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q5_K,
+        [](uint8_t * tile_dst, const block_q5_K * b, int row, int kt_local) {
+            const float d     = GGML_FP16_TO_FP32(b->d);
+            const float dmin  = GGML_FP16_TO_FP32(b->dmin);
+            uint8_t * plane   = tile_dst + 640;
+
+            uint8_t sc, m;
+            get_scale_min_k4(kt_local, b->scales, &sc, &m);
+
+            const uint8_t * qs_sub = b->qs + (kt_local / 2) * 32;
+            const int shift    = (kt_local & 1) ? 4 : 0;
+            const uint8_t hbit = (uint8_t) (1 << kt_local);
+
+            for (int cp = 0; cp < 16; cp++) {
+                const uint8_t q0 = (qs_sub[2 * cp + 0] >> shift) & 0x0F;
+                const uint8_t q1 = (qs_sub[2 * cp + 1] >> shift) & 0x0F;
+                tile_dst[cp * 32 + row] = (uint8_t) ((q1 << 4) | q0);
+
+                const int i    = cp / 4;
+                const int lane = (cp % 4) * 32 + row;
+                if (b->qh[2 * cp + 0] & hbit) {
+                    plane[lane] |= (uint8_t) (1 << (2 * i));
+                }
+                if (b->qh[2 * cp + 1] & hbit) {
+                    plane[lane] |= (uint8_t) (1 << (2 * i + 1));
+                }
+            }
+
+            ggml_half * scale_dst = (ggml_half *) (tile_dst + 512);
+            scale_dst[2 * row + 0] = GGML_FP32_TO_FP16(d * (float) sc);
+            scale_dst[2 * row + 1] = GGML_FP32_TO_FP16(-dmin * (float) m);
+        });
+}
+
+static void repack_tiled_q5_K(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+    unpack_kquant_from_tiled_grid<block_q5_K>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q5_K,
+        [](block_q5_K * b, const uint8_t * sb_tiles, size_t tile_size, int row) {
+            float sub_scales[8];
+            float sub_mins[8];
+
+            for (int kt_local = 0; kt_local < 8; kt_local++) {
+                const uint8_t *   tile_src  = sb_tiles + kt_local * tile_size;
+                const uint8_t *   plane     = tile_src + 640;
+                const ggml_half * scale_src = (const ggml_half *) (tile_src + 512);
+
+                uint8_t * qs_sub   = b->qs + (kt_local / 2) * 32;
+                const int shift    = (kt_local & 1) ? 4 : 0;
+                const uint8_t hbit = (uint8_t) (1 << kt_local);
+
+                for (int cp = 0; cp < 16; cp++) {
+                    const uint8_t val = tile_src[cp * 32 + row];
+                    qs_sub[2 * cp + 0] |= (uint8_t) ((val & 0x0F) << shift);
+                    qs_sub[2 * cp + 1] |= (uint8_t) ((val >> 4)   << shift);
+
+                    const int i    = cp / 4;
+                    const int lane = (cp % 4) * 32 + row;
+                    if (plane[lane] & (1 << (2 * i))) {
+                        b->qh[2 * cp + 0] |= hbit;
+                    }
+                    if (plane[lane] & (1 << (2 * i + 1))) {
+                        b->qh[2 * cp + 1] |= hbit;
+                    }
+                }
+
+                const float D = GGML_FP16_TO_FP32(scale_src[2 * row + 0]);
+                const float M = GGML_FP16_TO_FP32(scale_src[2 * row + 1]);
+                sub_scales[kt_local] = (D > 0.0f) ? D : 0.0f;
+                sub_mins[kt_local]   = (-M > 0.0f) ? -M : 0.0f;
+            }
+
+            pack_k4_scales_mins(b->scales, &b->d, &b->dmin, sub_scales, sub_mins);
+        });
+}
+
+// --- Top-level Repack / Unpack Dispatchers ---
+
+void htp_repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t size) {
+    switch (tensor->type) {
+        case GGML_TYPE_Q4_0:
+        case GGML_TYPE_IQ4_NL:
+            repack_q4_0_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_Q4_1:
+            repack_q4_1_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_Q8_0:
+            repack_q8_0_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_MXFP4:
+            repack_mxfp4_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_Q2_K:
+            repack_q2_K_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_Q3_K:
+            repack_q3_K_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_Q4_K:
+            repack_q4_K_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_Q5_K:
+            repack_q5_K_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_Q6_K:
+            repack_q6_K_tiled(tensor, data, 0, size);
+            break;
+        case GGML_TYPE_F16:
+            repack_float_to_f16_tiled<uint16_t>(tensor, data, 0, size,
+                [](uint16_t v) { return v; });
+            break;
+        case GGML_TYPE_F32:
+            repack_float_to_f16_tiled<float>(tensor, data, 0, size,
+                [](float v) { return (uint16_t) ggml_fp32_to_fp16(v); });
+            break;
+        case GGML_TYPE_BF16:
+            repack_float_to_f16_tiled<ggml_bf16_t>(tensor, data, 0, size,
+                [](ggml_bf16_t v) { return (uint16_t) ggml_fp32_to_fp16(ggml_bf16_to_fp32(v)); });
+            break;
+        default:
+            break;
+    }
+}
+
+void htp_unpack_tensor_tiled(void * data, const ggml_tensor * tensor, size_t offset, size_t size) {
+    GGML_ASSERT(offset + size <= ggml_nbytes(tensor));
+
+    switch (tensor->type) {
+        case GGML_TYPE_Q4_0:
+        case GGML_TYPE_IQ4_NL:
+            repack_tiled_q4_0(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_Q4_1:
+            repack_tiled_q4_1(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_Q8_0:
+            repack_tiled_q8_0(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_MXFP4:
+            repack_tiled_mxfp4(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_Q2_K:
+            repack_tiled_q2_K(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_Q3_K:
+            repack_tiled_q3_K(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_Q4_K:
+            repack_tiled_q4_K(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_Q5_K:
+            repack_tiled_q5_K(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_Q6_K:
+            repack_tiled_q6_K(data, tensor, offset, size);
+            break;
+        case GGML_TYPE_F16:
+            unpack_f16_tiled_to_float<uint16_t>(data, tensor, offset, size,
+                [](uint16_t v) { return v; });
+            break;
+        case GGML_TYPE_F32:
+            unpack_f16_tiled_to_float<float>(data, tensor, offset, size,
+                [](uint16_t v) { return ggml_fp16_to_fp32((ggml_fp16_t) v); });
+            break;
+        case GGML_TYPE_BF16:
+            unpack_f16_tiled_to_float<ggml_bf16_t>(data, tensor, offset, size,
+                [](uint16_t v) { return ggml_fp32_to_bf16(ggml_fp16_to_fp32((ggml_fp16_t) v)); });
+            break;
+        default:
+            memcpy(data, (const char *) tensor->data + offset, size);
+            break;
+    }
+}

ggml/src/ggml-hexagon/htp-repack.h

@@ -0,0 +1,64 @@
+#ifndef HTP_REPACK_H
+#define HTP_REPACK_H
+
+#include <stddef.h>
+#include <stdint.h>
+
+#include "ggml.h"
+#include "htp/htp-tensor.h"
+
+static inline bool htp_is_tiled_weight_type(enum ggml_type type) {
+    return type == GGML_TYPE_Q4_0  || type == GGML_TYPE_Q4_1   ||
+           type == GGML_TYPE_Q8_0  || type == GGML_TYPE_IQ4_NL ||
+           type == GGML_TYPE_MXFP4 || type == GGML_TYPE_Q6_K   ||
+           type == GGML_TYPE_Q4_K  || type == GGML_TYPE_Q5_K   ||
+           type == GGML_TYPE_Q3_K  || type == GGML_TYPE_Q2_K   ||
+           type == GGML_TYPE_F16   || type == GGML_TYPE_F32    ||
+           type == GGML_TYPE_BF16;
+}
+
+static inline bool htp_is_tiled_weight_tensor(const struct ggml_tensor * t) {
+    if (!ggml_is_contiguous(t) || ggml_is_view(t)) {
+        return false;
+    }
+    if (!htp_is_tiled_weight_type(t->type)) {
+        return false;
+    }
+    if (t->ne[0] < 32) {
+        return false;
+    }
+    if (!ggml_is_quantized(t->type)) {
+        return (t->ne[0] % 32 == 0) && (t->ne[1] % 32 == 0) && (t->ne[2] == 1) && (t->ne[3] == 1);
+    }
+    return t->ne[3] == 1;
+}
+
+// Size of one tiled row in the DSP tiled layout. The Q3_K..Q6_K tiles store uncompressed scales/mins,
+// so they are larger than the ggml blocks. For the other tiled types (Q2_K included) the tile has the same size.
+static inline size_t htp_tiled_row_size(enum ggml_type type, int64_t ne0) {
+    if (type == GGML_TYPE_F16 || type == GGML_TYPE_F32 || type == GGML_TYPE_BF16) {
+        return (size_t) ne0 * sizeof(ggml_fp16_t);
+    }
+    if (type == GGML_TYPE_Q6_K) {
+        return (size_t) (ne0 / 32) * (HTP_TENSOR_TILE_SIZE_Q6_K / 32);
+    }
+    if (type == GGML_TYPE_Q4_K) {
+        return (size_t) (ne0 / 32) * (HTP_TENSOR_TILE_SIZE_Q4_K / 32);
+    }
+    if (type == GGML_TYPE_Q5_K) {
+        return (size_t) (ne0 / 32) * (HTP_TENSOR_TILE_SIZE_Q5_K / 32);
+    }
+    if (type == GGML_TYPE_Q3_K) {
+        return (size_t) (ne0 / 32) * (HTP_TENSOR_TILE_SIZE_Q3_K / 32);
+    }
+    return ggml_row_size(type, ne0);
+}
+
+static inline bool htp_is_hmx_weight_type(enum ggml_type type) {
+    return htp_is_tiled_weight_type(type);
+}
+
+void htp_repack_tensor_tiled(struct ggml_tensor * t, const void * data, size_t size);
+void htp_unpack_tensor_tiled(void * data, const struct ggml_tensor * t, size_t offset, size_t size);
+
+#endif // HTP_REPACK_H

ggml/src/ggml-hexagon/htp/CMakeLists.txt

@@ -20,7 +20,7 @@ add_library(${HTP_LIB} SHARED
     work-queue.c
     dma-queue.c
     hmx-queue.c
-    htp-tensor.c
+    htp-dcache.c
     matmul-ops.c
     flash-attn-ops.c
     gated-delta-net-ops.c

ggml/src/ggml-hexagon/htp/flash-attn-ops.c

@@ -361,7 +361,7 @@ static void flash_attn_ext_f16_thread(unsigned int nth, unsigned int ith, void *
                 htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_FA_K_PREP, ir);
                 for (uint32_t r = 0; r < current_block_size; ++r) {
                     __fp16 * row_k = (__fp16 *)(k_base + r * factx->size_k_row_padded);
-                    hvx_dequantize_row_q8_0_f16(row_k, row_k, DK);
+                    hvx_dequantize_row_flat_q8_0_f16(row_k, row_k, DK);
                 }
                 htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_FA_K_PREP, ir);
             }
@@ -369,7 +369,7 @@ static void flash_attn_ext_f16_thread(unsigned int nth, unsigned int ith, void *
                 htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_FA_V_PREP, ir);
                 for (uint32_t r = 0; r < current_block_size; ++r) {
                     __fp16 * row_v = (__fp16 *)(v_base + r * factx->size_v_row_padded);
-                    hvx_dequantize_row_q8_0_f16(row_v, row_v, DV);
+                    hvx_dequantize_row_flat_q8_0_f16(row_v, row_v, DV);
                 }
                 htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_FA_V_PREP, ir);
             }
@@ -625,7 +625,7 @@ static void fa_k_interleave_thread(unsigned int n, unsigned int i, void * data)
     if (factx->k->type == HTP_TYPE_Q8_0) {
         for (uint32_t r = start; r < end; ++r) {
             __fp16 * row_k = (__fp16 *)((char *)args->curr_k + r * args->src_stride * sizeof(__fp16));
-            hvx_dequantize_row_q8_0_f16(row_k, row_k, factx->DK);
+            hvx_dequantize_row_flat_q8_0_f16(row_k, row_k, factx->DK);
         }
     }
     hmx_interleave_rows_to_tiles(factx->vtcm_k_tiles[args->buf_idx], (const __fp16 *) args->curr_k, total_rows, factx->DK_pad,
@@ -679,7 +679,7 @@ static void fa_v_interleave_thread(unsigned int n, unsigned int i, void * data)
     if (factx->v->type == HTP_TYPE_Q8_0) {
         for (uint32_t r = start; r < end; ++r) {
             __fp16 * row_v = (__fp16 *)((char *)args->v_src + r * args->src_stride * sizeof(__fp16));
-            hvx_dequantize_row_q8_0_f16(row_v, row_v, factx->DV);
+            hvx_dequantize_row_flat_q8_0_f16(row_v, row_v, factx->DV);
         }
     }
     hmx_interleave_cols_to_tiles(v_tiles_dst, (const __fp16 *) args->v_src, total_rows, factx->DV_pad,

ggml/src/ggml-hexagon/htp/get-rows-ops.c

@@ -17,21 +17,32 @@
 #include "htp-tensor.h"
 #include "hvx-utils.h"
 #include "hvx-quant.h"
-#include "matmul-ops.h"
 #include "get-rows-ops.h"
 #include "work-queue.h"
 
+struct get_rows_context;
+
+typedef uint32_t (*get_rows_read_index_fn_t)(const void * src1, uint32_t idx, int idx_type);
+typedef void (*get_rows_flat_compute_fn_t)(float * dst_spad, const void * src_spad, uint32_t cur_elems);
+typedef void (*get_rows_tiled_compute_fn_t)(float * dst_spad, const uint8_t * src_spad, uint32_t n_k_tiles, uint32_t tile_stride, uint32_t row);
+
 struct get_rows_context {
     struct htp_ops_context * octx;
     const struct htp_get_rows_kernel_params * kparams;
     struct htp_get_rows_vtcm_layout vtcm_layout;
-    uint8_t * vtcm_base;
+    uint8_t    * vtcm_base;
+    const void * src1_data;
+    get_rows_read_index_fn_t    read_index;
+    get_rows_flat_compute_fn_t  compute_flat;
+    get_rows_tiled_compute_fn_t compute_tiled;
     uint32_t task_start;
     uint32_t tasks;
     uint32_t tasks_per_thread;
     uint32_t tile_size;
     uint32_t tile_stride;
-    bool index_i32;
+    uint32_t src0_row_size;
+    uint32_t dst_row_size;
+    int idx_type;
 };
 
 #define get_rows_preamble                      \
@@ -64,229 +75,169 @@ struct get_rows_context {
                                                \
     const uint32_t nr = ne10 * ne11 * ne12;
 
-#define GET_ROWS_THREAD_ST_FN(IDX_TYPE)                                                                                 \
-static void get_rows_thread_st_##IDX_TYPE(unsigned int nth, unsigned int ith, void *data) {                             \
-    struct get_rows_context * grctx = (struct get_rows_context *)data;                                                  \
-    struct htp_ops_context * octx = grctx->octx;                                                                        \
-    const struct htp_get_rows_kernel_params * kparams = grctx->kparams;                                                 \
-    get_rows_preamble;                                                                                                  \
-    const uint32_t dr  = grctx->tasks_per_thread;                                                                       \
-    const uint32_t ir0 = grctx->task_start + dr * ith;                                                                  \
-    if (ir0 >= grctx->task_start + grctx->tasks) {                                                                      \
-        return;                                                                                                         \
-    }                                                                                                                   \
-    const uint32_t ir1 = MIN(ir0 + dr, grctx->task_start + grctx->tasks);                                               \
-    const uint32_t row_size_bytes = htp_tensor_get_row_size(octx->src[0]->type, ne00);                                  \
-    dma_queue * dma_q = octx->ctx->dma[ith];                                                                            \
-    for (uint32_t i = ir0; i < ir1; ++i) {                                                                              \
-        const uint32_t i12 = fastdiv(i, &kparams->div_ne10_ne11);                                                       \
-        const uint32_t rem = i - i12 * ne11 * ne10;                                                                     \
-        const uint32_t i11 = fastdiv(rem, &kparams->div_ne10);                                                          \
-        const uint32_t i10 = rem - i11 * ne10;                                                                          \
-        const IDX_TYPE * src1_ptr = (const IDX_TYPE *)(uintptr_t)(octx->src[1]->data + i10*nb10 + i11*nb11 + i12*nb12); \
-        const uint32_t i01 = (uint32_t)*src1_ptr;                                                                       \
-        assert(i01 < ne01);                                                                                             \
-        const uint32_t q02 = fastdiv(i11, &kparams->div_ne02);                                                          \
-        const uint32_t i02 = i11 - q02 * ne02;                                                                          \
-        const uint32_t q03 = fastdiv(i12, &kparams->div_ne03);                                                          \
-        const uint32_t i03 = i12 - q03 * ne03;                                                                          \
-        const dma_addr_t src0_data = octx->src[0]->data + i01*nb01 + i02*nb02 + i03*nb03;                               \
-        const dma_addr_t dst_data  = octx->dst->data    + i10*nb1  + i11*nb2  + i12*nb3;                                \
-        while (!dma_queue_push(dma_q, dma_make_data(dst_data, src0_data), nb1, nb01,                                    \
-                               row_size_bytes, 1)) {                                                                    \
-            dma_queue_pop(dma_q);                                                                                       \
-        }                                                                                                               \
-    }                                                                                                                   \
-    dma_queue_flush(dma_q);                                                                                             \
+static uint32_t get_rows_read_index_ddr(
+    const void * src1,
+    uint32_t idx,
+    int idx_type
+) {
+    const uint32_t off = idx * ((idx_type == HTP_TYPE_I64) ? sizeof(int64_t) : sizeof(int32_t));
+    return *(const uint32_t *) ((const uint8_t *) src1 + off);
 }
 
-GET_ROWS_THREAD_ST_FN(int32_t)
-GET_ROWS_THREAD_ST_FN(int64_t)
-
-#define GET_ROWS_THREAD_DT_FN(TYPE_NAME, SRC0_SIZE_EXPR, IDX_TYPE, COMPUTE_EXPR)                                                \
-static void get_rows_thread_##TYPE_NAME##_##IDX_TYPE(unsigned int nth, unsigned int ith, void *data) {                          \
-    struct get_rows_context * grctx = (struct get_rows_context *)data;                                                          \
-    struct htp_ops_context * octx = grctx->octx;                                                                                \
-    const struct htp_get_rows_kernel_params * kparams = grctx->kparams;                                                         \
-    get_rows_preamble;                                                                                                          \
-    struct htp_thread_trace * tr = &octx->ctx->trace[ith];                                                                      \
-    const uint32_t dr  = grctx->tasks_per_thread;                                                                               \
-    const uint32_t ir0 = grctx->task_start + dr * ith;                                                                          \
-    if (ir0 >= grctx->task_start + grctx->tasks) {                                                                              \
-        return;                                                                                                                 \
-    }                                                                                                                           \
-    const uint32_t ir1 = MIN(ir0 + dr, grctx->task_start + grctx->tasks);                                                       \
-    const uint32_t chunks_per_row = kparams->chunks_per_row;                                                                    \
-    const uint32_t chunk_size     = kparams->chunk_size;                                                                        \
-    dma_queue * dma_q = octx->ctx->dma[ith];                                                                                    \
-    const struct htp_get_rows_vtcm_layout * vtcm_layout = &grctx->vtcm_layout;                                                  \
-    uint8_t * vtcm_src0 = grctx->vtcm_base + vtcm_layout->off_src0 + ith * vtcm_layout->src0_bytes_per_thread;                  \
-    uint8_t * vtcm_dst  = grctx->vtcm_base + vtcm_layout->off_dst  + ith * vtcm_layout->dst_bytes_per_thread;                   \
-    for (uint32_t step = 0, spad_idx = 0; step < ir1 - ir0 && spad_idx < 2; ++step, spad_idx++) {                               \
-        const uint32_t i = ir0 + step;                                                                                          \
-        const uint32_t row_idx   = fastdiv(i, &kparams->div_chunks_per_row);                                                    \
-        const uint32_t chunk_idx = i - row_idx * chunks_per_row;                                                                \
-        const uint32_t i12 = fastdiv(row_idx, &kparams->div_ne10_ne11);                                                         \
-        const uint32_t rem = row_idx - i12 * ne11 * ne10;                                                                       \
-        const uint32_t i11 = fastdiv(rem, &kparams->div_ne10);                                                                  \
-        const uint32_t i10 = rem - i11 * ne10;                                                                                  \
-        const IDX_TYPE * src1_ptr = (const IDX_TYPE *)(uintptr_t)(octx->src[1]->data + i10*nb10 + i11*nb11 + i12*nb12);         \
-        const uint32_t i01 = (uint32_t)*src1_ptr;                                                                               \
-        assert(i01 < ne01);                                                                                                     \
-        const uint32_t q02 = fastdiv(i11, &kparams->div_ne02);                                                                  \
-        const uint32_t i02 = i11 - q02 * ne02;                                                                                  \
-        const uint32_t q03 = fastdiv(i12, &kparams->div_ne03);                                                                  \
-        const uint32_t i03 = i12 - q03 * ne03;                                                                                  \
-        const uint32_t offset = chunk_idx * chunk_size;                                                                         \
-        const uint32_t cur_elems = (offset < ne00) ? MIN(chunk_size, ne00 - offset) : 0;                                        \
-        const uint32_t cur_src0_bytes = SRC0_SIZE_EXPR(cur_elems);                                                              \
-        const uint32_t cur_dst_bytes  = cur_elems * sizeof(float);                                                              \
-        const dma_addr_t src0_data = octx->src[0]->data + i01*nb01 + i02*nb02 + i03*nb03 + SRC0_SIZE_EXPR(offset);              \
-        dma_queue_push(dma_q,                                                                                                   \
-                       dma_make_data(octx->dst->data,                                                                           \
-                                     vtcm_dst + spad_idx * vtcm_layout->dst_spad_half_size),                                    \
-                       cur_dst_bytes, vtcm_layout->dst_spad_half_size, cur_dst_bytes, 0);                                       \
-        dma_queue_push(dma_q,                                                                                                   \
-                       dma_make_data(vtcm_src0 + spad_idx * vtcm_layout->src0_spad_half_size, src0_data),                       \
-                       vtcm_layout->src0_spad_half_size, cur_src0_bytes, cur_src0_bytes, 1);                                    \
-    }                                                                                                                           \
-    for (uint32_t step = 0; step < ir1 - ir0; ++step) {                                                                         \
-        const uint32_t i = ir0 + step;                                                                                          \
-        void * dst_spad = (void *) dma_queue_pop(dma_q).src;                                                                    \
-        void * src_spad = (void *) dma_queue_pop(dma_q).dst;                                                                    \
-        const uint32_t row_idx   = fastdiv(i, &kparams->div_chunks_per_row);                                                    \
-        const uint32_t chunk_idx = i - row_idx * chunks_per_row;                                                                \
-        const uint32_t i12 = fastdiv(row_idx, &kparams->div_ne10_ne11);                                                         \
-        const uint32_t rem = row_idx - i12 * ne11 * ne10;                                                                       \
-        const uint32_t i11 = fastdiv(rem, &kparams->div_ne10);                                                                  \
-        const uint32_t i10 = rem - i11 * ne10;                                                                                  \
-        const uint32_t offset = chunk_idx * chunk_size;                                                                         \
-        const uint32_t cur_elems = (offset < ne00) ? MIN(chunk_size, ne00 - offset) : 0;                                        \
-        const uint32_t cur_dst_bytes  = cur_elems * sizeof(float);                                                              \
-        htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_COMP, i);                                                                   \
-        COMPUTE_EXPR;                                                                                                           \
-        htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_COMP, i);                                                                    \
-        const dma_addr_t dst_data  = octx->dst->data + i10*nb1 + i11*nb2 + i12*nb3 + offset * sizeof(float);                    \
-        dma_queue_push(dma_q,                                                                                                   \
-                       dma_make_data(dst_data, dst_spad),                                                                       \
-                       cur_dst_bytes, vtcm_layout->dst_spad_half_size, cur_dst_bytes, 1);                                       \
-        const uint32_t next_step = step + 2;                                                                                    \
-        if (next_step < ir1 - ir0) {                                                                                            \
-            const uint32_t pi = ir0 + next_step;                                                                                \
-            const uint32_t prow_idx   = fastdiv(pi, &kparams->div_chunks_per_row);                                              \
-            const uint32_t pchunk_idx = pi - prow_idx * chunks_per_row;                                                         \
-            const uint32_t pi12 = fastdiv(prow_idx, &kparams->div_ne10_ne11);                                                   \
-            const uint32_t prem = prow_idx - pi12 * ne11 * ne10;                                                                \
-            const uint32_t pi11 = fastdiv(prem, &kparams->div_ne10);                                                            \
-            const uint32_t pi10 = prem - pi11 * ne10;                                                                           \
-            const IDX_TYPE * psrc1_ptr = (const IDX_TYPE *)(uintptr_t)(octx->src[1]->data + pi10*nb10 + pi11*nb11 + pi12*nb12); \
-            const uint32_t pi01 = (uint32_t)*psrc1_ptr;                                                                         \
-            assert(pi01 < ne01);                                                                                                \
-            const uint32_t pq02 = fastdiv(pi11, &kparams->div_ne02);                                                            \
-            const uint32_t pi02 = pi11 - pq02 * ne02;                                                                           \
-            const uint32_t pq03 = fastdiv(pi12, &kparams->div_ne03);                                                            \
-            const uint32_t pi03 = pi12 - pq03 * ne03;                                                                           \
-            const uint32_t poffset = pchunk_idx * chunk_size;                                                                   \
-            const uint32_t pcur_elems = (poffset < ne00) ? MIN(chunk_size, ne00 - poffset) : 0;                                 \
-            const uint32_t pcur_src0_bytes = SRC0_SIZE_EXPR(pcur_elems);                                                        \
-            const dma_addr_t psrc0_data =                                                                                       \
-                octx->src[0]->data + pi01*nb01 + pi02*nb02 + pi03*nb03 + SRC0_SIZE_EXPR(poffset);                               \
-            dma_queue_push(dma_q,                                                                                               \
-                           dma_make_data(src_spad, psrc0_data),                                                                 \
-                           vtcm_layout->src0_spad_half_size, pcur_src0_bytes, pcur_src0_bytes, 1);                              \
-        }                                                                                                                       \
-    }                                                                                                                           \
-    dma_queue_flush(dma_q);                                                                                                     \
+static __attribute__((noinline)) uint32_t get_rows_read_index_vtcm(
+    const void * src1,
+    uint32_t idx,
+    int idx_type
+) {
+    const uint32_t off = idx * ((idx_type == HTP_TYPE_I64) ? sizeof(int64_t) : sizeof(int32_t));
+    const HVX_Vector v = hvx_vmem((const uint8_t *) src1 + off); // off & ~127u in HW
+    // vextract is normally slow, but here it avoids an HVX->L2 store and L1D miss
+    return (uint32_t) Q6_R_vextract_VR(v, off);
 }
 
-#define F16_BYTES(n)  ((n) * sizeof(__fp16))
-#define Q8_0_BYTES(n) (((n) / 32) * sizeof(block_q8_0))
+static void get_rows_thread_sametype(unsigned int nth, unsigned int ith, void * data) {
+    struct get_rows_context * grctx = (struct get_rows_context *) data;
+    struct htp_ops_context * octx = grctx->octx;
+    const struct htp_get_rows_kernel_params * kparams = grctx->kparams;
+    get_rows_preamble;
 
-static __attribute__((noinline)) void compute_get_rows_f16(float * dst_spad, const void * src_spad, uint32_t cur_elems) {
-    hvx_dequantize_row_f16_f32(dst_spad, src_spad, cur_elems);
-}
+    const uint32_t dr  = grctx->tasks_per_thread;
+    const uint32_t ir0 = grctx->task_start + dr * ith;
+    if (ir0 >= grctx->task_start + grctx->tasks) {
+        return;
+    }
+    const uint32_t ir1 = MIN(ir0 + dr, grctx->task_start + grctx->tasks);
+    const uint32_t row_size_bytes = grctx->src0_row_size;
+    dma_queue * dma_q = octx->ctx->dma[ith];
 
-static __attribute__((noinline)) void compute_get_rows_q8_0(float * dst_spad, const void * src_spad, uint32_t cur_elems) {
-    hvx_dequantize_row_q8_0_f32(dst_spad, src_spad, cur_elems);
-}
+    for (uint32_t i = ir0; i < ir1; ++i) {
+        const uint32_t i12 = fastdiv(i, &kparams->div_ne10_ne11);
+        const uint32_t rem = i - i12 * ne11 * ne10;
+        const uint32_t i11 = fastdiv(rem, &kparams->div_ne10);
+        const uint32_t i10 = rem - i11 * ne10;
 
-GET_ROWS_THREAD_DT_FN(f16,  F16_BYTES,  int32_t, { compute_get_rows_f16((float *)dst_spad, src_spad, cur_elems); })
-GET_ROWS_THREAD_DT_FN(f16,  F16_BYTES,  int64_t, { compute_get_rows_f16((float *)dst_spad, src_spad, cur_elems); })
+        const uint32_t i01 = grctx->read_index(grctx->src1_data, i, grctx->idx_type);
+        assert(i01 < ne01);
 
-GET_ROWS_THREAD_DT_FN(q8_0, Q8_0_BYTES, int32_t, { compute_get_rows_q8_0((float *)dst_spad, src_spad, cur_elems); })
-GET_ROWS_THREAD_DT_FN(q8_0, Q8_0_BYTES, int64_t, { compute_get_rows_q8_0((float *)dst_spad, src_spad, cur_elems); })
+        const uint32_t q02 = fastdiv(i11, &kparams->div_ne02);
+        const uint32_t i02 = i11 - q02 * ne02;
+        const uint32_t q03 = fastdiv(i12, &kparams->div_ne03);
+        const uint32_t i03 = i12 - q03 * ne03;
 
+        const dma_addr_t src0_data = octx->src[0]->data + i01*nb01 + i02*nb02 + i03*nb03;
+        const dma_addr_t dst_data  = octx->dst->data    + i10*nb1  + i11*nb2  + i12*nb3;
 
-static __attribute__((noinline)) void compute_get_rows_tiled(float * dst, const uint8_t * tile, uint32_t row, bool q4, bool q4_k) {
-    const HVX_VectorPred first2 = Q6_Q_vsetq_R(2);
-    const HVX_VectorPred first4 = Q6_Q_vsetq_R(4);
-    HVX_Vector vq = Q6_V_vzero();
-    if (q4) {
-        const HVX_VectorPred first1 = Q6_Q_vsetq_R(1);
-        const HVX_VectorPred first3 = Q6_Q_vsetq_R(3);
-        for (int group = 3; group >= 0; --group) {
-            const HVX_Vector v = Q6_V_vror_VR(hvx_vmem(tile + group * VLEN), row);
-            // Four planes contribute bytes at 0, 32, 64 and 96 after rotation.
-            HVX_Vector packed = Q6_V_vmux_QVV(first1, v, Q6_V_vror_VR(v, 31));
-            packed = Q6_V_vmux_QVV(first2, packed, Q6_V_vror_VR(v, 62));
-            packed = Q6_V_vmux_QVV(first3, packed, Q6_V_vror_VR(v, 93));
-            vq = Q6_V_vmux_QVV(first4, packed, Q6_V_vror_VR(vq, VLEN - 4));
-        }
-        const HVX_Vector lo = Q6_V_vand_VV(vq, Q6_Vb_vsplat_R(0x0F));
-        const HVX_Vector hi = Q6_Vub_vlsr_VubR(vq, 4);
-        vq = Q6_V_lo_W(Q6_W_vshuff_VVR(hi, lo, -1));
-        if (!q4_k) {
-            vq = Q6_Vb_vsub_VbVb(vq, Q6_Vb_vsplat_R(8));
-        }
-    } else {
-        for (int group = 7; group >= 0; --group) {
-            const HVX_Vector v = Q6_V_vror_VR(hvx_vmem(tile + group * VLEN), 2 * row);
-            // Two planes contribute halfwords at 0 and 64 after rotation.
-            const HVX_Vector packed = Q6_V_vmux_QVV(first2, v, Q6_V_vror_VR(v, 62));
-            vq = Q6_V_vmux_QVV(first4, packed, Q6_V_vror_VR(vq, VLEN - 4));
+        while (!dma_queue_push(dma_q, dma_make_data(dst_data, src0_data), nb1, nb01, row_size_bytes, 1)) {
+            dma_queue_pop(dma_q);
         }
     }
-    const HVX_Vector scales = hvx_vmem(tile + (q4 ? 512 : 1024));
-    const HVX_Vector scale_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, (q4_k ? 4 : 2) * row));
-    const HVX_Vector scale = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_hf));
-    const HVX_VectorPair p16 = Q6_Wh_vunpack_Vb(vq);
-    const HVX_VectorPair p32 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16));
-    HVX_Vector values = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(p32)), scale);
-    if (q4_k) {
-        const HVX_Vector offset_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 4 * row + 2));
-        const HVX_Vector offset = Q6_V_lo_W(hvx_vec_f16_to_f32(offset_hf));
-        values = hvx_vec_add_f32_f32(values, offset);
-    }
-    *(HVX_Vector *) dst = values;
+    dma_queue_flush(dma_q);
 }
 
-static __attribute__((noinline)) void compute_get_rows_q6_k(float * dst, const uint8_t * tile, uint32_t row) {
-    const HVX_VectorPred first4 = Q6_Q_vsetq_R(4);
-    const HVX_VectorPred first16 = Q6_Q_vsetq_R(16 * sizeof(float));
-    const HVX_Vector mask_0f = Q6_Vb_vsplat_R(0x0F);
-    const HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03);
-    HVX_Vector vq = Q6_V_vzero();
 
-    for (int group = 7; group >= 0; --group) {
-        const HVX_Vector lo_plane = Q6_V_vror_VR(hvx_vmem(tile + (group >> 1) * VLEN), 4 * row);
-        const HVX_Vector hi_plane = Q6_V_vror_VR(hvx_vmem(tile + 512 + (group >> 2) * VLEN), 4 * row);
-        const HVX_Vector lo = (group & 1) ? Q6_Vub_vlsr_VubR(lo_plane, 4) : Q6_V_vand_VV(lo_plane, mask_0f);
-        const HVX_Vector hi = Q6_Vub_vlsr_VubR(hi_plane, 2 * (group & 3));
-        const HVX_Vector packed = Q6_V_vor_VV(lo, Q6_Vw_vasl_VwR(Q6_V_vand_VV(hi, mask_03), 4));
-        vq = Q6_V_vmux_QVV(first4, packed, Q6_V_vror_VR(vq, VLEN - 4));
+
+struct get_rows_flat_task {
+    dma_addr_t src0_data;
+    dma_addr_t dst_data;
+};
+
+static inline struct get_rows_flat_task get_rows_flat_calc_task(
+    const struct htp_ops_context * octx,
+    const struct get_rows_context * grctx,
+    uint32_t i,
+    bool need_src0
+) {
+    const struct htp_get_rows_kernel_params * kparams = grctx->kparams;
+    get_rows_preamble;
+
+    const uint32_t i12 = fastdiv(i, &kparams->div_ne10_ne11);
+    const uint32_t rem = i - i12 * ne11 * ne10;
+    const uint32_t i11 = fastdiv(rem, &kparams->div_ne10);
+    const uint32_t i10 = rem - i11 * ne10;
+
+    struct get_rows_flat_task task;
+    task.src0_data = 0;
+    task.dst_data  = octx->dst->data + i10*nb1 + i11*nb2 + i12*nb3;
+
+    if (need_src0) {
+        const uint32_t i01 = grctx->read_index(grctx->src1_data, i, grctx->idx_type);
+        assert(i01 < ne01);
+
+        const uint32_t q02 = fastdiv(i11, &kparams->div_ne02);
+        const uint32_t i02 = i11 - q02 * ne02;
+        const uint32_t q03 = fastdiv(i12, &kparams->div_ne03);
+        const uint32_t i03 = i12 - q03 * ne03;
+
+        task.src0_data = octx->src[0]->data + i01*nb01 + i02*nb02 + i03*nb03;
     }
 
-    const HVX_Vector scales = hvx_vmem(tile + 768);
-    const HVX_Vector scale_lo_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 2 * row));
-    const HVX_Vector scale_hi_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 64 + 2 * row));
-    const HVX_Vector scale_lo = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_lo_hf));
-    const HVX_Vector scale_hi = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_hi_hf));
-    const HVX_Vector scale = Q6_V_vmux_QVV(first16, scale_lo, scale_hi);
-    const HVX_VectorPair p16 = Q6_Wh_vunpack_Vb(Q6_Vb_vsub_VbVb(vq, Q6_Vb_vsplat_R(32)));
-    const HVX_VectorPair p32 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16));
-    *(HVX_Vector *) dst = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(p32)), scale);
+    return task;
 }
 
+static void get_rows_thread_flat(unsigned int nth, unsigned int ith, void * data) {
+    struct get_rows_context * grctx = (struct get_rows_context *) data;
+    struct htp_ops_context * octx = grctx->octx;
+    get_rows_preamble;
+
+    const uint32_t dr  = grctx->tasks_per_thread;
+    const uint32_t ir0 = grctx->task_start + dr * ith;
+    if (ir0 >= grctx->task_start + grctx->tasks) {
+        return;
+    }
+    const uint32_t ir1 = MIN(ir0 + dr, grctx->task_start + grctx->tasks);
+
+    dma_queue * dma_q = octx->ctx->dma[ith];
+    struct htp_thread_trace * tr = &octx->ctx->trace[ith];
+    const struct htp_get_rows_vtcm_layout * vtcm_layout = &grctx->vtcm_layout;
+    uint8_t * vtcm_src0 = VTCM_LAYOUT_PTR(uint8_t, grctx->vtcm_base, vtcm_layout->off_src0) + ith * vtcm_layout->src0_bytes_per_thread;
+    uint8_t * vtcm_dst  = VTCM_LAYOUT_PTR(uint8_t, grctx->vtcm_base, vtcm_layout->off_dst)  + ith * vtcm_layout->dst_bytes_per_thread;
+    const uint32_t src0_bytes = grctx->src0_row_size;
+    const uint32_t dst_bytes  = grctx->dst_row_size;
+    get_rows_flat_compute_fn_t compute = grctx->compute_flat;
+
+    for (uint32_t step = 0, spad_idx = 0; step < ir1 - ir0 && spad_idx < 2; ++step, spad_idx++) {
+        const uint32_t i = ir0 + step;
+        struct get_rows_flat_task task = get_rows_flat_calc_task(octx, grctx, i, true);
+
+        dma_queue_push(dma_q,
+                       dma_make_data(octx->dst->data,
+                                     vtcm_dst + spad_idx * vtcm_layout->dst_spad_half_size),
+                       dst_bytes, vtcm_layout->dst_spad_half_size, dst_bytes, 0);
+        dma_queue_push(dma_q,
+                       dma_make_data(vtcm_src0 + spad_idx * vtcm_layout->src0_spad_half_size, task.src0_data),
+                       vtcm_layout->src0_spad_half_size, src0_bytes, src0_bytes, 1);
+    }
+
+    for (uint32_t step = 0; step < ir1 - ir0; ++step) {
+        const uint32_t i = ir0 + step;
+        void * dst_spad = (void *) dma_queue_pop(dma_q).src;
+        void * src_spad = (void *) dma_queue_pop(dma_q).dst;
+
+        struct get_rows_flat_task task = get_rows_flat_calc_task(octx, grctx, i, false);
+
+        htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_COMP, (uint16_t) i);
+        compute((float *) dst_spad, src_spad, ne00);
+        htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_COMP, (uint16_t) i);
+
+        dma_queue_push(dma_q,
+                       dma_make_data(task.dst_data, dst_spad),
+                       dst_bytes, vtcm_layout->dst_spad_half_size, dst_bytes, 1);
+
+        const uint32_t next_step = step + 2;
+        if (next_step < ir1 - ir0) {
+            const uint32_t pi = ir0 + next_step;
+            struct get_rows_flat_task next_task = get_rows_flat_calc_task(octx, grctx, pi, true);
+            dma_queue_push(dma_q,
+                           dma_make_data(src_spad, next_task.src0_data),
+                           vtcm_layout->src0_spad_half_size, src0_bytes, src0_bytes, 1);
+        }
+    }
+
+    dma_queue_flush(dma_q);
+}
+
+
+
 struct get_rows_tiled_task {
     dma_addr_t tile_src_base;
     dma_addr_t dst_data;
@@ -307,16 +258,15 @@ static inline struct get_rows_tiled_task get_rows_tiled_calc_task(
     const uint32_t rem = i - i12 * ne11 * ne10;
     const uint32_t i11 = fastdiv(rem, &kparams->div_ne10);
     const uint32_t i10 = rem - i11 * ne10;
-    const dma_addr_t src1_data = octx->src[1]->data + i10*nb10 + i11*nb11 + i12*nb12;
-    const uint32_t i01 = grctx->index_i32 ? *(const int32_t *)(uintptr_t) src1_data : (uint32_t) *(const int64_t *)(uintptr_t) src1_data;
+    const uint32_t i01 = grctx->read_index(grctx->src1_data, i, grctx->idx_type);
     assert(i01 < ne01);
 
     const uint32_t q02 = fastdiv(i11, &kparams->div_ne02);
     const uint32_t i02 = i11 - q02 * ne02;
     const uint32_t q03 = fastdiv(i12, &kparams->div_ne03);
     const uint32_t i03 = i12 - q03 * ne03;
-    const uint32_t column_tile = i01 / HTP_MM_HMX_TILE_N_ROWS;
-    const uint32_t row = i01 % HTP_MM_HMX_TILE_N_ROWS;
+    const uint32_t column_tile = i01 / HTP_TENSOR_TILE_N_ROWS;
+    const uint32_t row = i01 % HTP_TENSOR_TILE_N_ROWS;
     const dma_addr_t matrix = octx->src[0]->data + i02*nb02 + i03*nb03;
 
     struct get_rows_tiled_task task;
@@ -329,7 +279,6 @@ static inline struct get_rows_tiled_task get_rows_tiled_calc_task(
 static void get_rows_thread_tiled(unsigned int nth, unsigned int ith, void * data) {
     struct get_rows_context * grctx = (struct get_rows_context *) data;
     struct htp_ops_context * octx = grctx->octx;
-    const struct htp_get_rows_kernel_params * kparams = grctx->kparams;
     get_rows_preamble;
 
     const uint32_t dr  = grctx->tasks_per_thread;
@@ -339,19 +288,17 @@ static void get_rows_thread_tiled(unsigned int nth, unsigned int ith, void * dat
     }
 
     const uint32_t ir1 = MIN(ir0 + dr, grctx->task_start + grctx->tasks);
-    const uint32_t n_k_tiles = ne00 / HTP_MM_HMX_TILE_N_COLS;
+    const uint32_t n_k_tiles = ne00 / HTP_TENSOR_TILE_N_COLS;
     const struct htp_get_rows_vtcm_layout * vtcm_layout = &grctx->vtcm_layout;
-    uint8_t * src_spad_base = grctx->vtcm_base + vtcm_layout->off_src0 + ith * vtcm_layout->src0_bytes_per_thread;
-    uint8_t * dst_spad_base = grctx->vtcm_base + vtcm_layout->off_dst + ith * vtcm_layout->dst_bytes_per_thread;
+    uint8_t * src_spad_base = VTCM_LAYOUT_PTR(uint8_t, grctx->vtcm_base, vtcm_layout->off_src0) + ith * vtcm_layout->src0_bytes_per_thread;
+    uint8_t * dst_spad_base = VTCM_LAYOUT_PTR(uint8_t, grctx->vtcm_base, vtcm_layout->off_dst)  + ith * vtcm_layout->dst_bytes_per_thread;
     dma_queue * dma_q = octx->ctx->dma[ith];
     struct htp_thread_trace * tr = &octx->ctx->trace[ith];
 
     const uint32_t tile_size   = grctx->tile_size;
     const uint32_t tile_stride = grctx->tile_stride;
-    const uint32_t dst_bytes   = ne00 * sizeof(float);
-    const bool is_q4 = octx->src[0]->type == HTP_TYPE_Q4_0 || octx->src[0]->type == HTP_TYPE_Q4_K;
-    const bool is_q4_k = octx->src[0]->type == HTP_TYPE_Q4_K;
-    const bool is_q6_k = octx->src[0]->type == HTP_TYPE_Q6_K;
+    const uint32_t dst_bytes   = grctx->dst_row_size;
+    get_rows_tiled_compute_fn_t compute_tiled = grctx->compute_tiled;
 
     for (uint32_t step = 0, spad_idx = 0; step < ir1 - ir0 && spad_idx < 2; ++step, ++spad_idx) {
         const uint32_t i = ir0 + step;
@@ -376,15 +323,7 @@ static void get_rows_thread_tiled(unsigned int nth, unsigned int ith, void * dat
         struct get_rows_tiled_task task = get_rows_tiled_calc_task(octx, grctx, i, n_k_tiles, tile_size);
 
         htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_COMP, (uint16_t) i);
-        for (uint32_t k_tile = 0; k_tile < n_k_tiles; ++k_tile) {
-            const uint8_t * tile = src_spad + k_tile * tile_stride;
-            float * dst_block = dst_spad + k_tile * HTP_MM_HMX_TILE_N_COLS;
-            if (is_q6_k) {
-                compute_get_rows_q6_k(dst_block, tile, task.row);
-            } else {
-                compute_get_rows_tiled(dst_block, tile, task.row, is_q4, is_q4_k);
-            }
-        }
+        compute_tiled(dst_spad, src_spad, n_k_tiles, tile_stride, task.row);
         htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_COMP, (uint16_t) i);
 
         // Real writeback of dst_spad
@@ -405,16 +344,40 @@ static void get_rows_thread_tiled(unsigned int nth, unsigned int ith, void * dat
     dma_queue_flush(dma_q);
 }
 
+static __attribute__((noinline)) void get_rows_dma_src1(
+    struct htp_ops_context * octx,
+    uint8_t * vtcm_src1,
+    uint32_t elem_size
+) {
+    const uint32_t ne10 = octx->src[1]->ne[0];
+    const uint32_t ne11 = octx->src[1]->ne[1];
+    const uint32_t ne12 = octx->src[1]->ne[2];
+    const uint32_t ne13 = octx->src[1]->ne[3];
+    const uint32_t nb11 = octx->src[1]->nb[1];
+    const uint32_t nb12 = octx->src[1]->nb[2];
+    const uint32_t nb13 = octx->src[1]->nb[3];
+
+    dma_queue * dma_q = octx->ctx->dma[0];
+    uint8_t * cur_dst = vtcm_src1;
+    for (uint32_t i13 = 0; i13 < ne13; ++i13) {
+        for (uint32_t i12 = 0; i12 < ne12; ++i12) {
+            const dma_addr_t src_plane = octx->src[1]->data + i12 * nb12 + i13 * nb13;
+            while (!dma_queue_push(dma_q, dma_make_data(cur_dst, src_plane),
+                                   ne10 * elem_size, nb11, ne10 * elem_size, ne11)) {
+                dma_queue_pop(dma_q);
+            }
+            cur_dst += ne10 * ne11 * elem_size;
+        }
+    }
+    dma_queue_flush(dma_q);
+}
+
 int op_get_rows(struct htp_ops_context * octx) {
     const struct htp_get_rows_kernel_params * kparams = (const struct htp_get_rows_kernel_params *) octx->kernel_params;
 
-    if (octx->src[0]->type != HTP_TYPE_F32 &&
-        octx->src[0]->type != HTP_TYPE_F16 &&
-        octx->src[0]->type != HTP_TYPE_Q4_0 &&
-        octx->src[0]->type != HTP_TYPE_Q4_K &&
-        octx->src[0]->type != HTP_TYPE_Q6_K &&
-        octx->src[0]->type != HTP_TYPE_Q8_0 &&
-        octx->src[0]->type != HTP_TYPE_I32) {
+    const uint32_t ne00          = octx->src[0]->ne[0];
+    const uint32_t src0_row_size = htp_tensor_get_row_size(octx->src[0]->type, ne00);
+    if (src0_row_size == 0) {
         return HTP_STATUS_NO_SUPPORT;
     }
 
@@ -432,20 +395,16 @@ int op_get_rows(struct htp_ops_context * octx) {
         return HTP_STATUS_NO_SUPPORT;
     }
 
-    if (htp_tensor_is_extended(octx->src[1])) {
-        return HTP_STATUS_NO_SUPPORT;
-    }
-
     const struct htp_tensor * dst = octx->dst;
     const uint32_t total_tasks    = kparams->total_tasks;
-    const size_t dst_row_size     = htp_tensor_get_row_size(dst->type, dst->ne[0]);
+    const uint32_t dst_row_size   = (kparams->kernel_type == HTP_GET_ROWS_KERNEL_SAMETYPE) ? src0_row_size : ne00 * sizeof(float);
 
     uint32_t task_start = 0;
     uint32_t tasks      = total_tasks;
 
     if (octx->ctx->mdev.count > 1) {
         uint32_t tasks_per_chunk = 1;
-        htp_tensor_mdev_rows_per_chunk(dst, dst_row_size / dst->ne[0], (uint32_t) dst_row_size, &tasks_per_chunk);
+        htp_tensor_mdev_rows_per_chunk(dst, dst_row_size / dst->ne[0], dst_row_size, &tasks_per_chunk);
         const struct htp_tensor_mdev_range range = htp_tensor_mdev_partition(total_tasks, tasks_per_chunk, octx->ctx->mdev.idx, octx->ctx->mdev.count, &octx->ctx->mdev.count_div);
         task_start = range.start;
         tasks      = range.count;
@@ -465,15 +424,18 @@ int op_get_rows(struct htp_ops_context * octx) {
     grctx.octx = octx;
     grctx.kparams = kparams;
     grctx.vtcm_base = (uint8_t *)octx->ctx->vtcm_base;
+    grctx.compute_flat = NULL;
+    grctx.compute_tiled = NULL;
     grctx.task_start = task_start;
     grctx.tasks = tasks;
     grctx.tasks_per_thread = octx->ctx->mdev.count == 1 ? kparams->tasks_per_thread : fastdiv(tasks + n_threads - 1, &octx->n_threads_div);
-    grctx.tile_size = htp_mm_get_weight_tile_size(octx->src[0]->type);
-    grctx.tile_stride = (grctx.tile_size + 127) & ~127;
-    grctx.index_i32 = octx->src[1]->type == HTP_TYPE_I32;
+    grctx.tile_size   = 0;
+    grctx.tile_stride = 0;
+    grctx.src0_row_size = src0_row_size;
+    grctx.dst_row_size  = dst_row_size;
+    grctx.idx_type      = octx->src[1]->type;
 
-    const uint32_t ne00 = octx->src[0]->ne[0];
-    htp_get_rows_vtcm_layout_build(&grctx.vtcm_layout, kparams->kernel_type, octx->src[0]->type, ne00, n_threads);
+    htp_get_rows_vtcm_layout_build(&grctx.vtcm_layout, kparams->kernel_type, octx->src[0]->type, ne00, total_tasks, grctx.idx_type, n_threads);
 
     if (grctx.vtcm_layout.total_bytes > octx->ctx->vtcm_size) {
         FARF(ERROR, "get-rows: VTCM reservation %zu is too small, needed %zu\n",
@@ -481,27 +443,56 @@ int op_get_rows(struct htp_ops_context * octx) {
         return HTP_STATUS_INVAL_PARAMS;
     }
 
-    const bool is_i32 = (octx->src[1]->type == HTP_TYPE_I32);
-
     work_queue_func_t q_func = NULL;
     switch (kparams->kernel_type) {
         case HTP_GET_ROWS_KERNEL_SAMETYPE:
-            q_func = (work_queue_func_t)(is_i32 ? get_rows_thread_st_int32_t : get_rows_thread_st_int64_t);
+            q_func = get_rows_thread_sametype;
             break;
         case HTP_GET_ROWS_KERNEL_TILED:
+            switch (octx->src[0]->type) {
+                case HTP_TYPE_F16:  grctx.compute_tiled = hvx_dequantize_row_tiled_f16_f32; break;
+                case HTP_TYPE_Q8_0: grctx.compute_tiled = hvx_dequantize_row_tiled_q8_0_f32; break;
+                case HTP_TYPE_Q4_0: grctx.compute_tiled = hvx_dequantize_row_tiled_q4_0_f32; break;
+                case HTP_TYPE_Q4_K: grctx.compute_tiled = hvx_dequantize_row_tiled_q4_k_f32; break;
+                case HTP_TYPE_Q5_K: grctx.compute_tiled = hvx_dequantize_row_tiled_q5_k_f32; break;
+                case HTP_TYPE_Q6_K: grctx.compute_tiled = hvx_dequantize_row_tiled_q6_k_f32; break;
+                default:            return HTP_STATUS_NO_SUPPORT;
+            }
+            grctx.tile_size   = htp_tensor_get_tile_size(octx->src[0]->type);
+            grctx.tile_stride = htp_tensor_get_tile_padded_size(octx->src[0]->type);
             q_func = get_rows_thread_tiled;
             break;
         case HTP_GET_ROWS_KERNEL_FLAT:
             switch (octx->src[0]->type) {
-                case HTP_TYPE_F16:  q_func = (work_queue_func_t)(is_i32 ? get_rows_thread_f16_int32_t  : get_rows_thread_f16_int64_t);  break;
-                case HTP_TYPE_Q8_0: q_func = (work_queue_func_t)(is_i32 ? get_rows_thread_q8_0_int32_t : get_rows_thread_q8_0_int64_t); break;
+                case HTP_TYPE_F16:  grctx.compute_flat = hvx_dequantize_row_flat_f16_f32;  break;
+                case HTP_TYPE_BF16: grctx.compute_flat = hvx_dequantize_row_flat_bf16_f32; break;
+                case HTP_TYPE_Q8_0: grctx.compute_flat = hvx_dequantize_row_flat_q8_0_f32; break;
+                case HTP_TYPE_Q4_0: grctx.compute_flat = hvx_dequantize_row_flat_q4_0_f32; break;
+                case HTP_TYPE_Q4_K: grctx.compute_flat = hvx_dequantize_row_flat_q4_k_f32; break;
+                case HTP_TYPE_Q5_K: grctx.compute_flat = hvx_dequantize_row_flat_q5_k_f32; break;
+                case HTP_TYPE_Q6_K: grctx.compute_flat = hvx_dequantize_row_flat_q6_k_f32; break;
                 default:            return HTP_STATUS_NO_SUPPORT;
             }
+            q_func = get_rows_thread_flat;
             break;
         default:
             return HTP_STATUS_NO_SUPPORT;
     }
 
+    const uint32_t elem_size = (grctx.idx_type == HTP_TYPE_I64) ? sizeof(int64_t) : sizeof(int32_t);
+    const size_t src1_bytes  = (size_t) total_tasks * elem_size;
+
+    if (!htp_tensor_is_extended(octx->src[1]) && src1_bytes <= 128 &&
+        htp_tensor_is_contiguous(octx->src[1], elem_size)) {
+        grctx.src1_data  = (const void *) (uintptr_t) octx->src[1]->data;
+        grctx.read_index = get_rows_read_index_ddr;
+    } else {
+        uint8_t * vtcm_src1 = VTCM_LAYOUT_PTR(uint8_t, grctx.vtcm_base, grctx.vtcm_layout.off_src1);
+        get_rows_dma_src1(octx, vtcm_src1, elem_size);
+        grctx.src1_data  = vtcm_src1;
+        grctx.read_index = get_rows_read_index_vtcm;
+    }
+
     FARF(HIGH, "get-rows: (%ux%ux%ux%u) x (%ux%ux%ux%u) -> (%ux%ux%ux%u) : src0-vtcm-size %zu dst-vtcm-size %zu kernel-type %d n-threads %d\n",
          octx->src[0]->ne[0], octx->src[0]->ne[1], octx->src[0]->ne[2], octx->src[0]->ne[3],
          octx->src[1]->ne[0], octx->src[1]->ne[1], octx->src[1]->ne[2], octx->src[1]->ne[3],

ggml/src/ggml-hexagon/htp/get-rows-ops.h

@@ -5,7 +5,8 @@
 #include <string.h>
 
 #include "hex-fastdiv.h"
-#include "matmul-ops.h"
+#include "htp-tensor.h"
+#include "htp-vtcm.h"
 
 enum htp_get_rows_kernel_type {
     HTP_GET_ROWS_KERNEL_SAMETYPE = 0,
@@ -16,8 +17,6 @@ enum htp_get_rows_kernel_type {
 struct htp_get_rows_kernel_params {
     int32_t  n_threads;
     int32_t  kernel_type;
-    int32_t  chunks_per_row;
-    int32_t  chunk_size;
     int32_t  total_tasks;
     int32_t  tasks_per_thread;
     int32_t  vtcm_size;
@@ -25,13 +24,14 @@ struct htp_get_rows_kernel_params {
     // Fastdiv helpers
     struct fastdiv_values div_ne10;
     struct fastdiv_values div_ne10_ne11;
-    struct fastdiv_values div_chunks_per_row;
     struct fastdiv_values div_ne02;
     struct fastdiv_values div_ne03;
 };
 
 struct htp_get_rows_vtcm_layout {
     size_t total_bytes;
+    size_t off_src1;
+    size_t src1_bytes;
     size_t off_src0;
     size_t off_dst;
 
@@ -43,57 +43,44 @@ struct htp_get_rows_vtcm_layout {
 };
 
 static inline void htp_get_rows_vtcm_layout_build(
-    struct htp_get_rows_vtcm_layout * vtcm_layout,
+    struct htp_get_rows_vtcm_layout * L,
     int kernel_type,
     int type,
     uint32_t ne00,
+    uint32_t total_tasks,
+    int idx_type,
     uint32_t n_threads) {
 
-    if (kernel_type == HTP_GET_ROWS_KERNEL_SAMETYPE) {
-        memset(vtcm_layout, 0, sizeof(*vtcm_layout));
-        return;
+    memset(L, 0, sizeof(*L));
+
+    const size_t elem_size = (idx_type == HTP_TYPE_I64) ? sizeof(int64_t) : sizeof(int32_t);
+    const size_t src1_raw_bytes = total_tasks * elem_size;
+    const size_t src1_bytes = ((src1_raw_bytes + 127) & ~127) + 128;
+
+    L->src1_bytes = src1_bytes;
+    const size_t src1_alloc_bytes = (src1_bytes + 255) & ~255;
+
+    if (kernel_type != HTP_GET_ROWS_KERNEL_SAMETYPE) {
+        size_t src0_row_size = 0;
+        if (kernel_type == HTP_GET_ROWS_KERNEL_TILED) {
+            const size_t tile_stride = htp_tensor_get_tile_padded_size(type);
+            const uint32_t n_k_tiles = ne00 / HTP_TENSOR_TILE_N_COLS;
+            src0_row_size = n_k_tiles > 0 ? (n_k_tiles * tile_stride) : tile_stride;
+        } else {
+            src0_row_size = htp_tensor_get_row_size(type, ne00);
+        }
+
+        L->src0_spad_half_size   = (src0_row_size + 255) & ~255;
+        L->dst_spad_half_size    = (ne00 * sizeof(float) + 255) & ~255;
+        L->src0_bytes_per_thread = 2 * L->src0_spad_half_size;
+        L->dst_bytes_per_thread  = 2 * L->dst_spad_half_size;
     }
 
-    if (kernel_type == HTP_GET_ROWS_KERNEL_TILED) {
-        const size_t tile_size   = htp_mm_get_weight_tile_size(type);
-        const size_t tile_stride = (tile_size + 127) & ~127;
-        const uint32_t n_k_tiles = ne00 / HTP_MM_HMX_TILE_N_COLS;
-        const size_t row_tiles_size = n_k_tiles > 0 ? (n_k_tiles * tile_stride) : tile_stride;
-        vtcm_layout->src0_spad_half_size = (row_tiles_size + 255) & ~255;
-        vtcm_layout->dst_spad_half_size  = (ne00 * sizeof(float) + 255) & ~255;
-        vtcm_layout->src0_bytes_per_thread = 2 * vtcm_layout->src0_spad_half_size;
-        vtcm_layout->dst_bytes_per_thread  = 2 * vtcm_layout->dst_spad_half_size;
-        vtcm_layout->off_src0 = 0;
-        vtcm_layout->off_dst  = vtcm_layout->src0_bytes_per_thread * n_threads;
-        vtcm_layout->total_bytes = vtcm_layout->off_dst + vtcm_layout->dst_bytes_per_thread * n_threads;
-        return;
-    }
-
-    uint32_t src0_row_size = 0;
-    switch (type) {
-        case HTP_TYPE_F16:
-            src0_row_size = ne00 * 2;
-            break;
-        case HTP_TYPE_Q8_0:
-            src0_row_size = (ne00 / 32) * 34;
-            break;
-        default:
-            src0_row_size = 0;
-            break;
-    }
-
-    size_t src0_row_size_aligned = (src0_row_size + 255) & ~255;
-    size_t dst_row_size_aligned  = (ne00 * sizeof(float) + 255) & ~255;
-
-    vtcm_layout->src0_spad_half_size = src0_row_size_aligned;
-    vtcm_layout->dst_spad_half_size  = dst_row_size_aligned;
-
-    vtcm_layout->src0_bytes_per_thread = src0_row_size_aligned * 2;
-    vtcm_layout->dst_bytes_per_thread  = dst_row_size_aligned * 2;
-
-    vtcm_layout->off_src0 = 0;
-    vtcm_layout->off_dst  = vtcm_layout->off_src0 + vtcm_layout->src0_bytes_per_thread * n_threads;
-    vtcm_layout->total_bytes = vtcm_layout->off_dst + vtcm_layout->dst_bytes_per_thread * n_threads;
+    size_t off = 0;
+    VTCM_LAYOUT_ALLOC(off, off_src1, src1_alloc_bytes);
+    VTCM_LAYOUT_ALLOC(off, off_src0, L->src0_bytes_per_thread * n_threads);
+    VTCM_LAYOUT_ALLOC(off, off_dst,  L->dst_bytes_per_thread  * n_threads);
+    L->total_bytes = off;
 }
 
 #if defined(__cplusplus)

ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h

@@ -14,11 +14,15 @@ static const __fp16 iq4_nl_to_fp16_lut[64] __attribute__((aligned(VLEN))) = {
 
 // --- tiled format dequantizers ---
 
-typedef struct {
+typedef struct tiled_dequantize_state tiled_dequantize_state_t;
+typedef void (*tiled_dequant_task_fn_t)(const tiled_dequantize_state_t * state, uint32_t start_tile, uint32_t end_tile);
+
+struct tiled_dequantize_state {
     struct htp_context      * ctx;
     struct htp_thread_trace * traces;
     __fp16                  * dst;
     const uint8_t           * src;
+    tiled_dequant_task_fn_t   task_fn;
 
     struct fastdiv_values     n_k_tiles_div;
     uint32_t                  n_k_tiles;
@@ -31,7 +35,7 @@ typedef struct {
     uint32_t                  k_block;
     size_t                    row_stride;
     uint32_t                  weight_type;
-} tiled_dequantize_state_t;
+};
 
 // Dequantize a single tile from tiled weight data (already in VTCM) to tile-major FP16.
 static void dequantize_tiled_weight_to_fp16_task_q4_0(
@@ -611,7 +615,7 @@ static void dequantize_tiled_weight_to_fp16_task_q5_k(
     }
 }
 
-// Q6_K stores 6-bit weights and one fp16 scale per 16 k, see HTP_MM_WEIGHT_TILE_SIZE_Q6_K.
+// Q6_K stores 6-bit weights and one fp16 scale per 16 k, see HTP_TENSOR_TILE_SIZE_Q6_K.
 // A k-group holds 4 k per row, the HMX tile holds 2, so each group is dealt into two tiles.
 static void dequantize_tiled_weight_to_fp16_task_q6_k(
         const tiled_dequantize_state_t *state,
@@ -669,7 +673,7 @@ static void dequantize_tiled_weight_to_fp16_task_q6_k(
     }
 }
 
-// Q3_K stores 3-bit weights and one fp16 scale per 16 k, see HTP_MM_WEIGHT_TILE_SIZE_Q3_K.
+// Q3_K stores 3-bit weights and one fp16 scale per 16 k, see HTP_TENSOR_TILE_SIZE_Q3_K.
 static void dequantize_tiled_weight_to_fp16_task_q3_k(
         const tiled_dequantize_state_t *state,
         uint32_t start_tile, uint32_t end_tile) {
@@ -701,7 +705,7 @@ static void dequantize_tiled_weight_to_fp16_task_q3_k(
     }
 }
 
-// Q2_K stores 2-bit weights, the ggml scale bytes and d / -dmin per super-block, see HTP_MM_WEIGHT_TILE_SIZE_Q2_K.
+// Q2_K stores 2-bit weights, the ggml scale bytes and d / -dmin per super-block, see HTP_TENSOR_TILE_SIZE_Q2_K.
 static void dequantize_tiled_weight_to_fp16_task_q2_k(
         const tiled_dequantize_state_t *state,
         uint32_t start_tile, uint32_t end_tile) {
@@ -837,6 +841,56 @@ void quantize_f32_weight_to_fp16_tiles_task(
     }
 }
 
+static __attribute__((noinline))
+void convert_bf16_weight_to_fp16_tiles_task(
+        const tiled_dequantize_state_t *state,
+        uint32_t start_tile, uint32_t end_tile) {
+
+    const uint32_t n_k_tiles = state->n_k_tiles;
+    const struct fastdiv_values n_k_tiles_div = state->n_k_tiles_div;
+
+    const HVX_Vector v_scat_base  = hvx_vmem(hmx_transpose_scatter_offsets);
+    const HVX_Vector v_scat_step  = Q6_V_vsplat_R(4);
+    const HVX_VectorPred q_mask64 = Q6_Q_vsetq_R(64);
+
+    unsigned ct = fastdiv((unsigned)start_tile, &n_k_tiles_div);
+    unsigned kt = fastmodulo((unsigned)start_tile, n_k_tiles, &n_k_tiles_div);
+
+    for (unsigned t = start_tile; t < (unsigned)end_tile; ) {
+        if (kt >= (unsigned)n_k_tiles) { kt = 0; ct++; }
+
+        __fp16 *tile_base = state->dst + t * HTP_MM_HMX_TILE_N_ELMS;
+        {
+            uint32_t byte_off = kt * 32 * sizeof(uint16_t);
+
+            HVX_Vector v_off = v_scat_base;
+            for (uint32_t r = 0; r < HTP_MM_HMX_TILE_N_ROWS; r += 2) {
+                uint32_t row0 = ct * HTP_MM_HMX_TILE_N_COLS + r;
+                uint32_t row1 = row0 + 1;
+
+                const uint8_t *r0 = state->src + row0 * state->row_stride;
+                const uint8_t *r1 = state->src + row1 * state->row_stride;
+
+                HVX_Vector v0_raw = (row0 < state->n_cols) ? hvx_vmemu((const void *)(r0 + byte_off)) : Q6_V_vzero();
+                HVX_Vector v1_raw = (row1 < state->n_cols) ? hvx_vmemu((const void *)(r1 + byte_off)) : Q6_V_vzero();
+
+                HVX_Vector v0_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(Q6_Ww_vunpack_Vh(v0_raw)), 16);
+                HVX_Vector v1_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(Q6_Ww_vunpack_Vh(v1_raw)), 16);
+
+                HVX_Vector v_out = hvx_vec_f32_to_f16(v0_f32, v1_f32);
+
+                Q6_vscatter_QRMVwV(q_mask64, (size_t)tile_base, HTP_MM_HMX_TILE_SIZE - 1, v_off, v_out);
+                v_off = Q6_Vw_vadd_VwVw(v_off, v_scat_step);
+
+                HVX_Vector v_out_hi = Q6_V_vror_VR(v_out, 64);
+                Q6_vscatter_QRMVwV(q_mask64, (size_t)tile_base, HTP_MM_HMX_TILE_SIZE - 1, v_off, v_out_hi);
+                v_off = Q6_Vw_vadd_VwVw(v_off, v_scat_step);
+            }
+        }
+        ++t; ++kt;
+    }
+}
+
 // --- End tiled dequantizers ---
 
 // dot-chunk functions require external HMX lock

ggml/src/ggml-hexagon/htp/htp-ctx.h

@@ -3,6 +3,7 @@
 
 #include "dma-queue.h"
 #include "hmx-queue.h"
+#include "htp-dcache.h"
 #include "htp-ops.h"
 #include "hex-profile.h"
 #include "work-queue.h"
@@ -18,8 +19,7 @@
 #define HTP_MAX_NTHREADS 10
 #endif
 
-#define HTP_MAX_MMAPS        64
-#define HTP_MAX_DIRTY_RANGES 64
+#define HTP_MAX_MMAPS 64
 
 // Memory mapping
 struct htp_mmap {
@@ -29,11 +29,6 @@ struct htp_mmap {
     uint32_t flags;
 };
 
-struct htp_dirty_range {
-    uint32_t start;
-    uint32_t end;
-};
-
 // Scratchpad state
 struct htp_spad {
     const struct htp_tensor * src;             // original src of the data (for reuse)

ggml/src/ggml-hexagon/htp/htp-dcache.c

@@ -1,4 +1,4 @@
-#include "htp-tensor.h"
+#include "htp-dcache.h"
 
 #include <qurt.h>
 #include <qurt_memory.h>
@@ -94,7 +94,7 @@ static void merge_dirty_ranges(struct htp_context * ctx) {
     }
 }
 
-void htp_tensor_dirty_all(struct htp_context * ctx, const struct htp_tensor * const * tensors, uint32_t n) {
+void htp_dcache_mark_dirty_tensors(struct htp_context * ctx, const struct htp_tensor * const * tensors, uint32_t n) {
     const struct htp_tensor * pending[HTP_OP_MAX_OUTPUTS];
     uint32_t n_pending = 0;
 
@@ -301,7 +301,7 @@ static void flush_dirty_ranges(struct htp_context * ctx, const struct htp_dirty_
     }
 }
 
-void htp_flush_dirty_ranges(struct htp_context * ctx) {
+void htp_dcache_flush_dirty_ranges(struct htp_context * ctx) {
     struct htp_dirty_range ranges[HTP_MAX_DIRTY_RANGES];
     uint32_t n_ranges = 0;
     uint64_t total_dirty = 0;
@@ -328,7 +328,7 @@ void htp_flush_dirty_ranges(struct htp_context * ctx) {
     memset(ctx->dirty_ranges, 0, sizeof(ctx->dirty_ranges));
 }
 
-void htp_tensor_flush_all(struct htp_context * ctx, const struct htp_tensor * const * tensors, uint32_t n) {
+void htp_dcache_flush_dirty_tensors(struct htp_context * ctx, const struct htp_tensor * const * tensors, uint32_t n) {
     const struct htp_tensor * dirty_tensors[HTP_OP_MAX_INPUTS];
     struct htp_dirty_range ranges[HTP_OP_MAX_INPUTS];
     uint32_t n_dirty = 0;

ggml/src/ggml-hexagon/htp/htp-dcache.h

@@ -0,0 +1,20 @@
+#ifndef HTP_DCACHE_H
+#define HTP_DCACHE_H
+
+#include <stdint.h>
+#include "htp-ops.h"
+
+#define HTP_MAX_DIRTY_RANGES 64
+
+struct htp_dirty_range {
+    uint32_t start;
+    uint32_t end;
+};
+
+struct htp_context;
+
+void htp_dcache_flush_dirty_ranges(struct htp_context * ctx);
+void htp_dcache_flush_dirty_tensors(struct htp_context * ctx, const struct htp_tensor * const * tensors, uint32_t n);
+void htp_dcache_mark_dirty_tensors(struct htp_context * ctx, const struct htp_tensor * const * tensors, uint32_t n);
+
+#endif // HTP_DCACHE_H

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

@@ -30,21 +30,14 @@ enum htp_data_type {
     HTP_TYPE_IQ4_NL = 20,
     HTP_TYPE_I32    = 26,
     HTP_TYPE_I64    = 27,
+    HTP_TYPE_BF16   = 30,
     HTP_TYPE_MXFP4  = 39,
 
-    // types used internally for repack, dyn.quant, etc
-    HTP_TYPE_Q4_0_TILED = 200,
-    HTP_TYPE_Q4_1_TILED,
-    HTP_TYPE_Q8_0_TILED,
-    HTP_TYPE_MXFP4_TILED,
-
     HTP_TYPE_INVALID
 };
 
-// Constats for internal types
-#define QK_Q4_0_TILED  256  // 32x32 Q4_0 tiled layout
+// Constants for internal types
 #define QK_Q8_0_TILED  128  // 32x32 Q8_0 tiled layout
-#define QK_MXFP4_TILED 256  // 32x32 MXFP4 tiled layout
 
 
 
@@ -141,7 +134,7 @@ enum htp_op_code {
 
 enum htp_tensor_flags {
     HTP_TENSOR_WEIGHT  = (1U << 0), // Tensor buffer model weight data (not compute)
-    HTP_TENSOR_REPACK  = (1U << 1), // Tensor is in repacked tiled format
+    HTP_TENSOR_TILED   = (1U << 1), // Tensor is in tiled format
     HTP_TENSOR_FENCE   = (1U << 2)  // Tensor is synchronization fence (explicitly managed)
 };
 

ggml/src/ggml-hexagon/htp/htp-tensor.h

@@ -8,27 +8,59 @@
 #include "hex-common.h"
 #include "hex-fastdiv.h"
 
-enum {
-    HTP_TENSOR_MDEV_LINE_SIZE = 128,
-};
+#define HTP_TENSOR_MDEV_LINE_SIZE   128
+
+// Currently HMX requires 32x32 tile formats (do not change)
+#define HTP_TENSOR_TILE_N_COLS      32
+#define HTP_TENSOR_TILE_N_ROWS      32
+
+#define HTP_TENSOR_TILE_SIZE_F16    2048
+#define HTP_TENSOR_TILE_SIZE_Q4_0   576
+#define HTP_TENSOR_TILE_SIZE_Q4_1   640
+#define HTP_TENSOR_TILE_SIZE_Q8_0   1088
+#define HTP_TENSOR_TILE_SIZE_IQ4_NL 576
+#define HTP_TENSOR_TILE_SIZE_MXFP4  544
+#define HTP_TENSOR_TILE_SIZE_Q4_K   640
+#define HTP_TENSOR_TILE_SIZE_Q5_K   768
+#define HTP_TENSOR_TILE_SIZE_Q6_K   896
+#define HTP_TENSOR_TILE_SIZE_Q3_K   512
+#define HTP_TENSOR_TILE_SIZE_Q2_K   336
+
+// Q2_K, Q3_K, Q6_K use vrmpy byte order within each 128B HVX vector:
+//   k = 0..31 is split into 8 groups g = k/4 of 4 k each; byte 4*r + (k%4) holds row r
+//
+// Q2_K (336B, lossless tiling with same size as block_q2_K):
+//   bytes   0..255 (vec 0..1): unsigned 2-bit quants, vec m holds groups 4m..4m+3 at bit offsets 0,2,4,6
+//   bytes 256..319:            raw scales[] bytes (min << 4 | sc): [0..31] for k 0..15, [32..63] for k 16..31
+//   bytes 320..335:            8 fp16 in tile j (0..7) of super-block: d (j < 4) or -dmin (j >= 4) for rows 8*(j%4)..8*(j%4)+7
+//
+// Q3_K (512B):
+//   bytes   0..255 (vec 0..1): low 2 bits, vec m holds groups 4m..4m+3 at bit offsets 0,2,4,6
+//   bytes 256..383 (vec 2):    sign mask, bit g set when hmask is 0 (quant = low2 - 4)
+//   bytes 384..511 (vec 3):    fp16 scales d * (scales[] - 32): fp16[0..31] for k 0..15, fp16[32..63] for k 16..31
+//
+// Q4_K (640B, same layout as Q4_1 tile; each 256-k super-block splits into 8 tiles of 32 k):
+//   bytes   0..511 (vec 0..3): 4-bit quants, byte 32*cp + r holds k = 2*cp (low nibble) and 2*cp + 1 (high nibble)
+//   bytes 512..639 (vec 4):    32x2 fp16 per row: (d * sc, -dmin * m)
+//
+// Q5_K (768B):
+//   bytes   0..639 (vec 0..4): Q4_K tile (low 4 bits + fp16 scale/min pairs)
+//   bytes 640..767 (vec 5):    5th-bit plane; byte l holds the 8 high bits for lane l across vec i = 0..3 (bit 2i = lo, 2i+1 = hi)
+//
+// Q6_K (896B):
+//   bytes   0..511 (vec 0..3): low 4 bits, vec i holds group 2i (low nibble) and 2i+1 (high nibble)
+//   bytes 512..767 (vec 4..5): high 2 bits, vec 4+m holds groups 4m..4m+3 at bit offsets 0,2,4,6
+//   bytes 768..895 (vec 6):    fp16 scales d * scales[]: fp16[0..31] for k 0..15, fp16[32..63] for k 16..31
 
 struct htp_tensor_mdev_range {
     uint32_t start;
     uint32_t count;
 };
 
-static inline void * htp_tensor_data(const struct htp_tensor * t) {
-    return (void *) (uintptr_t) t->data;
-}
-
 static inline bool htp_tensor_is_extended(const struct htp_tensor * t) {
     return t && (t->data >> 32) != 0;
 }
 
-static inline uint32_t * htp_tensor_flags(const struct htp_tensor * t) {
-    return (uint32_t *) &t->flags;
-}
-
 static inline bool htp_tensor_is_contiguous(const struct htp_tensor * t, uint32_t type_size) {
     uint32_t next_nb = type_size;
     if (t->ne[0] != 1 && t->nb[0] != next_nb) {
@@ -125,19 +157,50 @@ static inline struct htp_tensor_mdev_range htp_tensor_mdev_partition(uint32_t to
     return range;
 }
 
-static inline uint32_t htp_tensor_get_row_size(int type, uint32_t ne00) {
+static inline uint32_t htp_tensor_get_tile_size(int type) {
     switch (type) {
-        case HTP_TYPE_F32:  return ne00 * 4;
-        case HTP_TYPE_F16:  return ne00 * 2;
-        case HTP_TYPE_Q8_0: return (ne00 / 32) * 34;
-        case HTP_TYPE_I32:  return ne00 * 4;
-        default:            return 0;
+        case HTP_TYPE_Q4_0:   return HTP_TENSOR_TILE_SIZE_Q4_0;
+        case HTP_TYPE_IQ4_NL: return HTP_TENSOR_TILE_SIZE_IQ4_NL;
+        case HTP_TYPE_Q4_1:   return HTP_TENSOR_TILE_SIZE_Q4_1;
+        case HTP_TYPE_Q4_K:   return HTP_TENSOR_TILE_SIZE_Q4_K;
+        case HTP_TYPE_Q8_0:   return HTP_TENSOR_TILE_SIZE_Q8_0;
+        case HTP_TYPE_Q5_K:   return HTP_TENSOR_TILE_SIZE_Q5_K;
+        case HTP_TYPE_Q6_K:   return HTP_TENSOR_TILE_SIZE_Q6_K;
+        case HTP_TYPE_Q3_K:   return HTP_TENSOR_TILE_SIZE_Q3_K;
+        case HTP_TYPE_Q2_K:   return HTP_TENSOR_TILE_SIZE_Q2_K;
+        case HTP_TYPE_MXFP4:  return HTP_TENSOR_TILE_SIZE_MXFP4;
+        case HTP_TYPE_F16:    return HTP_TENSOR_TILE_SIZE_F16;
+        default:              return 0;
     }
 }
 
-struct htp_context;
-void htp_flush_dirty_ranges(struct htp_context * ctx);
-void htp_tensor_flush_all(struct htp_context * ctx, const struct htp_tensor * const * tensors, uint32_t n);
-void htp_tensor_dirty_all(struct htp_context * ctx, const struct htp_tensor * const * tensors, uint32_t n);
+static inline uint32_t htp_tensor_get_tile_padded_size(int type) {
+    return (htp_tensor_get_tile_size(type) + 127u) & ~127u;
+}
+
+static inline bool htp_tensor_is_tiled(const struct htp_tensor * t) {
+    return t && (t->flags & HTP_TENSOR_TILED) != 0;
+}
+
+static inline uint32_t htp_tensor_get_row_size(int type, uint32_t ne00) {
+    switch (type) {
+        case HTP_TYPE_F32:    return ne00 * 4;
+        case HTP_TYPE_F16:
+        case HTP_TYPE_BF16:   return ne00 * 2;
+        case HTP_TYPE_Q8_0:   return (ne00 / 32) * 34;
+        case HTP_TYPE_Q4_0:
+        case HTP_TYPE_IQ4_NL: return (ne00 / 32) * 18;
+        case HTP_TYPE_Q4_1:   return (ne00 / 32) * 20;
+        case HTP_TYPE_MXFP4:  return (ne00 / 32) * 17;
+        case HTP_TYPE_Q2_K:   return (ne00 / 256) * 84;
+        case HTP_TYPE_Q3_K:   return (ne00 / 256) * 110;
+        case HTP_TYPE_Q4_K:   return (ne00 / 256) * 144;
+        case HTP_TYPE_Q5_K:   return (ne00 / 256) * 176;
+        case HTP_TYPE_Q6_K:   return (ne00 / 256) * 210;
+        case HTP_TYPE_I32:    return ne00 * 4;
+        case HTP_TYPE_I64:    return ne00 * 8;
+        default:              return 0;
+    }
+}
 
 #endif // HTP_TENSOR_H

ggml/src/ggml-hexagon/htp/hvx-mm-kernels-float.h

@@ -76,7 +76,7 @@ static inline void quantize_f16_f16_kernel(
 #define HVX_OP_MUL_F32(a, b) Q6_Vsf_vmpy_VsfVsf(a, b)
 #endif
 
-static inline void vec_dot_f32_f32_aa_1x1(const uint32_t n, float * restrict s, const void * restrict vx, const void * restrict vy) {
+static inline void flat_vec_dot_f32_f32_aa_1x1(const uint32_t n, float * restrict s, const void * restrict vx, const void * restrict vy) {
     const HVX_Vector * restrict x = (const HVX_Vector *) vx;
     const HVX_Vector * restrict y = (const HVX_Vector *) vy;
 
@@ -104,7 +104,7 @@ static inline void vec_dot_f32_f32_aa_1x1(const uint32_t n, float * restrict s,
     *s = hvx_vec_get_f32(hvx_vec_reduce_sum_f32(rsum));
 }
 
-static inline void vec_dot_f32_f32_aa_2x1(const uint32_t n, float * restrict s0,
+static inline void flat_vec_dot_f32_f32_aa_2x1(const uint32_t n, float * restrict s0,
                                 const void * restrict vx0, const void * restrict vx1,
                                 const void * restrict vy0) {
     const HVX_Vector * restrict x0 = (const HVX_Vector *) vx0;
@@ -143,7 +143,7 @@ static inline void vec_dot_f32_f32_aa_2x1(const uint32_t n, float * restrict s0,
     hvx_vec_store_u(s0, 8, rsum);
 }
 
-static inline void vec_dot_f32_f32_aa_2x2(const uint32_t n, float * restrict s0, float * restrict s1,
+static inline void flat_vec_dot_f32_f32_aa_2x2(const uint32_t n, float * restrict s0, float * restrict s1,
                                 const void * restrict vx0, const void * restrict vx1,
                                 const void * restrict vy0, const void * restrict vy1) {
     const HVX_Vector * restrict x0 = (const HVX_Vector *) vx0;
@@ -196,10 +196,203 @@ static inline void vec_dot_f32_f32_aa_2x2(const uint32_t n, float * restrict s0,
     hvx_vec_store_u(s1, 8, r0_r1_c1_sum);
 }
 
+static inline void flat_vec_dot_bf16_f32_aa_1x1(const uint32_t n, float * restrict s, const void * restrict vx, const void * restrict vy) {
+    const HVX_Vector * restrict x = (const HVX_Vector *) vx;
+    const HVX_Vector * restrict y = (const HVX_Vector *) vy;
+
+    uint32_t nvec = n / 64;
+    uint32_t rem  = n % 64;
+
+    HVX_Vector rsum0 = Q6_V_vzero();
+    HVX_Vector rsum1 = Q6_V_vzero();
+
+    uint32_t i = 0;
+    #pragma unroll(2)
+    for (i = 0; i < nvec; i++) {
+        HVX_Vector v_bf16 = x[i];
+        HVX_VectorPair p = Q6_Ww_vunpack_Vh(v_bf16);
+        HVX_Vector w0_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p), 16);
+        HVX_Vector w1_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p), 16);
+
+        rsum0 = HVX_OP_ADD_F32(rsum0, HVX_OP_MUL_F32(w0_f32, y[2 * i + 0]));
+        rsum1 = HVX_OP_ADD_F32(rsum1, HVX_OP_MUL_F32(w1_f32, y[2 * i + 1]));
+    }
+
+    if (rem) {
+        HVX_VectorPred bmask = Q6_Q_vsetq2_R(rem * 2);
+        HVX_Vector v_bf16 = Q6_V_vand_QV(bmask, x[i]);
+        HVX_VectorPair p = Q6_Ww_vunpack_Vh(v_bf16);
+        HVX_Vector w0_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p), 16);
+        HVX_Vector w1_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p), 16);
+
+        if (rem <= 32) {
+            HVX_VectorPred a_mask = Q6_Q_vsetq2_R(rem * 4);
+            HVX_Vector y0 = Q6_V_vand_QV(a_mask, y[2 * i + 0]);
+            rsum0 = HVX_OP_ADD_F32(rsum0, HVX_OP_MUL_F32(w0_f32, y0));
+        } else {
+            HVX_Vector y0 = y[2 * i + 0];
+            rsum0 = HVX_OP_ADD_F32(rsum0, HVX_OP_MUL_F32(w0_f32, y0));
+
+            HVX_VectorPred a_mask = Q6_Q_vsetq2_R((rem - 32) * 4);
+            HVX_Vector y1 = Q6_V_vand_QV(a_mask, y[2 * i + 1]);
+            rsum1 = HVX_OP_ADD_F32(rsum1, HVX_OP_MUL_F32(w1_f32, y1));
+        }
+    }
+
+    HVX_Vector rsum = HVX_OP_ADD_F32(rsum0, rsum1);
+    *s = hvx_vec_get_f32(hvx_vec_reduce_sum_f32(rsum));
+}
+
+static inline void flat_vec_dot_bf16_f32_aa_2x1(const uint32_t n, float * restrict s0,
+                                const void * restrict vx0, const void * restrict vx1,
+                                const void * restrict vy0) {
+    const HVX_Vector * restrict x0 = (const HVX_Vector *) vx0;
+    const HVX_Vector * restrict x1 = (const HVX_Vector *) vx1;
+    const HVX_Vector * restrict y  = (const HVX_Vector *) vy0;
+
+    uint32_t nvec = n / 64;
+    uint32_t rem  = n % 64;
+
+    HVX_Vector rsum0 = Q6_V_vzero();
+    HVX_Vector rsum1 = Q6_V_vzero();
+
+    uint32_t i = 0;
+    #pragma unroll(2)
+    for (i = 0; i < nvec; i++) {
+        HVX_VectorPair p0 = Q6_Ww_vunpack_Vh(x0[i]);
+        HVX_VectorPair p1 = Q6_Ww_vunpack_Vh(x1[i]);
+        HVX_Vector w00_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p0), 16);
+        HVX_Vector w01_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p0), 16);
+        HVX_Vector w10_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p1), 16);
+        HVX_Vector w11_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p1), 16);
+
+        HVX_Vector y0 = y[2 * i + 0];
+        HVX_Vector y1 = y[2 * i + 1];
+
+        rsum0 = HVX_OP_ADD_F32(rsum0, HVX_OP_MUL_F32(w00_f32, y0));
+        rsum1 = HVX_OP_ADD_F32(rsum1, HVX_OP_MUL_F32(w10_f32, y0));
+        rsum0 = HVX_OP_ADD_F32(rsum0, HVX_OP_MUL_F32(w01_f32, y1));
+        rsum1 = HVX_OP_ADD_F32(rsum1, HVX_OP_MUL_F32(w11_f32, y1));
+    }
+
+    if (rem) {
+        HVX_VectorPred bmask = Q6_Q_vsetq2_R(rem * 2);
+        HVX_VectorPair p0 = Q6_Ww_vunpack_Vh(Q6_V_vand_QV(bmask, x0[i]));
+        HVX_VectorPair p1 = Q6_Ww_vunpack_Vh(Q6_V_vand_QV(bmask, x1[i]));
+        HVX_Vector w00_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p0), 16);
+        HVX_Vector w01_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p0), 16);
+        HVX_Vector w10_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p1), 16);
+        HVX_Vector w11_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p1), 16);
+
+        if (rem <= 32) {
+            HVX_VectorPred a_mask = Q6_Q_vsetq2_R(rem * 4);
+            HVX_Vector y0 = Q6_V_vand_QV(a_mask, y[2 * i + 0]);
+            rsum0 = HVX_OP_ADD_F32(rsum0, HVX_OP_MUL_F32(w00_f32, y0));
+            rsum1 = HVX_OP_ADD_F32(rsum1, HVX_OP_MUL_F32(w10_f32, y0));
+        } else {
+            HVX_Vector y0 = y[2 * i + 0];
+            rsum0 = HVX_OP_ADD_F32(rsum0, HVX_OP_MUL_F32(w00_f32, y0));
+            rsum1 = HVX_OP_ADD_F32(rsum1, HVX_OP_MUL_F32(w10_f32, y0));
+
+            HVX_VectorPred a_mask = Q6_Q_vsetq2_R((rem - 32) * 4);
+            HVX_Vector y1 = Q6_V_vand_QV(a_mask, y[2 * i + 1]);
+            rsum0 = HVX_OP_ADD_F32(rsum0, HVX_OP_MUL_F32(w01_f32, y1));
+            rsum1 = HVX_OP_ADD_F32(rsum1, HVX_OP_MUL_F32(w11_f32, y1));
+        }
+    }
+
+    HVX_Vector rsum = hvx_vec_reduce_sum_f32x2(rsum0, rsum1);
+    hvx_vec_store_u(s0, 8, rsum);
+}
+
+static inline void flat_vec_dot_bf16_f32_aa_2x2(const uint32_t n, float * restrict s0, float * restrict s1,
+                                const void * restrict vx0, const void * restrict vx1,
+                                const void * restrict vy0, const void * restrict vy1) {
+    const HVX_Vector * restrict x0 = (const HVX_Vector *) vx0;
+    const HVX_Vector * restrict x1 = (const HVX_Vector *) vx1;
+    const HVX_Vector * restrict y0 = (const HVX_Vector *) vy0;
+    const HVX_Vector * restrict y1 = (const HVX_Vector *) vy1;
+
+    uint32_t nvec = n / 64;
+    uint32_t rem  = n % 64;
+
+    HVX_Vector r0_c0_sum = Q6_V_vzero();
+    HVX_Vector r0_c1_sum = Q6_V_vzero();
+    HVX_Vector r1_c0_sum = Q6_V_vzero();
+    HVX_Vector r1_c1_sum = Q6_V_vzero();
+
+    uint32_t i = 0;
+    #pragma unroll(2)
+    for (i = 0; i < nvec; i++) {
+        HVX_VectorPair p0 = Q6_Ww_vunpack_Vh(x0[i]);
+        HVX_VectorPair p1 = Q6_Ww_vunpack_Vh(x1[i]);
+        HVX_Vector w00_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p0), 16);
+        HVX_Vector w01_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p0), 16);
+        HVX_Vector w10_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p1), 16);
+        HVX_Vector w11_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p1), 16);
+
+        HVX_Vector c0_0 = y0[2 * i + 0];
+        HVX_Vector c0_1 = y0[2 * i + 1];
+        HVX_Vector c1_0 = y1[2 * i + 0];
+        HVX_Vector c1_1 = y1[2 * i + 1];
+
+        r0_c0_sum = HVX_OP_ADD_F32(r0_c0_sum, HVX_OP_MUL_F32(w00_f32, c0_0));
+        r0_c1_sum = HVX_OP_ADD_F32(r0_c1_sum, HVX_OP_MUL_F32(w00_f32, c1_0));
+        r1_c0_sum = HVX_OP_ADD_F32(r1_c0_sum, HVX_OP_MUL_F32(w10_f32, c0_0));
+        r1_c1_sum = HVX_OP_ADD_F32(r1_c1_sum, HVX_OP_MUL_F32(w10_f32, c1_0));
+
+        r0_c0_sum = HVX_OP_ADD_F32(r0_c0_sum, HVX_OP_MUL_F32(w01_f32, c0_1));
+        r0_c1_sum = HVX_OP_ADD_F32(r0_c1_sum, HVX_OP_MUL_F32(w01_f32, c1_1));
+        r1_c0_sum = HVX_OP_ADD_F32(r1_c0_sum, HVX_OP_MUL_F32(w11_f32, c0_1));
+        r1_c1_sum = HVX_OP_ADD_F32(r1_c1_sum, HVX_OP_MUL_F32(w11_f32, c1_1));
+    }
+
+    if (rem) {
+        HVX_VectorPred bmask = Q6_Q_vsetq2_R(rem * 2);
+        HVX_VectorPair p0 = Q6_Ww_vunpack_Vh(Q6_V_vand_QV(bmask, x0[i]));
+        HVX_VectorPair p1 = Q6_Ww_vunpack_Vh(Q6_V_vand_QV(bmask, x1[i]));
+        HVX_Vector w00_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p0), 16);
+        HVX_Vector w01_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p0), 16);
+        HVX_Vector w10_f32 = Q6_Vw_vasl_VwR(Q6_V_lo_W(p1), 16);
+        HVX_Vector w11_f32 = Q6_Vw_vasl_VwR(Q6_V_hi_W(p1), 16);
+
+        if (rem <= 32) {
+            HVX_VectorPred a_mask = Q6_Q_vsetq2_R(rem * 4);
+            HVX_Vector c0_0 = Q6_V_vand_QV(a_mask, y0[2 * i + 0]);
+            HVX_Vector c1_0 = Q6_V_vand_QV(a_mask, y1[2 * i + 0]);
+            r0_c0_sum = HVX_OP_ADD_F32(r0_c0_sum, HVX_OP_MUL_F32(w00_f32, c0_0));
+            r0_c1_sum = HVX_OP_ADD_F32(r0_c1_sum, HVX_OP_MUL_F32(w00_f32, c1_0));
+            r1_c0_sum = HVX_OP_ADD_F32(r1_c0_sum, HVX_OP_MUL_F32(w10_f32, c0_0));
+            r1_c1_sum = HVX_OP_ADD_F32(r1_c1_sum, HVX_OP_MUL_F32(w10_f32, c1_0));
+        } else {
+            HVX_Vector c0_0 = y0[2 * i + 0];
+            HVX_Vector c1_0 = y1[2 * i + 0];
+            r0_c0_sum = HVX_OP_ADD_F32(r0_c0_sum, HVX_OP_MUL_F32(w00_f32, c0_0));
+            r0_c1_sum = HVX_OP_ADD_F32(r0_c1_sum, HVX_OP_MUL_F32(w00_f32, c1_0));
+            r1_c0_sum = HVX_OP_ADD_F32(r1_c0_sum, HVX_OP_MUL_F32(w10_f32, c0_0));
+            r1_c1_sum = HVX_OP_ADD_F32(r1_c1_sum, HVX_OP_MUL_F32(w10_f32, c1_0));
+
+            HVX_VectorPred a_mask = Q6_Q_vsetq2_R((rem - 32) * 4);
+            HVX_Vector c0_1 = Q6_V_vand_QV(a_mask, y0[2 * i + 1]);
+            HVX_Vector c1_1 = Q6_V_vand_QV(a_mask, y1[2 * i + 1]);
+            r0_c0_sum = HVX_OP_ADD_F32(r0_c0_sum, HVX_OP_MUL_F32(w01_f32, c0_1));
+            r0_c1_sum = HVX_OP_ADD_F32(r0_c1_sum, HVX_OP_MUL_F32(w01_f32, c1_1));
+            r1_c0_sum = HVX_OP_ADD_F32(r1_c0_sum, HVX_OP_MUL_F32(w11_f32, c0_1));
+            r1_c1_sum = HVX_OP_ADD_F32(r1_c1_sum, HVX_OP_MUL_F32(w11_f32, c1_1));
+        }
+    }
+
+    HVX_Vector r0_r1_c0_sum = hvx_vec_reduce_sum_f32x2(r0_c0_sum, r1_c0_sum);
+    HVX_Vector r0_r1_c1_sum = hvx_vec_reduce_sum_f32x2(r0_c1_sum, r1_c1_sum);
+
+    hvx_vec_store_u(s0, 8, r0_r1_c0_sum);
+    hvx_vec_store_u(s1, 8, r0_r1_c1_sum);
+}
+
 #undef HVX_OP_ADD_F32
 #undef HVX_OP_MUL_F32
 
-static inline void vec_dot_f16_f16_aa_1x1(const uint32_t n, float * restrict s, const void * restrict vx, const void * restrict vy) {
+static inline void flat_vec_dot_f16_f16_aa_1x1(const uint32_t n, float * restrict s, const void * restrict vx, const void * restrict vy) {
     const HVX_Vector * restrict x = (const HVX_Vector *) vx;
     const HVX_Vector * restrict y = (const HVX_Vector *) vy;
 
@@ -226,7 +419,7 @@ static inline void vec_dot_f16_f16_aa_1x1(const uint32_t n, float * restrict s,
     hvx_vec_store_u(s, 4, hvx_vec_reduce_sum_f32(rsum));
 }
 
-static inline void vec_dot_f16_f16_aa_2x1(const uint32_t n, float * restrict s0,
+static inline void flat_vec_dot_f16_f16_aa_2x1(const uint32_t n, float * restrict s0,
                                 const void * restrict vx0, const void * restrict vx1,
                                 const void * restrict vy0) {
     const HVX_Vector * restrict x0 = (const HVX_Vector *) vx0;
@@ -263,7 +456,7 @@ static inline void vec_dot_f16_f16_aa_2x1(const uint32_t n, float * restrict s0,
     hvx_vec_store_u(s0, 8, rsum);
 }
 
-static inline void vec_dot_f16_f16_aa_2x2(const uint32_t n, float * restrict s0, float * restrict s1,
+static inline void flat_vec_dot_f16_f16_aa_2x2(const uint32_t n, float * restrict s0, float * restrict s1,
                                 const void * restrict vx0, const void * restrict vx1,
                                 const void * restrict vy0, const void * restrict vy1) {
     const HVX_Vector * restrict x0 = (const HVX_Vector *) vx0;

ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h

@@ -274,6 +274,72 @@ static void quantize_row_f32_q8_1_s16_tiled(float * restrict x, uint8_t * restri
     }
 }
 
+static inline void quantize_half_f16_tiled_splat(HVX_Vector v_src, HVX_Vector * restrict dst) {
+    static const uint8_t __attribute__((aligned(128))) repl[128] = {
+        0x00, 0x00, 0x00, 0x00, 0x04, 0x04, 0x04, 0x04, 0x08, 0x08, 0x08, 0x08, 0x04, 0x04, 0x04, 0x04,
+        0x10, 0x10, 0x10, 0x10, 0x04, 0x04, 0x04, 0x04, 0x08, 0x08, 0x08, 0x08, 0x04, 0x04, 0x04, 0x04,
+        0x20, 0x20, 0x20, 0x20, 0x04, 0x04, 0x04, 0x04, 0x08, 0x08, 0x08, 0x08, 0x04, 0x04, 0x04, 0x04,
+        0x10, 0x10, 0x10, 0x10, 0x04, 0x04, 0x04, 0x04, 0x08, 0x08, 0x08, 0x08, 0x04, 0x04, 0x04, 0x04,
+        0x40, 0x40, 0x40, 0x40, 0x04, 0x04, 0x04, 0x04, 0x08, 0x08, 0x08, 0x08, 0x04, 0x04, 0x04, 0x04,
+        0x10, 0x10, 0x10, 0x10, 0x04, 0x04, 0x04, 0x04, 0x08, 0x08, 0x08, 0x08, 0x04, 0x04, 0x04, 0x04,
+        0x20, 0x20, 0x20, 0x20, 0x04, 0x04, 0x04, 0x04, 0x08, 0x08, 0x08, 0x08, 0x04, 0x04, 0x04, 0x04,
+        0x10, 0x10, 0x10, 0x10, 0x04, 0x04, 0x04, 0x04, 0x08, 0x08, 0x08, 0x08, 0x04, 0x04, 0x04, 0x04,
+    };
+    HVX_Vector v_repl_ctrl = *(const HVX_Vector *) repl;
+
+    #pragma unroll(4)
+    for (int cp = 0; cp < 32; cp++) {
+        dst[cp] = Q6_V_vdelta_VV(cp == 0 ? v_src : Q6_V_vror_VR(v_src, cp * 4), v_repl_ctrl);
+    }
+}
+
+static inline void quantize_block_f32_f16_tiled(const float * restrict x, uint8_t * restrict y_block) {
+    assert((unsigned long) x % 128 == 0);
+    assert((unsigned long) y_block % 128 == 0);
+
+    const HVX_Vector * vx = (const HVX_Vector *) x;
+    HVX_Vector * restrict dst = (HVX_Vector *) y_block;
+
+    HVX_Vector vx01_hf = hvx_vec_f32_to_f16(vx[0], vx[1]);
+    quantize_half_f16_tiled_splat(vx01_hf, dst);
+
+    HVX_Vector vx23_hf = hvx_vec_f32_to_f16(vx[2], vx[3]);
+    quantize_half_f16_tiled_splat(vx23_hf, dst + 32);
+}
+
+static inline void quantize_block_f16_f16_tiled(const __fp16 * restrict x, uint8_t * restrict y_block) {
+    assert((unsigned long) x % 128 == 0);
+    assert((unsigned long) y_block % 128 == 0);
+
+    const HVX_Vector * vx = (const HVX_Vector *) x;
+    HVX_Vector * restrict dst = (HVX_Vector *) y_block;
+
+    quantize_half_f16_tiled_splat(vx[0], dst);
+    quantize_half_f16_tiled_splat(vx[1], dst + 32);
+}
+
+static void quantize_row_f32_f16_tiled(const float * restrict x, uint8_t * restrict y, uint32_t k) {
+    assert(k % 32 == 0);
+    const uint32_t qk = QK_Q8_0_TILED;
+    const uint32_t nb = (k + qk - 1) / qk;
+
+    for (uint32_t i = 0; i < nb; i++) {
+        uint8_t * restrict y_block = y + i * 4 * 2048;
+        quantize_block_f32_f16_tiled(x + i * qk, y_block);
+    }
+}
+
+static void quantize_row_f16_f16_tiled(const __fp16 * restrict x, uint8_t * restrict y, uint32_t k) {
+    assert(k % 32 == 0);
+    const uint32_t qk = QK_Q8_0_TILED;
+    const uint32_t nb = (k + qk - 1) / qk;
+
+    for (uint32_t i = 0; i < nb; i++) {
+        uint8_t * restrict y_block = y + i * 4 * 2048;
+        quantize_block_f16_f16_tiled(x + i * qk, y_block);
+    }
+}
+
 // Dot kernels & helpers that consume tiled activations
 
 static inline HVX_Vector hvx_vec_mul_f16_f16_to_f32_lower32(HVX_Vector v1, HVX_Vector v2) {
@@ -454,7 +520,7 @@ static inline HVX_VectorPair accum_q8_0_32x2(
     return Q6_W_vcombine_VV(v_sum1, v_sum0);
 }
 
-// Q5_K: OR 0x10 into every lane of v whose flag j is set in the plane (see HTP_MM_WEIGHT_TILE_SIZE_Q5_K)
+// Q5_K: OR 0x10 into every lane of v whose flag j is set in the plane (see HTP_TENSOR_TILE_SIZE_Q5_K)
 static inline HVX_Vector hvx_q5k_or_hibit(HVX_Vector v, HVX_Vector v_plane, int j) {
     HVX_VectorPred q = Q6_Q_vand_VR(v_plane, 0x01010101u << j);
     return Q6_V_vandor_VQR(v, q, 0x10101010);
@@ -526,7 +592,7 @@ static inline HVX_VectorPair accum_5bit_32x2(
     return Q6_W_vcombine_VV(v_sum1, v_sum0);
 }
 
-// Q6_K weights are stored unsigned (0..63), see HTP_MM_WEIGHT_TILE_SIZE_Q6_K. Unpack k-group g of a tile to signed bytes (q - 32)
+// Q6_K weights are stored unsigned (0..63), see HTP_TENSOR_TILE_SIZE_Q6_K. Unpack k-group g of a tile to signed bytes (q - 32)
 static inline HVX_Vector unpack_q6_k_group(const HVX_Vector * restrict vptr, int g, HVX_Vector mask_0f, HVX_Vector mask_03, HVX_Vector i32) {
     HVX_Vector v_lo = (g & 1) ? Q6_Vub_vlsr_VubR(vptr[g >> 1], 4) : Q6_V_vand_VV(vptr[g >> 1], mask_0f);
     HVX_Vector v_hi = (g & 3) ? Q6_Vub_vlsr_VubR(vptr[4 + (g >> 2)], 2 * (g & 3)) : vptr[4 + (g >> 2)];
@@ -593,7 +659,7 @@ static inline HVX_Vector scale_q6_k_32x1(HVX_VectorPair v_sums, HVX_Vector v_sca
     return hvx_vec_add_f32_f32(v_lo, v_hi);
 }
 
-// Q3_K / Q2_K: low 2 bits of k-group g, see HTP_MM_WEIGHT_TILE_SIZE_Q3_K
+// Q3_K / Q2_K: low 2 bits of k-group g, see HTP_TENSOR_TILE_SIZE_Q3_K
 static inline HVX_Vector unpack_q3_k_low2(const HVX_Vector * restrict vptr, int g, HVX_Vector mask_03) {
     HVX_Vector v = vptr[g >> 2];
     if ((g & 3) == 3) {
@@ -1482,6 +1548,59 @@ static void tiled_vec_dot_mxfp4_32x2(const uint32_t n, float * restrict s0, floa
     }
 }
 
+static void tiled_vec_dot_f16_32x1(const uint32_t n, float * restrict s, const void * restrict vx, const void * restrict vy, uint32_t valid_rows, const float * restrict sz) {
+    const HVX_Vector * restrict vptr  = (const HVX_Vector *) vx;
+    const HVX_Vector * restrict v_act = (const HVX_Vector *) vy;
+    const uint32_t n_vecs = ((n + 31) / 32) * 16;
+
+    HVX_VectorPair acc = Q6_W_vzero();
+
+    #pragma unroll(4)
+    for (uint32_t i = 0; i < n_vecs; i++) {
+        acc = hvx_vec_mpyacc_f32_f16(acc, vptr[i], v_act[i]);
+    }
+
+    HVX_Vector v_sum = hvx_vec_add_f32_f32(Q6_V_lo_W(acc), Q6_V_hi_W(acc));
+
+    if (sz) {
+        hvx_vec_store_u(s, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum, hvx_vmemu(sz)));
+    } else {
+        hvx_vec_store_u(s, valid_rows * sizeof(float), v_sum);
+    }
+}
+
+static void tiled_vec_dot_f16_32x2(const uint32_t n, float * restrict s0, float * restrict s1, const void * restrict vx, const void * restrict vy0, const void * restrict vy1, uint32_t valid_rows, const float * restrict sz0, const float * restrict sz1) {
+    const HVX_Vector * restrict vptr   = (const HVX_Vector *) vx;
+    const HVX_Vector * restrict v_act0 = (const HVX_Vector *) vy0;
+    const HVX_Vector * restrict v_act1 = (const HVX_Vector *) vy1;
+    const uint32_t n_vecs = ((n + 31) / 32) * 16;
+
+    HVX_VectorPair acc0 = Q6_W_vzero();
+    HVX_VectorPair acc1 = Q6_W_vzero();
+
+    #pragma unroll(4)
+    for (uint32_t i = 0; i < n_vecs; i++) {
+        HVX_Vector v_w = vptr[i];
+        acc0 = hvx_vec_mpyacc_f32_f16(acc0, v_w, v_act0[i]);
+        acc1 = hvx_vec_mpyacc_f32_f16(acc1, v_w, v_act1[i]);
+    }
+
+    HVX_Vector v_sum0 = hvx_vec_add_f32_f32(Q6_V_lo_W(acc0), Q6_V_hi_W(acc0));
+    HVX_Vector v_sum1 = hvx_vec_add_f32_f32(Q6_V_lo_W(acc1), Q6_V_hi_W(acc1));
+
+    if (sz0) {
+        hvx_vec_store_u(s0, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum0, hvx_vmemu(sz0)));
+    } else {
+        hvx_vec_store_u(s0, valid_rows * sizeof(float), v_sum0);
+    }
+    if (sz1) {
+        hvx_vec_store_u(s1, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum1, hvx_vmemu(sz1)));
+    } else {
+        hvx_vec_store_u(s1, valid_rows * sizeof(float), v_sum1);
+    }
+}
+
+
 static inline void quantize_f32_q8_0_tiled_kernel(
     const uint8_t * restrict src_data,
     uint8_t * restrict dst_data,
@@ -1533,6 +1652,40 @@ static inline void quantize_f32_q8_1_s16_tiled_kernel(
     }
 }
 
+static inline void quantize_f32_f16_tiled_kernel(
+    const uint8_t * restrict src_data,
+    uint8_t * restrict dst_data,
+    uint8_t * restrict tmp_data,
+    uint32_t ne0,
+    uint32_t nrows,
+    size_t src_row_size,
+    size_t dst_row_size
+) {
+    (void) tmp_data;
+    for (uint32_t i = 0; i < nrows; ++i) {
+        quantize_row_f32_f16_tiled((const float *) src_data, dst_data, ne0);
+        dst_data += dst_row_size;
+        src_data += src_row_size;
+    }
+}
+
+static inline void quantize_f16_f16_tiled_kernel(
+    const uint8_t * restrict src_data,
+    uint8_t * restrict dst_data,
+    uint8_t * restrict tmp_data,
+    uint32_t ne0,
+    uint32_t nrows,
+    size_t src_row_size,
+    size_t dst_row_size
+) {
+    (void) tmp_data;
+    for (uint32_t i = 0; i < nrows; ++i) {
+        quantize_row_f16_f16_tiled((const __fp16 *) src_data, dst_data, ne0);
+        dst_data += dst_row_size;
+        src_data += src_row_size;
+    }
+}
+
 static inline void quantize_f32_q8_0_tiled_block_kernel(
     const float * restrict src,
     uint8_t * restrict dst,
@@ -1622,3 +1775,63 @@ static inline void quantize_f32_q8_1_s16_tiled_block_kernel(
         }
     }
 }
+
+static inline void quantize_f32_f16_tiled_block_kernel(
+    const float * restrict src,
+    uint8_t * restrict dst,
+    uint8_t * restrict tmp_data,
+    uint32_t ne0,
+    uint32_t ib_first,
+    uint32_t ib_last,
+    size_t src_row_size,
+    size_t dst_row_size,
+    uint32_t r,
+    uint32_t c
+) {
+    (void) tmp_data;
+    const uint32_t qk = QK_Q8_0_TILED;
+    const uint32_t nb = (ne0 + qk - 1) / qk;
+
+    for (uint32_t ib = ib_first; ib < ib_last; ++ib) {
+        const float * restrict src_ptr = (const float *) ((const uint8_t *) src + r * src_row_size + c * qk * sizeof(float));
+        uint8_t * restrict dst_ptr = dst + r * dst_row_size + c * 4 * 2048;
+
+        quantize_block_f32_f16_tiled(src_ptr, dst_ptr);
+
+        c++;
+        if (c == nb) {
+            c = 0;
+            r++;
+        }
+    }
+}
+
+static inline void quantize_f16_f16_tiled_block_kernel(
+    const __fp16 * restrict src,
+    uint8_t * restrict dst,
+    uint8_t * restrict tmp_data,
+    uint32_t ne0,
+    uint32_t ib_first,
+    uint32_t ib_last,
+    size_t src_row_size,
+    size_t dst_row_size,
+    uint32_t r,
+    uint32_t c
+) {
+    (void) tmp_data;
+    const uint32_t qk = QK_Q8_0_TILED;
+    const uint32_t nb = (ne0 + qk - 1) / qk;
+
+    for (uint32_t ib = ib_first; ib < ib_last; ++ib) {
+        const __fp16 * restrict src_ptr = (const __fp16 *) ((const uint8_t *) src + r * src_row_size + c * qk * sizeof(__fp16));
+        uint8_t * restrict dst_ptr = dst + r * dst_row_size + c * 4 * 2048;
+
+        quantize_block_f16_f16_tiled(src_ptr, dst_ptr);
+
+        c++;
+        if (c == nb) {
+            c = 0;
+            r++;
+        }
+    }
+}

ggml/src/ggml-hexagon/htp/hvx-quant.h

@@ -5,6 +5,7 @@
 #include <stdint.h>
 #include <string.h>
 
+#include "htp-tensor.h"
 #include "hvx-arith.h"
 #include "hvx-base.h"
 #include "hvx-reduce.h"
@@ -96,7 +97,7 @@ static inline void hvx_quantize_row_q8_0_f32(void * restrict dst_ptr, const floa
     }
 }
 
-static inline void hvx_dequantize_row_q8_0_f32(float * restrict dst_ptr, const void * restrict src_ptr, int n) {
+static inline void hvx_dequantize_row_flat_q8_0_f32(float * restrict dst_ptr, const void * restrict src_ptr, uint32_t n) {
     const int nb = n / QK8_0;
     const block_q8_0 * src = (const block_q8_0 *) src_ptr;
 
@@ -120,7 +121,7 @@ static inline void hvx_dequantize_row_q8_0_f32(float * restrict dst_ptr, const v
     }
 }
 
-static inline void hvx_dequantize_row_q8_0_f16(__fp16 * restrict dst_ptr, const void * restrict src_ptr, int n) {
+static inline void hvx_dequantize_row_flat_q8_0_f16(__fp16 * restrict dst_ptr, const void * restrict src_ptr, uint32_t n) {
     const int nb = n / QK8_0;
     const block_q8_0 * src = (const block_q8_0 *) src_ptr;
 
@@ -146,20 +147,460 @@ static inline void hvx_dequantize_row_q8_0_f16(__fp16 * restrict dst_ptr, const
     }
 }
 
-static inline void hvx_dequantize_row_f16_f32(float * restrict dst_ptr, const void * restrict src_ptr, int n) {
-    const int nb = n / 32;
+static inline void hvx_dequantize_row_flat_f16_f32(float * restrict dst_ptr, const void * restrict src_ptr, uint32_t n) {
+    const uint32_t nvec = n / VLEN_FP32;
+    const uint32_t nloe = n % VLEN_FP32;
     const _Float16 * src = (const _Float16 *) src_ptr;
 
-    for (int i = 0; i < nb; i++) {
-        HVX_Vector v_f16 = *(const HVX_UVector *) (src + i * 32);
+    for (uint32_t i = 0; i < nvec; i++) {
+        HVX_Vector v_f16 = *(const HVX_UVector *) (src + i * VLEN_FP32);
         HVX_VectorPair vp_f32 = hvx_vec_f16_to_f32(v_f16);
         HVX_Vector res = Q6_V_lo_W(vp_f32);
 
-        float * block_dst = dst_ptr + i * 32;
-        hvx_vmem(block_dst) = res;
+        hvx_vmem(dst_ptr + i * VLEN_FP32) = res;
+    }
+
+    if (nloe) {
+        HVX_Vector v_f16 = *(const HVX_UVector *) (src + nvec * VLEN_FP32);
+        HVX_VectorPair vp_f32 = hvx_vec_f16_to_f32(v_f16);
+        HVX_Vector res = Q6_V_lo_W(vp_f32);
+
+        hvx_vec_store_a(dst_ptr + nvec * VLEN_FP32, nloe * sizeof(float), res);
     }
 }
 
+static inline void hvx_dequantize_row_flat_bf16_f32(float * restrict dst_ptr, const void * restrict src_ptr, uint32_t n) {
+    const uint32_t nvec = n / VLEN_FP32;
+    const uint32_t nloe = n % VLEN_FP32;
+    const uint16_t * src = (const uint16_t *) src_ptr;
 
+    for (uint32_t i = 0; i < nvec; i++) {
+        HVX_Vector v_bf16 = *(const HVX_UVector *) (src + i * VLEN_FP32);
+        HVX_VectorPair p = Q6_Ww_vunpack_Vh(v_bf16);
+        HVX_Vector res = Q6_Vw_vasl_VwR(Q6_V_lo_W(p), 16);
+
+        hvx_vmem(dst_ptr + i * VLEN_FP32) = res;
+    }
+
+    if (nloe) {
+        HVX_Vector v_bf16 = *(const HVX_UVector *) (src + nvec * VLEN_FP32);
+        HVX_VectorPair p = Q6_Ww_vunpack_Vh(v_bf16);
+        HVX_Vector res = Q6_Vw_vasl_VwR(Q6_V_lo_W(p), 16);
+
+        hvx_vec_store_a(dst_ptr + nvec * VLEN_FP32, nloe * sizeof(float), res);
+    }
+}
+
+static inline void hvx_dequantize_row_flat_q4_0_f32(float * restrict dst_ptr, const void * restrict src_ptr, uint32_t n) {
+    const int nb = n / QK4_0;
+    const block_q4_0 * src = (const block_q4_0 *) src_ptr;
+
+    const HVX_VectorPred q16 = Q6_Q_vsetq_R(16);
+    const HVX_Vector v_eight = Q6_Vb_vsplat_R(8);
+    const HVX_Vector v_zero  = Q6_V_vzero();
+
+    for (int i = 0; i < nb; i++) {
+        HVX_Vector vd_f16     = Q6_Vh_vsplat_R(*(const int16_t *) &src[i].d);
+        HVX_VectorPair vp_f32 = hvx_vec_f16_to_f32(vd_f16);
+        HVX_Vector vd         = Q6_V_lo_W(vp_f32);
+
+        HVX_Vector vq = *(const HVX_UVector *) src[i].qs;
+        HVX_Vector v_lo = Q6_V_vand_VV(vq, Q6_Vb_vsplat_R(0x0F));
+        HVX_Vector v_hi = Q6_Vub_vlsr_VubR(vq, 4);
+
+        HVX_Vector v_lo_s8 = Q6_Vb_vsub_VbVb(v_lo, v_eight);
+        HVX_Vector v_hi_s8 = Q6_Vb_vsub_VbVb(v_hi, v_eight);
+
+        HVX_Vector v32 = Q6_V_vmux_QVV(q16, v_lo_s8, Q6_V_vlalign_VVR(v_hi_s8, v_zero, 16));
+
+        HVX_VectorPair p16   = Q6_Wh_vunpack_Vb(v32);
+        HVX_Vector     v_i16 = Q6_V_lo_W(p16);
+        HVX_VectorPair p32   = Q6_Ww_vunpack_Vh(v_i16);
+        HVX_Vector     v_i32 = Q6_V_lo_W(p32);
+
+        HVX_Vector v_f32 = Q6_Vsf_equals_Vw(v_i32);
+        HVX_Vector res   = hvx_vec_mul_f32_f32(v_f32, vd);
+
+        hvx_vmem(dst_ptr + i * QK4_0) = res;
+    }
+}
+
+static inline void hvx_dequantize_row_flat_q4_k_f32(float * restrict dst_ptr, const void * restrict src_ptr, uint32_t n) {
+    const int nb = n / QK_K;
+    const block_q4_K * src = (const block_q4_K *) src_ptr;
+
+    for (int i = 0; i < nb; i++) {
+        const float d   = GGML_FP16_TO_FP32(src[i].d);
+        const float min = GGML_FP16_TO_FP32(src[i].dmin);
+        const uint8_t * sc = src[i].scales;
+
+        float * y = dst_ptr + i * QK_K;
+
+        for (int step = 0; step < 4; ++step) {
+            uint8_t sc1, m1_u, sc2, m2_u;
+            if (step < 2) {
+                sc1  = sc[2 * step + 0] & 63;
+                m1_u = sc[2 * step + 4] & 63;
+                sc2  = sc[2 * step + 1] & 63;
+                m2_u = sc[2 * step + 5] & 63;
+            } else {
+                const int s = step - 2;
+                sc1  = (sc[8 + 2 * s + 0] & 0xF) | ((sc[2 * s + 0] >> 6) << 4);
+                m1_u = (sc[8 + 2 * s + 0] >>  4) | ((sc[4 + 2 * s + 0] >> 6) << 4);
+                sc2  = (sc[8 + 2 * s + 1] & 0xF) | ((sc[2 * s + 1] >> 6) << 4);
+                m2_u = (sc[8 + 2 * s + 1] >>  4) | ((sc[4 + 2 * s + 1] >> 6) << 4);
+            }
+
+            HVX_Vector vq = *(const HVX_UVector *) (src[i].qs + step * 32);
+            HVX_Vector v_lo = Q6_V_vand_VV(vq, Q6_Vb_vsplat_R(0x0F));
+            HVX_Vector v_hi = Q6_Vub_vlsr_VubR(vq, 4);
+
+            HVX_VectorPair p16_lo = Q6_Wuh_vunpack_Vub(v_lo);
+            HVX_VectorPair p32_lo = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16_lo));
+            HVX_Vector     v_i32_lo = Q6_V_lo_W(p32_lo);
+            HVX_Vector     v_f32_lo = Q6_Vsf_equals_Vw(v_i32_lo);
+
+            HVX_Vector vd1 = hvx_vec_splat_f32(d * sc1);
+            HVX_Vector vm1 = hvx_vec_splat_f32(min * m1_u);
+            HVX_Vector v_out1 = hvx_vec_sub_f32_f32(hvx_vec_mul_f32_f32(v_f32_lo, vd1), vm1);
+            hvx_vmem(y + 0) = v_out1;
+
+            HVX_VectorPair p16_hi = Q6_Wuh_vunpack_Vub(v_hi);
+            HVX_VectorPair p32_hi = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16_hi));
+            HVX_Vector     v_i32_hi = Q6_V_lo_W(p32_hi);
+            HVX_Vector     v_f32_hi = Q6_Vsf_equals_Vw(v_i32_hi);
+
+            HVX_Vector vd2 = hvx_vec_splat_f32(d * sc2);
+            HVX_Vector vm2 = hvx_vec_splat_f32(min * m2_u);
+            HVX_Vector v_out2 = hvx_vec_sub_f32_f32(hvx_vec_mul_f32_f32(v_f32_hi, vd2), vm2);
+            hvx_vmem(y + 32) = v_out2;
+
+            y += 64;
+        }
+    }
+}
+
+static inline void hvx_dequantize_row_flat_q5_k_f32(float * restrict dst_ptr, const void * restrict src_ptr, uint32_t n) {
+    const int nb = n / QK_K;
+    const block_q5_K * src = (const block_q5_K *) src_ptr;
+
+    for (int i = 0; i < nb; i++) {
+        const float d   = GGML_FP16_TO_FP32(src[i].d);
+        const float min = GGML_FP16_TO_FP32(src[i].dmin);
+        const uint8_t * sc = src[i].scales;
+
+        HVX_Vector v_qh = *(const HVX_UVector *) src[i].qh;
+        float * y = dst_ptr + i * QK_K;
+
+        for (int step = 0; step < 4; ++step) {
+            uint8_t sc1, m1_u, sc2, m2_u;
+            if (step < 2) {
+                sc1  = sc[2 * step + 0] & 63;
+                m1_u = sc[2 * step + 4] & 63;
+                sc2  = sc[2 * step + 1] & 63;
+                m2_u = sc[2 * step + 5] & 63;
+            } else {
+                const int s = step - 2;
+                sc1  = (sc[8 + 2 * s + 0] & 0xF) | ((sc[2 * s + 0] >> 6) << 4);
+                m1_u = (sc[8 + 2 * s + 0] >>  4) | ((sc[4 + 2 * s + 0] >> 6) << 4);
+                sc2  = (sc[8 + 2 * s + 1] & 0xF) | ((sc[2 * s + 1] >> 6) << 4);
+                m2_u = (sc[8 + 2 * s + 1] >>  4) | ((sc[4 + 2 * s + 1] >> 6) << 4);
+            }
+
+            HVX_Vector v_ql = *(const HVX_UVector *) (src[i].qs + step * 32);
+            HVX_Vector v_lo = Q6_V_vand_VV(v_ql, Q6_Vb_vsplat_R(0x0F));
+            HVX_Vector v_hi = Q6_Vub_vlsr_VubR(v_ql, 4);
+
+            HVX_Vector v_h1 = Q6_Vw_vasl_VwR(Q6_V_vand_VV(Q6_Vub_vlsr_VubR(v_qh, 2 * step), Q6_Vb_vsplat_R(1)), 4);
+            HVX_Vector v_val1 = Q6_V_vor_VV(v_lo, v_h1);
+
+            HVX_VectorPair p16_1 = Q6_Wuh_vunpack_Vub(v_val1);
+            HVX_VectorPair p32_1 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16_1));
+            HVX_Vector     v_i32_1 = Q6_V_lo_W(p32_1);
+            HVX_Vector     v_f32_1 = Q6_Vsf_equals_Vw(v_i32_1);
+
+            HVX_Vector vd1 = hvx_vec_splat_f32(d * sc1);
+            HVX_Vector vm1 = hvx_vec_splat_f32(min * m1_u);
+            HVX_Vector v_out1 = hvx_vec_sub_f32_f32(hvx_vec_mul_f32_f32(v_f32_1, vd1), vm1);
+            hvx_vmem(y + 0) = v_out1;
+
+            HVX_Vector v_h2 = Q6_Vw_vasl_VwR(Q6_V_vand_VV(Q6_Vub_vlsr_VubR(v_qh, 2 * step + 1), Q6_Vb_vsplat_R(1)), 4);
+            HVX_Vector v_val2 = Q6_V_vor_VV(v_hi, v_h2);
+
+            HVX_VectorPair p16_2 = Q6_Wuh_vunpack_Vub(v_val2);
+            HVX_VectorPair p32_2 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16_2));
+            HVX_Vector     v_i32_2 = Q6_V_lo_W(p32_2);
+            HVX_Vector     v_f32_2 = Q6_Vsf_equals_Vw(v_i32_2);
+
+            HVX_Vector vd2 = hvx_vec_splat_f32(d * sc2);
+            HVX_Vector vm2 = hvx_vec_splat_f32(min * m2_u);
+            HVX_Vector v_out2 = hvx_vec_sub_f32_f32(hvx_vec_mul_f32_f32(v_f32_2, vd2), vm2);
+            hvx_vmem(y + 32) = v_out2;
+
+            y += 64;
+        }
+    }
+}
+
+static inline void hvx_dequantize_row_flat_q6_k_f32(float * restrict dst_ptr, const void * restrict src_ptr, uint32_t n) {
+    const int nb = n / QK_K;
+    const block_q6_K * src = (const block_q6_K *) src_ptr;
+
+    const HVX_VectorPred q64 = Q6_Q_vsetq_R(64);
+    const HVX_Vector v32_b   = Q6_Vb_vsplat_R(32);
+    const HVX_Vector mask_0f = Q6_Vb_vsplat_R(0x0F);
+    const HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03);
+
+    for (int i = 0; i < nb; i++) {
+        const float d = GGML_FP16_TO_FP32(src[i].d);
+
+        const uint8_t * restrict ql = src[i].ql;
+        const uint8_t * restrict qh = src[i].qh;
+        const int8_t  * restrict sc = src[i].scales;
+
+        float * y = dst_ptr + i * QK_K;
+
+        for (int half = 0; half < 2; ++half) {
+            float s[8];
+            for (int j = 0; j < 8; ++j) {
+                s[j] = d * sc[j];
+            }
+
+            HVX_Vector vscale1 = Q6_V_vmux_QVV(q64, hvx_vec_splat_f32(s[0]), hvx_vec_splat_f32(s[1]));
+            HVX_Vector vscale2 = Q6_V_vmux_QVV(q64, hvx_vec_splat_f32(s[2]), hvx_vec_splat_f32(s[3]));
+            HVX_Vector vscale3 = Q6_V_vmux_QVV(q64, hvx_vec_splat_f32(s[4]), hvx_vec_splat_f32(s[5]));
+            HVX_Vector vscale4 = Q6_V_vmux_QVV(q64, hvx_vec_splat_f32(s[6]), hvx_vec_splat_f32(s[7]));
+
+            HVX_Vector v_ql0 = *(const HVX_UVector *) (ql);
+            HVX_Vector v_ql1 = *(const HVX_UVector *) (ql + 32);
+            HVX_Vector v_qh  = *(const HVX_UVector *) (qh);
+
+            HVX_Vector ql0_lo = Q6_V_vand_VV(v_ql0, mask_0f);
+            HVX_Vector qh_0   = Q6_Vw_vasl_VwR(Q6_V_vand_VV(v_qh, mask_03), 4);
+            HVX_Vector v_q1   = Q6_Vb_vsub_VbVb(Q6_V_vor_VV(ql0_lo, qh_0), v32_b);
+
+            HVX_Vector ql1_lo = Q6_V_vand_VV(v_ql1, mask_0f);
+            HVX_Vector qh_1   = Q6_Vw_vasl_VwR(Q6_V_vand_VV(Q6_Vub_vlsr_VubR(v_qh, 2), mask_03), 4);
+            HVX_Vector v_q2   = Q6_Vb_vsub_VbVb(Q6_V_vor_VV(ql1_lo, qh_1), v32_b);
+
+            HVX_Vector ql0_hi = Q6_Vub_vlsr_VubR(v_ql0, 4);
+            HVX_Vector qh_2   = Q6_Vw_vasl_VwR(Q6_V_vand_VV(Q6_Vub_vlsr_VubR(v_qh, 4), mask_03), 4);
+            HVX_Vector v_q3   = Q6_Vb_vsub_VbVb(Q6_V_vor_VV(ql0_hi, qh_2), v32_b);
+
+            HVX_Vector ql1_hi = Q6_Vub_vlsr_VubR(v_ql1, 4);
+            HVX_Vector qh_3   = Q6_Vw_vasl_VwR(Q6_V_vand_VV(Q6_Vub_vlsr_VubR(v_qh, 6), mask_03), 4);
+            HVX_Vector v_q4   = Q6_Vb_vsub_VbVb(Q6_V_vor_VV(ql1_hi, qh_3), v32_b);
+
+            HVX_VectorPair p16_1 = Q6_Wh_vunpack_Vb(v_q1);
+            HVX_VectorPair p32_1 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16_1));
+            HVX_Vector     v_f32_1 = Q6_Vsf_equals_Vw(Q6_V_lo_W(p32_1));
+            hvx_vmem(y + 0) = hvx_vec_mul_f32_f32(v_f32_1, vscale1);
+
+            HVX_VectorPair p16_2 = Q6_Wh_vunpack_Vb(v_q2);
+            HVX_VectorPair p32_2 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16_2));
+            HVX_Vector     v_f32_2 = Q6_Vsf_equals_Vw(Q6_V_lo_W(p32_2));
+            hvx_vmem(y + 32) = hvx_vec_mul_f32_f32(v_f32_2, vscale2);
+
+            HVX_VectorPair p16_3 = Q6_Wh_vunpack_Vb(v_q3);
+            HVX_VectorPair p32_3 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16_3));
+            HVX_Vector     v_f32_3 = Q6_Vsf_equals_Vw(Q6_V_lo_W(p32_3));
+            hvx_vmem(y + 64) = hvx_vec_mul_f32_f32(v_f32_3, vscale3);
+
+            HVX_VectorPair p16_4 = Q6_Wh_vunpack_Vb(v_q4);
+            HVX_VectorPair p32_4 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16_4));
+            HVX_Vector     v_f32_4 = Q6_Vsf_equals_Vw(Q6_V_lo_W(p32_4));
+            hvx_vmem(y + 96) = hvx_vec_mul_f32_f32(v_f32_4, vscale4);
+
+            y  += 128;
+            ql += 64;
+            qh += 32;
+            sc += 8;
+        }
+    }
+}
+
+// bytes row, row + 32, row + 64 and row + 96 of v, to bytes 0..3
+static inline HVX_Vector hvx_dequantize_gather4(HVX_Vector v, uint32_t row) {
+    v = Q6_V_vror_VR(v, row);
+    HVX_Vector packed = Q6_V_vmux_QVV(Q6_Q_vsetq_R(1), v, Q6_V_vror_VR(v, 31));
+    packed = Q6_V_vmux_QVV(Q6_Q_vsetq_R(2), packed, Q6_V_vror_VR(v, 62));
+    return Q6_V_vmux_QVV(Q6_Q_vsetq_R(3), packed, Q6_V_vror_VR(v, 93));
+}
+
+static inline void hvx_dequantize_row_tiled_q8_0_f32(float * dst_spad, const uint8_t * src_spad, uint32_t n_k_tiles, uint32_t tile_stride, uint32_t row) {
+    const HVX_VectorPred first2 = Q6_Q_vsetq_R(2);
+    const HVX_VectorPred first4 = Q6_Q_vsetq_R(4);
+
+    for (uint32_t k_tile = 0; k_tile < n_k_tiles; ++k_tile) {
+        const uint8_t * tile = src_spad + k_tile * tile_stride;
+        float * dst = dst_spad + k_tile * HTP_TENSOR_TILE_N_COLS;
+
+        HVX_Vector vq = Q6_V_vzero();
+        for (int group = 7; group >= 0; --group) {
+            const HVX_Vector v = Q6_V_vror_VR(hvx_vmem(tile + group * VLEN), 2 * row);
+            const HVX_Vector packed = Q6_V_vmux_QVV(first2, v, Q6_V_vror_VR(v, 62));
+            vq = Q6_V_vmux_QVV(first4, packed, Q6_V_vror_VR(vq, VLEN - 4));
+        }
+
+        const HVX_Vector scales = hvx_vmem(tile + 1024);
+        const HVX_Vector scale_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 2 * row));
+        const HVX_Vector scale = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_hf));
+        const HVX_VectorPair p16 = Q6_Wh_vunpack_Vb(vq);
+        const HVX_VectorPair p32 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16));
+        HVX_Vector values = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(p32)), scale);
+        *(HVX_Vector *) dst = values;
+    }
+}
+
+static inline void hvx_dequantize_row_tiled_q4_0_f32(float * dst_spad, const uint8_t * src_spad, uint32_t n_k_tiles, uint32_t tile_stride, uint32_t row) {
+    const HVX_VectorPred first4 = Q6_Q_vsetq_R(4);
+
+    for (uint32_t k_tile = 0; k_tile < n_k_tiles; ++k_tile) {
+        const uint8_t * tile = src_spad + k_tile * tile_stride;
+        float * dst = dst_spad + k_tile * HTP_TENSOR_TILE_N_COLS;
+
+        HVX_Vector vq = Q6_V_vzero();
+        for (int group = 3; group >= 0; --group) {
+            vq = Q6_V_vmux_QVV(first4, hvx_dequantize_gather4(hvx_vmem(tile + group * VLEN), row), Q6_V_vror_VR(vq, VLEN - 4));
+        }
+
+        const HVX_Vector lo = Q6_V_vand_VV(vq, Q6_Vb_vsplat_R(0x0F));
+        const HVX_Vector hi = Q6_Vub_vlsr_VubR(vq, 4);
+        vq = Q6_V_lo_W(Q6_W_vshuff_VVR(hi, lo, -1));
+        vq = Q6_Vb_vsub_VbVb(vq, Q6_Vb_vsplat_R(8));
+
+        const HVX_Vector scales = hvx_vmem(tile + 512);
+        const HVX_Vector scale_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 2 * row));
+        const HVX_Vector scale = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_hf));
+        const HVX_VectorPair p16 = Q6_Wh_vunpack_Vb(vq);
+        const HVX_VectorPair p32 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16));
+        HVX_Vector values = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(p32)), scale);
+        *(HVX_Vector *) dst = values;
+    }
+}
+
+static inline void hvx_dequantize_row_tiled_q4_k_f32(float * dst_spad, const uint8_t * src_spad, uint32_t n_k_tiles, uint32_t tile_stride, uint32_t row) {
+    const HVX_VectorPred first4 = Q6_Q_vsetq_R(4);
+
+    for (uint32_t k_tile = 0; k_tile < n_k_tiles; ++k_tile) {
+        const uint8_t * tile = src_spad + k_tile * tile_stride;
+        float * dst = dst_spad + k_tile * HTP_TENSOR_TILE_N_COLS;
+
+        HVX_Vector vq = Q6_V_vzero();
+        for (int group = 3; group >= 0; --group) {
+            vq = Q6_V_vmux_QVV(first4, hvx_dequantize_gather4(hvx_vmem(tile + group * VLEN), row), Q6_V_vror_VR(vq, VLEN - 4));
+        }
+
+        const HVX_Vector lo = Q6_V_vand_VV(vq, Q6_Vb_vsplat_R(0x0F));
+        const HVX_Vector hi = Q6_Vub_vlsr_VubR(vq, 4);
+        vq = Q6_V_lo_W(Q6_W_vshuff_VVR(hi, lo, -1));
+
+        const HVX_Vector scales = hvx_vmem(tile + 512);
+        const HVX_Vector scale_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 4 * row));
+        const HVX_Vector scale = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_hf));
+        const HVX_Vector offset_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 4 * row + 2));
+        const HVX_Vector offset = Q6_V_lo_W(hvx_vec_f16_to_f32(offset_hf));
+
+        const HVX_VectorPair p16 = Q6_Wh_vunpack_Vb(vq);
+        const HVX_VectorPair p32 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16));
+        HVX_Vector values = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(p32)), scale);
+        *(HVX_Vector *) dst = hvx_vec_add_f32_f32(values, offset);
+    }
+}
+
+static inline void hvx_dequantize_row_tiled_q5_k_f32(float * dst_spad, const uint8_t * src_spad, uint32_t n_k_tiles, uint32_t tile_stride, uint32_t row) {
+    const HVX_VectorPred first4 = Q6_Q_vsetq_R(4);
+    const HVX_Vector one = Q6_Vb_vsplat_R(1);
+
+    for (uint32_t k_tile = 0; k_tile < n_k_tiles; ++k_tile) {
+        const uint8_t * tile = src_spad + k_tile * tile_stride;
+        float * dst = dst_spad + k_tile * HTP_TENSOR_TILE_N_COLS;
+
+        const HVX_Vector p5 = hvx_dequantize_gather4(hvx_vmem(tile + 640), row);
+        HVX_Vector vq = Q6_V_vzero();
+        HVX_Vector vh = Q6_V_vzero();
+        for (int group = 3; group >= 0; --group) {
+            vq = Q6_V_vmux_QVV(first4, hvx_dequantize_gather4(hvx_vmem(tile + group * VLEN), row), Q6_V_vror_VR(vq, VLEN - 4));
+            vh = Q6_V_vmux_QVV(first4, Q6_Vub_vlsr_VubR(p5, 2 * group), Q6_V_vror_VR(vh, VLEN - 4));
+        }
+
+        const HVX_Vector lo = Q6_V_vor_VV(Q6_V_vand_VV(vq, Q6_Vb_vsplat_R(0x0F)), Q6_Vw_vasl_VwR(Q6_V_vand_VV(vh, one), 4));
+        const HVX_Vector hi = Q6_V_vor_VV(Q6_Vub_vlsr_VubR(vq, 4), Q6_Vw_vasl_VwR(Q6_V_vand_VV(Q6_Vub_vlsr_VubR(vh, 1), one), 4));
+        vq = Q6_V_lo_W(Q6_W_vshuff_VVR(hi, lo, -1));
+
+        const HVX_Vector scales = hvx_vmem(tile + 512);
+        const HVX_Vector scale_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 4 * row));
+        const HVX_Vector scale = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_hf));
+        const HVX_Vector offset_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 4 * row + 2));
+        const HVX_Vector offset = Q6_V_lo_W(hvx_vec_f16_to_f32(offset_hf));
+
+        const HVX_VectorPair p16 = Q6_Wh_vunpack_Vb(vq);
+        const HVX_VectorPair p32 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16));
+        HVX_Vector values = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(p32)), scale);
+        *(HVX_Vector *) dst = hvx_vec_add_f32_f32(values, offset);
+    }
+}
+
+static inline void hvx_dequantize_row_tiled_q6_k_f32(float * dst_spad, const uint8_t * src_spad, uint32_t n_k_tiles, uint32_t tile_stride, uint32_t row) {
+    const HVX_VectorPred first4 = Q6_Q_vsetq_R(4);
+    const HVX_VectorPred first16 = Q6_Q_vsetq_R(16 * sizeof(float));
+    const HVX_Vector mask_0f = Q6_Vb_vsplat_R(0x0F);
+    const HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03);
+
+    for (uint32_t k_tile = 0; k_tile < n_k_tiles; ++k_tile) {
+        const uint8_t * tile = src_spad + k_tile * tile_stride;
+        float * dst = dst_spad + k_tile * HTP_TENSOR_TILE_N_COLS;
+
+        HVX_Vector vq = Q6_V_vzero();
+        for (int group = 7; group >= 0; --group) {
+            const HVX_Vector lo_plane = Q6_V_vror_VR(hvx_vmem(tile + (group >> 1) * VLEN), 4 * row);
+            const HVX_Vector hi_plane = Q6_V_vror_VR(hvx_vmem(tile + 512 + (group >> 2) * VLEN), 4 * row);
+            const HVX_Vector lo = (group & 1) ? Q6_Vub_vlsr_VubR(lo_plane, 4) : Q6_V_vand_VV(lo_plane, mask_0f);
+            const HVX_Vector hi = Q6_Vub_vlsr_VubR(hi_plane, 2 * (group & 3));
+            const HVX_Vector packed = Q6_V_vor_VV(lo, Q6_Vw_vasl_VwR(Q6_V_vand_VV(hi, mask_03), 4));
+            vq = Q6_V_vmux_QVV(first4, packed, Q6_V_vror_VR(vq, VLEN - 4));
+        }
+
+        const HVX_Vector scales = hvx_vmem(tile + 768);
+        const HVX_Vector scale_lo_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 2 * row));
+        const HVX_Vector scale_hi_hf = hvx_vec_repl_f16(Q6_V_vror_VR(scales, 64 + 2 * row));
+        const HVX_Vector scale_lo = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_lo_hf));
+        const HVX_Vector scale_hi = Q6_V_lo_W(hvx_vec_f16_to_f32(scale_hi_hf));
+        const HVX_Vector scale = Q6_V_vmux_QVV(first16, scale_lo, scale_hi);
+        const HVX_VectorPair p16 = Q6_Wh_vunpack_Vb(Q6_Vb_vsub_VbVb(vq, Q6_Vb_vsplat_R(32)));
+        const HVX_VectorPair p32 = Q6_Ww_vunpack_Vh(Q6_V_lo_W(p16));
+        *(HVX_Vector *) dst = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(p32)), scale);
+    }
+}
+
+static inline HVX_Vector hvx_dequantize_tiled_f16_gather4(const uint8_t * p, uint32_t r_byte) {
+    HVX_Vector w0  = Q6_V_vror_VR(hvx_vmem(p + 0 * VLEN), r_byte);
+    HVX_Vector w1  = Q6_V_vror_VR(hvx_vmem(p + 1 * VLEN), r_byte);
+    HVX_Vector w01 = Q6_V_lo_W(Q6_W_vshuff_VVR(w1, w0, 4));
+    HVX_Vector w2  = Q6_V_vror_VR(hvx_vmem(p + 2 * VLEN), r_byte);
+    HVX_Vector w3  = Q6_V_vror_VR(hvx_vmem(p + 3 * VLEN), r_byte);
+    HVX_Vector w23 = Q6_V_lo_W(Q6_W_vshuff_VVR(w3, w2, 4));
+    return Q6_V_lo_W(Q6_W_vshuff_VVR(w23, w01, 8));
+}
+
+static inline void hvx_dequantize_row_tiled_f16_f32(float * dst_spad, const uint8_t * src_spad, uint32_t n_k_tiles, uint32_t tile_stride, uint32_t row) {
+    const uint32_t r_byte = 4 * row;
+
+    for (uint32_t k_tile = 0; k_tile < n_k_tiles; ++k_tile) {
+        const uint8_t * tile = src_spad + k_tile * tile_stride;
+        float * dst = dst_spad + k_tile * HTP_TENSOR_TILE_N_COLS;
+
+        HVX_Vector g0 = hvx_dequantize_tiled_f16_gather4(tile + 0 * VLEN, r_byte);
+        HVX_Vector g1 = hvx_dequantize_tiled_f16_gather4(tile + 4 * VLEN, r_byte);
+        HVX_Vector g01 = Q6_V_lo_W(Q6_W_vshuff_VVR(g1, g0, 16));
+
+        HVX_Vector g2 = hvx_dequantize_tiled_f16_gather4(tile + 8 * VLEN, r_byte);
+        HVX_Vector g3 = hvx_dequantize_tiled_f16_gather4(tile + 12 * VLEN, r_byte);
+        HVX_Vector g23 = Q6_V_lo_W(Q6_W_vshuff_VVR(g3, g2, 16));
+
+        HVX_Vector g = Q6_V_lo_W(Q6_W_vshuff_VVR(g23, g01, 32));
+        *(HVX_Vector *) dst = Q6_V_lo_W(hvx_vec_f16_to_f32(g));
+    }
+}
 
 #endif // HVX_QUANT_H

ggml/src/ggml-hexagon/htp/main.c

@@ -28,8 +28,8 @@
 #include "ggml-common.h"
 #include "hex-bitmap.h"
 #include "htp-ctx.h"
+#include "htp-dcache.h"
 #include "htp-ops.h"
-#include "htp-tensor.h"
 #include "htp_iface.h"
 #include "work-queue.h"
 #include "hex-profile.h"
@@ -835,7 +835,7 @@ static int op_fence(struct htp_ops_context * octx) {
     atomic_uint * sync_fence = (atomic_uint *) (uintptr_t) sync->data;
 
     if (mode == 1) {
-        htp_flush_dirty_ranges(ctx);
+        htp_dcache_flush_dirty_ranges(ctx);
 
         htp_mdev_group_barrier(octx);
 
@@ -1211,7 +1211,7 @@ static int proc_op_req(struct htp_ops_context * octx, struct htp_buf_desc * bufs
             src->ne[0], src->ne[1], src->ne[2], src->ne[3]);
     }
 
-    htp_tensor_flush_all(octx->ctx, octx->src, HTP_OP_MAX_INPUTS);
+    htp_dcache_flush_dirty_tensors(octx->ctx, octx->src, HTP_OP_MAX_INPUTS);
 
     // Prep output tensors
     for (uint32_t i = 0; i < HTP_OP_MAX_OUTPUTS; i++) {
@@ -1227,7 +1227,7 @@ static int proc_op_req(struct htp_ops_context * octx, struct htp_buf_desc * bufs
             dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3]);
     }
 
-    htp_tensor_dirty_all(octx->ctx, octx->dsts, HTP_OP_MAX_OUTPUTS);
+    htp_dcache_mark_dirty_tensors(octx->ctx, octx->dsts, HTP_OP_MAX_OUTPUTS);
 
     htp_mdev_group_barrier(octx);
 

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

@@ -4,6 +4,7 @@
 #include <stdint.h>
 #include <stddef.h>
 #include "htp-ops.h"
+#include "htp-tensor.h"
 #include "hex-fastdiv.h"
 #include "hex-common.h"
 #include "htp-vtcm.h"
@@ -19,45 +20,10 @@ extern "C" {
 #define HTP_MM_HMX_TILE_N_ELMS 1024
 #define HTP_MM_HMX_MIN_NROWS   4
 
-// --- Weight Repacked Tile Sizes ---
-#define HTP_MM_WEIGHT_TILE_SIZE_Q4_0   576
-#define HTP_MM_WEIGHT_TILE_SIZE_Q4_1   640
-#define HTP_MM_WEIGHT_TILE_SIZE_Q8_0   1088
-#define HTP_MM_WEIGHT_TILE_SIZE_IQ4_NL 576
-#define HTP_MM_WEIGHT_TILE_SIZE_MXFP4  544
-// Q5_K: the Q4_1 tile (640) followed by a 128-byte plane with the 5th bit of every quant, transposed so that
-//   plane byte l holds the eight flags of lane l: bit 2i = low nibble of nibble vector i, bit 2i+1 = high nibble
-#define HTP_MM_WEIGHT_TILE_SIZE_Q5_K   768
-// Q6_K native 6-bit tile (32 rows x 32 k), vrmpy-ready: byte 4*row+b of a vector holds k = 4*group+b
-//   vectors 0..3: low nibbles, vector i holds group 2i (low nibble) and group 2i+1 (high nibble)
-//   vectors 4..5: high 2 bits, vector m holds groups 4m..4m+3 at bit offsets 0,2,4,6
-//   vector 6: fp16 scales per row, d * scales[]: k 0..15 in lanes 0..31, k 16..31 in lanes 32..63
-#define HTP_MM_WEIGHT_TILE_SIZE_Q6_K   896
-// Q3_K native 3-bit tile, vrmpy-ready like Q6_K
-//   vectors 0..1: low 2 bits, vector m holds groups 4m..4m+3 at bit offsets 0,2,4,6
-//   vector 2: bit g set where the hmask bit of group g is clear (quant = low 2 bits - 4)
-//   vector 3: fp16 scales per row, d * (scales[] - 32): k 0..15 in lanes 0..31, k 16..31 in lanes 32..63
-#define HTP_MM_WEIGHT_TILE_SIZE_Q3_K   512
-// Q2_K native 2-bit tile with the ggml scale bytes, vrmpy-ready like Q6_K (same bits per weight as block_q2_K)
-//   bytes   0..255: unsigned 2-bit quants, vector m holds groups 4m..4m+3 at bit offsets 0,2,4,6
-//   bytes 256..319: scales[] bytes (min << 4 | scale), sub-block s of row r at 32 * s + r
-//   bytes 320..335: 8 fp16, tile j of a super-block: d (j < 4) or -dmin (j >= 4) of rows 8 * (j % 4) .. 8 * (j % 4) + 7
-#define HTP_MM_WEIGHT_TILE_SIZE_Q2_K   336
-
-// --- Weight Repacked Aligned Tile Sizes ---
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q4_0   640
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q4_1   640
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q8_0   1152
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_IQ4_NL 640
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_MXFP4  640
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q5_K   768
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q6_K   896
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q3_K   512
-#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q2_K   384
-
 // --- Activation Tiled Block Sizes (including padding) ---
 #define HTP_MM_ACT_TILE_SIZE_Q8_0      1152
 #define HTP_MM_ACT_TILE_SIZE_Q8_1      1280
+#define HTP_MM_ACT_TILE_SIZE_F16       2048
 
 #define HTP_MM_MAX_PREFETCH 16
 
@@ -74,11 +40,12 @@ enum htp_mm_kernel_type {
 
     // HMX paths
     HTP_MM_KERNEL_HMX_2D,
-    HTP_MM_KERNEL_HMX_F16_BATCHED,
+    HTP_MM_KERNEL_HMX_BATCHED,
 
     // HVX floating-point paths
     HTP_MM_KERNEL_HVX_F16_F16_VTCM,
     HTP_MM_KERNEL_HVX_F32_F32_VTCM,
+    HTP_MM_KERNEL_HVX_BF16_F32_VTCM,
 
     // HVX quantized paths
     HTP_MM_KERNEL_HVX_QUANT_ROW,      // standard row-wise parallel quantization
@@ -212,57 +179,6 @@ next_nc:
     return 0;
 }
 
-// --- Tile Size Helpers ---
-static inline uint32_t htp_mm_get_weight_tile_size(int weight_type) {
-    switch (weight_type) {
-        case HTP_TYPE_Q4_0:
-        case HTP_TYPE_IQ4_NL:
-            return HTP_MM_WEIGHT_TILE_SIZE_Q4_0;
-        case HTP_TYPE_Q4_1:
-        case HTP_TYPE_Q4_K:
-            return HTP_MM_WEIGHT_TILE_SIZE_Q4_1;
-        case HTP_TYPE_Q8_0:
-            return HTP_MM_WEIGHT_TILE_SIZE_Q8_0;
-        case HTP_TYPE_Q5_K:
-            return HTP_MM_WEIGHT_TILE_SIZE_Q5_K;
-        case HTP_TYPE_Q6_K:
-            return HTP_MM_WEIGHT_TILE_SIZE_Q6_K;
-        case HTP_TYPE_Q3_K:
-            return HTP_MM_WEIGHT_TILE_SIZE_Q3_K;
-        case HTP_TYPE_Q2_K:
-            return HTP_MM_WEIGHT_TILE_SIZE_Q2_K;
-        case HTP_TYPE_MXFP4:
-            return HTP_MM_WEIGHT_TILE_SIZE_MXFP4;
-        default:
-            return 0;
-    }
-}
-
-static inline uint32_t htp_mm_get_weight_aligned_tile_size(int weight_type) {
-    switch (weight_type) {
-        case HTP_TYPE_Q4_0:
-        case HTP_TYPE_IQ4_NL:
-            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q4_0;
-        case HTP_TYPE_Q4_1:
-        case HTP_TYPE_Q4_K:
-            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q4_1;
-        case HTP_TYPE_Q8_0:
-            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q8_0;
-        case HTP_TYPE_Q5_K:
-            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q5_K;
-        case HTP_TYPE_Q6_K:
-            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q6_K;
-        case HTP_TYPE_Q3_K:
-            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q3_K;
-        case HTP_TYPE_Q2_K:
-            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q2_K;
-        case HTP_TYPE_MXFP4:
-            return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_MXFP4;
-        default:
-            return 0;
-    }
-}
-
 // weight types whose tiles carry a per-block offset (x = d * q + m): the activations need block sums (q8_1)
 // (Q2_K: per-16 k sums, q8_1_s16)
 static inline bool htp_mm_weight_has_offset(int weight_type) {
@@ -283,27 +199,17 @@ static inline size_t htp_mm_q8_1_tiled_row_size(uint32_t ne) {
     return nb_32 * HTP_MM_ACT_TILE_SIZE_Q8_1;
 }
 
-static inline size_t htp_mm_get_tiled_row_stride(int weight_type, uint32_t k) {
-    uint32_t nb = (k + QK_Q4_0_TILED - 1) / QK_Q4_0_TILED;
-    switch (weight_type) {
-        case HTP_TYPE_Q4_0:
-        case HTP_TYPE_IQ4_NL:
-        case HTP_TYPE_Q4_1:
-        case HTP_TYPE_Q4_K:
-        case HTP_TYPE_Q8_0:
-        case HTP_TYPE_Q5_K:
-        case HTP_TYPE_Q6_K:
-        case HTP_TYPE_Q3_K:
-        case HTP_TYPE_Q2_K:
-        case HTP_TYPE_MXFP4:
-            return (size_t) nb * htp_mm_get_weight_tile_size(weight_type);
-        case HTP_TYPE_F16:
-            return (size_t) k * sizeof(__fp16);
-        case HTP_TYPE_F32:
-            return (size_t) k * sizeof(float);
-        default:
-            return 0;
+static inline size_t htp_mm_f16_tiled_row_size(uint32_t ne) {
+    const uint32_t ne_padded = ((ne + 127) / 128) * 128;
+    const uint32_t nb_32 = ne_padded / 32;
+    return nb_32 * HTP_MM_ACT_TILE_SIZE_F16;
+}
+
+static inline size_t htp_mm_tiled_act_row_size(int weight_type, uint32_t ne) {
+    if (weight_type == HTP_TYPE_F16) {
+        return htp_mm_f16_tiled_row_size(ne);
     }
+    return htp_mm_weight_has_offset(weight_type) ? htp_mm_q8_1_tiled_row_size(ne) : htp_mm_q8_0_tiled_row_size(ne);
 }
 
 static inline size_t htp_mm_round_up(size_t n, size_t m) {
@@ -318,8 +224,8 @@ static inline void htp_mm_hmx_get_2d_chunk_costs(
     int wtype, uint32_t k, bool pipeline, uint32_t aligned_tile_size,
     size_t * size_per_n_out, size_t * size_per_m_out, size_t * size_per_mn_out
 ) {
-    const bool is_quant = (wtype != HTP_TYPE_F16 && wtype != HTP_TYPE_F32);
-    const size_t row_stride = htp_mm_get_tiled_row_stride(wtype, k);
+    const bool is_quant = (wtype != HTP_TYPE_F16 && wtype != HTP_TYPE_F32 && wtype != HTP_TYPE_BF16);
+    const size_t row_stride = htp_tensor_get_row_size(wtype, k);
     const size_t vec_dot_size = k * sizeof(uint16_t);
     const uint32_t n_k_tiles = k / HTP_MM_HMX_TILE_N_COLS;
     const size_t qweight_row_stride = is_quant ? (size_t)(n_k_tiles * aligned_tile_size) / 32 : 0;
@@ -404,7 +310,7 @@ static inline void htp_mm_hmx_vtcm_layout_build(
 ) {
     size_t off = 0;
 
-    if (kernel_type == HTP_MM_KERNEL_HMX_F16_BATCHED) {
+    if (kernel_type == HTP_MM_KERNEL_HMX_BATCHED) {
         const size_t vec_dot_size     = k * sizeof(uint16_t);
         const size_t act_head_stride   = mc * k;
         const size_t weight_area_size  = hex_align_up(nc * vec_dot_size, HTP_MM_HMX_TILE_SIZE);
@@ -450,8 +356,8 @@ static inline void htp_mm_hmx_vtcm_layout_build(
         off = off_group_a + hex_smax(group_b_size, group_c_size);
     } else {
         // HTP_MM_KERNEL_HMX_2D
-        const bool is_quant = (wtype != HTP_TYPE_F16 && wtype != HTP_TYPE_F32);
-        const size_t row_stride = htp_mm_get_tiled_row_stride(wtype, k);
+        const bool is_quant = (wtype != HTP_TYPE_F16 && wtype != HTP_TYPE_F32 && wtype != HTP_TYPE_BF16);
+        const size_t row_stride = htp_tensor_get_row_size(wtype, k);
         const size_t vec_dot_size = k * sizeof(uint16_t);
         const uint32_t n_k_tiles = k / HTP_MM_HMX_TILE_N_COLS;
 
@@ -527,28 +433,13 @@ static inline void htp_mm_hvx_vtcm_layout_build(
     size_t dst_sz     = 0;
     size_t act_raw_sz = 0;
 
-    const bool is_repack = (wtype == HTP_TYPE_Q4_0 || wtype == HTP_TYPE_Q4_1 ||
-                            wtype == HTP_TYPE_Q8_0 || wtype == HTP_TYPE_IQ4_NL ||
-                            wtype == HTP_TYPE_MXFP4 || wtype == HTP_TYPE_Q6_K ||
-                            wtype == HTP_TYPE_Q4_K || wtype == HTP_TYPE_Q5_K ||
-                            wtype == HTP_TYPE_Q3_K || wtype == HTP_TYPE_Q2_K);
-
     if (is_fused_nx) {
-        const size_t src0_row_size_padded = hex_round_up(src0_row_size, 128);
+        uint32_t aligned_tile_size = htp_tensor_get_tile_padded_size(wtype);
+        uint32_t n_k_tiles = hex_round_up(ne10, 32) / 32;
+        uint32_t tile_row_size = n_k_tiles * aligned_tile_size;
+        size_t weight_sz_per_thread = hex_round_up(n_prefetch * tile_row_size, 128);
 
-        size_t weight_sz_per_thread = 0;
-
-        if (is_repack) {
-            uint32_t aligned_tile_size = htp_mm_get_weight_aligned_tile_size(wtype);
-            uint32_t n_k_tiles = hex_round_up(ne10, 32) / 32;
-            uint32_t tile_row_size = n_k_tiles * aligned_tile_size;
-
-            weight_sz_per_thread = hex_round_up(n_prefetch * tile_row_size, 128);
-        } else {
-            weight_sz_per_thread = hex_round_up(n_prefetch * src0_row_size_padded, 128);
-        }
-
-        size_t tiled_act_row_size = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10);
+        size_t tiled_act_row_size = htp_mm_tiled_act_row_size(wtype, ne10);
         size_t q_act_sz = hex_round_up(tiled_act_row_size * act_nrows, 128);
         size_t raw_row_size = hex_round_up(ne10 * sizeof(float), QK_Q8_0_TILED * sizeof(float));
 
@@ -558,20 +449,13 @@ static inline void htp_mm_hvx_vtcm_layout_build(
         dst_sz     = 0;
         act_raw_sz = hex_round_up(raw_row_size * act_nrows, 128);
     } else if (is_matmul_id) {
-        const size_t src0_row_size_padded = htp_mm_round_up(src0_row_size, 128);
-        const size_t act_row_size_tiled = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10)
-                                                                          : htp_mm_q8_0_tiled_row_size(ne10);
+        const size_t act_row_size_tiled = htp_mm_tiled_act_row_size(wtype, ne10);
+        const uint32_t aligned_tile_size = htp_tensor_get_tile_padded_size(wtype);
+        const uint32_t n_k_tiles         = ne10 / 32;
+        const uint32_t tile_row_size     = n_k_tiles * aligned_tile_size;
+        size_t src0_sz_per_thread        = htp_mm_round_up(n_prefetch * tile_row_size, 256);
 
-        size_t src0_sz_per_thread = htp_mm_round_up(n_prefetch * src0_row_size_padded, 256);
-        act_sz                    = htp_mm_round_up(act_row_size_tiled * act_nrows, 256);
-
-        if (is_repack) {
-            const uint32_t aligned_tile_size = htp_mm_get_weight_aligned_tile_size(wtype);
-            const uint32_t n_k_tiles         = ne10 / 32;
-            const uint32_t tile_row_size     = n_k_tiles * aligned_tile_size;
-            size_t repacked_vtcm_size        = htp_mm_round_up(n_prefetch * tile_row_size, 256);
-            src0_sz_per_thread               = repacked_vtcm_size;
-        }
+        act_sz = htp_mm_round_up(act_row_size_tiled * act_nrows, 256);
 
         size_t raw_row_size = hex_round_up(ne10 * sizeof(float), QK_Q8_0_TILED * sizeof(float));
 
@@ -592,7 +476,8 @@ static inline void htp_mm_hvx_vtcm_layout_build(
                 act_raw_sz = hex_round_up(hex_round_up(ne10 * sizeof(float), 128) * act_nrows, 128);
                 break;
             }
-            case HTP_MM_KERNEL_HVX_F32_F32_VTCM: {
+            case HTP_MM_KERNEL_HVX_F32_F32_VTCM:
+            case HTP_MM_KERNEL_HVX_BF16_F32_VTCM: {
                 size_t f32_act_row_size = htp_mm_round_up(ne10 * 4, 128);
                 act_sz     = htp_mm_round_up(f32_act_row_size * act_nrows, 256);
                 src0_sz    = htp_mm_round_up(n_prefetch * src0_row_size_padded, 256) * n_threads;
@@ -602,21 +487,15 @@ static inline void htp_mm_hvx_vtcm_layout_build(
             }
             case HTP_MM_KERNEL_HVX_QUANT_BLOCK:
             case HTP_MM_KERNEL_HVX_QUANT_ROW: {
-                size_t q_act_row_size = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10);
+                size_t q_act_row_size = htp_mm_tiled_act_row_size(wtype, ne10);
+                uint32_t aligned_tile_size = htp_tensor_get_tile_padded_size(wtype);
+                uint32_t n_k_tiles = (ne10 + 31) / 32;
+                uint32_t tile_row_size = n_k_tiles * aligned_tile_size;
+                size_t tiled_vtcm_size = htp_mm_round_up(n_prefetch * tile_row_size, 256);
 
-                src0_sz = htp_mm_round_up(n_prefetch * src0_row_size_padded, 256);
+                src0_sz = tiled_vtcm_size * n_threads;
                 act_sz  = htp_mm_round_up(q_act_row_size * act_nrows, 256);
 
-                src0_sz = src0_sz * n_threads;
-
-                if (is_repack) {
-                    uint32_t aligned_tile_size = htp_mm_get_weight_aligned_tile_size(wtype);
-                    uint32_t n_k_tiles = ne10 / 32;
-                    uint32_t tile_row_size = n_k_tiles * aligned_tile_size;
-                    size_t repacked_vtcm_size = htp_mm_round_up(n_prefetch * tile_row_size, 256);
-                    src0_sz = repacked_vtcm_size * n_threads;
-                }
-
                 size_t dst_slice_per_thread = (dst_nrows > 0 && act_nrows == 1) ? htp_mm_round_up((dst_row_size + n_threads - 1) / n_threads, 128) : 0;
                 dst_sz = dst_slice_per_thread * n_threads;
                 size_t raw_row_size = hex_round_up(ne10 * sizeof(float), QK_Q8_0_TILED * sizeof(float));
@@ -687,9 +566,7 @@ static inline bool htp_mm_hvx_solve_vtcm_params(
     const size_t avail_act = vtcm_budget - fixed_bytes;
     size_t row_size = 0;
     if (kernel_type == HTP_MM_KERNEL_HVX_QUANT_ROW || kernel_type == HTP_MM_KERNEL_HVX_QUANT_BLOCK) {
-        row_size = htp_mm_weight_has_offset(wtype)
-                 ? htp_mm_q8_1_tiled_row_size(ne10)
-                 : htp_mm_q8_0_tiled_row_size(ne10);
+        row_size = htp_mm_tiled_act_row_size(wtype, ne10);
     } else if (kernel_type == HTP_MM_KERNEL_HVX_F16_F16_VTCM) {
         row_size = hex_round_up(ne10 * 2, 128);
     } else {
@@ -749,7 +626,7 @@ static inline size_t htp_mm_hmx_get_batched_vtcm_size(
     int wtype, uint32_t k, size_t mc, size_t nc, uint32_t group_size, bool pipeline, uint32_t act_threads, size_t bias_size) {
     (void)pipeline;
     struct htp_mm_hmx_vtcm_layout L;
-    htp_mm_hmx_vtcm_layout_build(&L, HTP_MM_KERNEL_HMX_F16_BATCHED, wtype, k, mc, nc, group_size, false, act_threads, 0, bias_size);
+    htp_mm_hmx_vtcm_layout_build(&L, HTP_MM_KERNEL_HMX_BATCHED, wtype, k, mc, nc, group_size, false, act_threads, 0, bias_size);
     return L.total_bytes;
 }
 

scripts/snapdragon/run.py

@@ -27,6 +27,7 @@ MANAGED_ENV_NAMES = (
     "GGML_HEXAGON_OPPOLL",
     "GGML_HEXAGON_OPFILTER",
     "GGML_HEXAGON_OPFUSION",
+    "GGML_HEXAGON_RT_REPACK",
     "GGML_HEXAGON_VMEM",
     "GGML_HEXAGON_MBUF",
     "GGML_HEXAGON_MM_SELECT",
@@ -167,6 +168,7 @@ def main():
     parser.add_argument("--hex-oppoll", default="1", help="Enable (1) or Disable (0) polling for NPU opbatch completion (GGML_HEXAGON_OPPOLL) (default: 1)")
     parser.add_argument("--hex-opfilter", help="Regex pattern to filter/select which operators are offloaded to NPU (GGML_HEXAGON_OPFILTER)")
     parser.add_argument("--hex-opfusion", help="NPU graph node fusion optimization level (0: disabled, 1: enabled) (GGML_HEXAGON_OPFUSION)")
+    parser.add_argument("--hex-rt-repack", help="Runtime weight repacking mode (0: disabled, 1: quantized only, 2: always) (GGML_HEXAGON_RT_REPACK)")
     parser.add_argument("--hex-vmem", help="Maximum NPU VMEM size limit in MB to allocate (GGML_HEXAGON_VMEM)")
     parser.add_argument("--hex-mbuf", help="Host buffer size limits in MB (supports K/M/G suffix): <dyn>[,<static>[,<total>]] (default: 512,1024,0) (GGML_HEXAGON_MBUF)")
     parser.add_argument("--hex-mm-select", help="Select MUL_MAT and MUL_MAT_ID kernel (GGML_HEXAGON_MM_SELECT) 2:HMX,1:HVX,0:disable")
@@ -311,6 +313,7 @@ def main():
     set_env("GGML_HEXAGON_OPPOLL", args.hex_oppoll)
     set_env("GGML_HEXAGON_OPFILTER", args.hex_opfilter)
     set_env("GGML_HEXAGON_OPFUSION", args.hex_opfusion)
+    set_env("GGML_HEXAGON_RT_REPACK", args.hex_rt_repack)
     set_env("GGML_HEXAGON_VMEM", args.hex_vmem)
     set_env("GGML_HEXAGON_MBUF", args.hex_mbuf)
     set_env("GGML_HEXAGON_MM_SELECT", args.hex_mm_select)

tests/test-backend-ops.cpp

@@ -10741,11 +10741,14 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
 
     // testing very large N
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F32, 16, 2105344, 32, {1, 1}, {1, 1}));
+
+#if 0
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 16, 8400000, 32, {1, 1}, {1, 1}));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 16, 8400000, 32, {1, 1}, {1, 1}));
     // same shapes with a repeat in dim 3
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 16, 8400000, 32, {1, 1}, {2, 1}));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q4_0, GGML_TYPE_F32, 16, 8400000, 32, {1, 1}, {2, 1}));
+#endif
 
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 31, 509, 2051, {1, 1}, {1, 1}));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 32, 509, 2112, {1, 1}, {1, 1}));