Commit 61dbf4922 for llama.cpp
commit 61dbf49228c8a96b72e40c89d93cb7c003411f8c
Author: Ehsan Bateni <ebateni@qti.qualcomm.com>
Date: Fri Nov 13 23:57:05 2026 -0500
Hexagon: add support for IQ2_XXS and NVFP4 (#31954)
* ggml-hexagon: IQ2_XXS and NVFP4 tiled repack layouts
* ggml-hexagon: IQ2_XXS HVX matmul kernel
* ggml-hexagon: NVFP4 HVX matmul kernel
* ggml-hexagon: NVFP4 per-tensor scale
* ggml-hexagon: HMX paths for IQ2_XXS and NVFP4
* ggml-hexagon: supports_op coverage for NVFP4 and IQ2_XXS
* fix HMX IQ2_XXS dequant scale underflow in fp16
* hexagon: IQ2_XXS 2-gather tile decode with compact grid, small-weight tests
* IQ2_XXS vpopcount sign parity and gather pipelining
* pipeline IQ2_XXS HMX dequant gathers and add large_w tests
* hex-iq2-xxs: rebase and cleanup
Remove unused helper functions.
Add lut_size to the overall vtcm required size.
Trim down (way way down) the new tests.
* hex-iq2-xxs: actually trim down the tests
---------
Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
diff --git a/docs/backend/snapdragon/developer.md b/docs/backend/snapdragon/developer.md
index 69a8d5f7d..c220ad08a 100644
--- a/docs/backend/snapdragon/developer.md
+++ b/docs/backend/snapdragon/developer.md
@@ -117,6 +117,13 @@ Block-32 formats use a tile size of 32 rows x 32 columns (`K_cols = 32`):
Kernel converts E8M0 bias to IEEE FP16 bias via `vh = clamp(e - 112, 0, 30) << 10` to form a power-of-two FP16 scale.
Each 4-bit FP4 nibble is unpacked via `mxfp4_to_fp16_lut` and multiplied by this scale.
+- **NVFP4** (576 bytes per tile, `HTP_TENSOR_TILE_SIZE_NVFP4`, lossless tiling with same size as `block_nvfp4`):
+ - A 64-element block spans 2 consecutive k-tiles; k-tile `h` holds sub-blocks `2h` and `2h+1`.
+ - **Quants** (bytes 0..511): same nibble transposition as `Q4_0`, holding E2M1 values in element order
+ (within a 16-element sub-block, `qs` byte `j` holds elements `j` and `j+8`).
+ - **Scales** (bytes 512..575): UE4M3 bytes, `tile[512 + 32*u + row]` for sub-block `u = 0..1` of the k-tile.
+ - The per-tensor F32 scale is not part of the tile.
+
- **Q4_1** (640 bytes per tile, `HTP_TENSOR_TILE_SIZE_Q4_1`):
- **Quants** (bytes 0..511, 512 bytes):
Same 4-bit nibble transposition and storage layout as `Q4_0`.
@@ -192,6 +199,10 @@ Byte ordering: `Q2_K`, `Q3_K`, and `Q6_K` use `vrmpy` byte order within each 128
- `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`.
+- **IQ2_XXS** (264 bytes per tile, `HTP_TENSOR_TILE_SIZE_IQ2_XXS`, lossless tiling with same size as `block_iq2_xxs`):
+ - `bytes 0..255`: `qs[4*kt .. 4*kt+3]` (grid indices, signs and scales of one 32-element group) of row `r` at `8*r`.
+ - `bytes 256..263`: 4 FP16 `d` in tile `j` (0..7) of super-block: rows `4*j..4*j+3`.
+
- **Q3_K** (512 bytes per tile, `HTP_TENSOR_TILE_SIZE_Q3_K`):
- `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`).
diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
index ab3854703..f9daf10f2 100644
--- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp
+++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
@@ -3507,6 +3507,10 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) {
"please update hexagon_type to match ggml_type");
static_assert((unsigned int) HTP_TYPE_MXFP4 == (unsigned int) GGML_TYPE_MXFP4,
"please update hexagon_type to match ggml_type");
+ static_assert((unsigned int) HTP_TYPE_NVFP4 == (unsigned int) GGML_TYPE_NVFP4,
+ "please update hexagon_type to match ggml_type");
+ static_assert((unsigned int) HTP_TYPE_IQ2_XXS == (unsigned int) GGML_TYPE_IQ2_XXS,
+ "please update hexagon_type to match ggml_type");
static_assert((unsigned int) HTP_TYPE_IQ4_NL == (unsigned int) GGML_TYPE_IQ4_NL,
"please update hexagon_type to match ggml_type");
static_assert((unsigned int) HTP_TYPE_Q4_K == (unsigned int) GGML_TYPE_Q4_K,
diff --git a/ggml/src/ggml-hexagon/htp-kparams.cpp b/ggml/src/ggml-hexagon/htp-kparams.cpp
index bc473defd..4c771946a 100644
--- a/ggml/src/ggml-hexagon/htp-kparams.cpp
+++ b/ggml/src/ggml-hexagon/htp-kparams.cpp
@@ -962,6 +962,8 @@ static bool htp_matmul_check_types(
case GGML_TYPE_Q8_0:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_MXFP4:
+ case GGML_TYPE_IQ2_XXS:
+ case GGML_TYPE_NVFP4:
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q6_K:
@@ -1006,6 +1008,8 @@ static bool htp_matmul_check_types(
case GGML_TYPE_Q8_0:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_MXFP4:
+ case GGML_TYPE_IQ2_XXS:
+ case GGML_TYPE_NVFP4:
case GGML_TYPE_Q4_K:
case GGML_TYPE_Q5_K:
case GGML_TYPE_Q6_K:
diff --git a/ggml/src/ggml-hexagon/htp-repack.cpp b/ggml/src/ggml-hexagon/htp-repack.cpp
index b5413791c..b60c1da63 100644
--- a/ggml/src/ggml-hexagon/htp-repack.cpp
+++ b/ggml/src/ggml-hexagon/htp-repack.cpp
@@ -1681,6 +1681,142 @@ static void repack_tiled_q1_0(void * data, const ggml_tensor * t, size_t offset,
}
}
+// --- IQ2_XXS (lossless 264 B tile: 32 rows x 8 B of qs + 4 FP16 d for rows 4*kt_local..4*kt_local+3) ---
+
+#define IQ2_XXS_QS_PER_KT 4
+#define IQ2_XXS_QS_PL_BYTES 256
+#define IQ2_XXS_D_PER_KT 4
+
+static void repack_iq2_xxs_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+ repack_kquant_to_tiled_grid<block_iq2_xxs>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_IQ2_XXS,
+ [](uint8_t * tile, const block_iq2_xxs * b, int row, int kt_local) {
+ ggml_half * dd_pl = (ggml_half *) (tile + IQ2_XXS_QS_PL_BYTES);
+ memcpy(tile + row * sizeof(uint16_t) * IQ2_XXS_QS_PER_KT, &b->qs[kt_local * IQ2_XXS_QS_PER_KT],
+ sizeof(uint16_t) * IQ2_XXS_QS_PER_KT);
+ if (row / IQ2_XXS_D_PER_KT == kt_local) {
+ dd_pl[row % IQ2_XXS_D_PER_KT] = b->d;
+ }
+ });
+}
+
+static void repack_tiled_iq2_xxs(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+ unpack_kquant_from_tiled_grid<block_iq2_xxs>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_IQ2_XXS,
+ [](block_iq2_xxs * b, const uint8_t * sb_tiles, size_t tile_size, int row) {
+ for (int kt_local = 0; kt_local < QK_K / 32; kt_local++) {
+ const uint8_t * tile = sb_tiles + kt_local * tile_size;
+ const ggml_half * dd_pl = (const ggml_half *) (tile + IQ2_XXS_QS_PL_BYTES);
+ memcpy(&b->qs[kt_local * IQ2_XXS_QS_PER_KT], tile + row * sizeof(uint16_t) * IQ2_XXS_QS_PER_KT,
+ sizeof(uint16_t) * IQ2_XXS_QS_PER_KT);
+ if (row / IQ2_XXS_D_PER_KT == kt_local) {
+ b->d = dd_pl[row % IQ2_XXS_D_PER_KT];
+ }
+ }
+ });
+}
+
+// --- NVFP4 (576 B tile: Q4_0-style nibble transpose + 2 UE4M3 scales per row, one block-64 = 2 k-tiles) ---
+
+#define NVFP4_TILE_NIBBLE_BYTES 512
+#define NVFP4_KT_PER_BLOCK (QK_NVFP4 / 32)
+#define NVFP4_SUB_PER_KT (32 / QK_NVFP4_SUB)
+
+// element e (0..31) of k-tile kt_half lives in sub-block 2*kt_half + e/16, low nibble if (e%16) < 8
+static inline uint8_t nvfp4_get_nibble(const block_nvfp4 & b, int kt_half, int e) {
+ const int sub = kt_half * NVFP4_SUB_PER_KT + e / QK_NVFP4_SUB;
+ const int j = e % (QK_NVFP4_SUB / 2);
+ const uint8_t q = b.qs[sub * (QK_NVFP4_SUB / 2) + j];
+ return (e % QK_NVFP4_SUB) < QK_NVFP4_SUB / 2 ? (q & 0x0F) : (q >> 4);
+}
+
+static inline void nvfp4_set_nibble(block_nvfp4 & b, int kt_half, int e, uint8_t v) {
+ const int sub = kt_half * NVFP4_SUB_PER_KT + e / QK_NVFP4_SUB;
+ const int j = e % (QK_NVFP4_SUB / 2);
+ uint8_t & q = b.qs[sub * (QK_NVFP4_SUB / 2) + j];
+ q |= (e % QK_NVFP4_SUB) < QK_NVFP4_SUB / 2 ? v : (uint8_t) (v << 4);
+}
+
+static void repack_nvfp4_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+ GGML_ASSERT(offset == 0);
+ GGML_ASSERT(t->ne[0] % QK_NVFP4 == 0);
+ GGML_UNUSED(size);
+
+ const block_nvfp4 * src = (const block_nvfp4 *) data;
+ const int64_t blocks_per_row = t->ne[0] / QK_NVFP4;
+ const int64_t n_k_tiles = blocks_per_row * NVFP4_KT_PER_BLOCK;
+ const int64_t n_col_tiles = hex_round_up(t->ne[1], 32) / 32;
+ const size_t matrix_size = (size_t) n_col_tiles * n_k_tiles * HTP_TENSOR_TILE_SIZE_NVFP4;
+ const size_t ct_bytes = (size_t) n_k_tiles * HTP_TENSOR_TILE_SIZE_NVFP4;
+
+ for (int64_t slice = 0; slice < t->ne[2] * t->ne[3]; ++slice) {
+ const block_nvfp4 * src_slice = src + slice * t->ne[1] * blocks_per_row;
+ uint8_t * dst_slice = (uint8_t *) t->data + slice * matrix_size;
+
+ repack_parallel_for_ct(matrix_size, (int) n_col_tiles, [&](int ct0, int ct1) {
+ // Padding rows and OR-ed nibbles require zeroed tiles
+ memset(dst_slice + (size_t) ct0 * ct_bytes, 0, (size_t) (ct1 - ct0) * ct_bytes);
+
+ const int64_t r0 = (int64_t) ct0 * 32;
+ const int64_t r1 = (std::min)((int64_t) ct1 * 32, t->ne[1]);
+ for (int64_t r = r0; r < r1; ++r) {
+ const int64_t ct = r / 32;
+ const int row = (int) (r % 32);
+ for (int64_t blk = 0; blk < blocks_per_row; ++blk) {
+ const block_nvfp4 & b = src_slice[r * blocks_per_row + blk];
+ for (int h = 0; h < NVFP4_KT_PER_BLOCK; ++h) {
+ uint8_t * tile = dst_slice + (ct * n_k_tiles + blk * NVFP4_KT_PER_BLOCK + h) * HTP_TENSOR_TILE_SIZE_NVFP4;
+ for (int e = 0; e < 32; ++e) {
+ tile[(e / 2) * 32 + row] |= (uint8_t) (nvfp4_get_nibble(b, h, e) << (4 * (e & 1)));
+ }
+ for (int u = 0; u < NVFP4_SUB_PER_KT; ++u) {
+ tile[NVFP4_TILE_NIBBLE_BYTES + u * 32 + row] = b.d[h * NVFP4_SUB_PER_KT + u];
+ }
+ }
+ }
+ }
+ });
+ }
+}
+
+static void repack_tiled_nvfp4(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+ GGML_ASSERT(t->ne[0] % QK_NVFP4 == 0);
+ GGML_ASSERT(offset + size <= ggml_nbytes(t));
+
+ const size_t blocks_per_row = t->ne[0] / QK_NVFP4;
+ const size_t blocks_per_slice = t->ne[1] * blocks_per_row;
+ const size_t n_k_tiles = blocks_per_row * NVFP4_KT_PER_BLOCK;
+ const size_t n_col_tiles = hex_round_up(t->ne[1], 32) / 32;
+ const size_t matrix_size = n_col_tiles * n_k_tiles * HTP_TENSOR_TILE_SIZE_NVFP4;
+
+ const size_t first_block = offset / sizeof(block_nvfp4);
+ const size_t end_block = (offset + size + sizeof(block_nvfp4) - 1) / sizeof(block_nvfp4);
+
+ for (size_t block_index = first_block; block_index < end_block; ++block_index) {
+ const size_t slice = block_index / blocks_per_slice;
+ const size_t in_slice = block_index % blocks_per_slice;
+ const size_t r = in_slice / blocks_per_row;
+ const size_t blk = in_slice % blocks_per_row;
+ const size_t ct = r / 32;
+ const int row = (int) (r % 32);
+
+ block_nvfp4 b = {};
+ for (int h = 0; h < NVFP4_KT_PER_BLOCK; ++h) {
+ const uint8_t * tile = (const uint8_t *) t->data + slice * matrix_size +
+ (ct * n_k_tiles + blk * NVFP4_KT_PER_BLOCK + h) * HTP_TENSOR_TILE_SIZE_NVFP4;
+ for (int e = 0; e < 32; ++e) {
+ nvfp4_set_nibble(b, h, e, (tile[(e / 2) * 32 + row] >> (4 * (e & 1))) & 0x0F);
+ }
+ for (int u = 0; u < NVFP4_SUB_PER_KT; ++u) {
+ b.d[h * NVFP4_SUB_PER_KT + u] = tile[NVFP4_TILE_NIBBLE_BYTES + u * 32 + row];
+ }
+ }
+
+ const size_t block_begin = block_index * sizeof(block_nvfp4);
+ const size_t copy_begin = (std::max)(block_begin, offset);
+ const size_t copy_end = (std::min)(block_begin + sizeof(block_nvfp4), offset + size);
+ memcpy((uint8_t *) data + copy_begin - offset, (const uint8_t *) &b + copy_begin - block_begin, copy_end - copy_begin);
+ }
+}
+
// --- Top-level Repack / Unpack Dispatchers ---
void htp_repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t size) {
@@ -1707,6 +1843,12 @@ void htp_repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t siz
case GGML_TYPE_MXFP4:
repack_mxfp4_tiled(tensor, data, 0, size);
break;
+ case GGML_TYPE_NVFP4:
+ repack_nvfp4_tiled(tensor, data, 0, size);
+ break;
+ case GGML_TYPE_IQ2_XXS:
+ repack_iq2_xxs_tiled(tensor, data, 0, size);
+ break;
case GGML_TYPE_Q2_0:
repack_q2_0_tiled(tensor, data, 0, size);
break;
@@ -1768,6 +1910,12 @@ void htp_unpack_tensor_tiled(void * data, const ggml_tensor * tensor, size_t off
case GGML_TYPE_MXFP4:
repack_tiled_mxfp4(data, tensor, offset, size);
break;
+ case GGML_TYPE_NVFP4:
+ repack_tiled_nvfp4(data, tensor, offset, size);
+ break;
+ case GGML_TYPE_IQ2_XXS:
+ repack_tiled_iq2_xxs(data, tensor, offset, size);
+ break;
case GGML_TYPE_Q2_0:
repack_tiled_q2_0(data, tensor, offset, size);
break;
diff --git a/ggml/src/ggml-hexagon/htp-repack.h b/ggml/src/ggml-hexagon/htp-repack.h
index 5b71f9c5a..fc1c0d7b5 100644
--- a/ggml/src/ggml-hexagon/htp-repack.h
+++ b/ggml/src/ggml-hexagon/htp-repack.h
@@ -12,6 +12,7 @@ static inline bool htp_is_tiled_weight_type(enum ggml_type type) {
type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q4_1 ||
type == GGML_TYPE_Q5_0 || type == GGML_TYPE_Q5_1 ||
type == GGML_TYPE_MXFP4 || type == GGML_TYPE_IQ4_NL ||
+ type == GGML_TYPE_NVFP4 || type == GGML_TYPE_IQ2_XXS ||
type == GGML_TYPE_Q8_0 || type == GGML_TYPE_Q6_K ||
type == GGML_TYPE_Q2_K || type == GGML_TYPE_Q3_K ||
type == GGML_TYPE_Q4_K || type == GGML_TYPE_Q5_K ||
diff --git a/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h b/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h
index a828012dd..37c529647 100644
--- a/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h
+++ b/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h
@@ -805,6 +805,110 @@ static void dequantize_tiled_weight_to_fp16_task_q6_k(
}
}
+#define NVFP4_UE4M3_F16_MAN_SHIFT 7 // 10 - 3 mantissa bits
+#define NVFP4_UE4M3_F16_REBIAS 0x2000 // (15 - 7) << 10
+#define NVFP4_UE4M3_F16_SUB_SCALE 0x1800 // 2^-9 as fp16 bits
+#define IQ2XXS_D_SCALE_F16 0x3000 // 0.125 as fp16 bits
+
+// UE4M3 codes in the low byte of each halfword -> fp16
+static inline HVX_Vector nvfp4_ue4m3_to_f16(HVX_Vector v_code) {
+ HVX_Vector v_mask = Q6_Vh_vsplat_R(NVFP4_UE4M3_MASK);
+ HVX_Vector v_x = Q6_V_vand_VV(v_code, v_mask);
+
+ HVX_Vector v_norm = Q6_Vh_vadd_VhVh(Q6_Vh_vasl_VhR(v_x, NVFP4_UE4M3_F16_MAN_SHIFT), Q6_Vh_vsplat_R(NVFP4_UE4M3_F16_REBIAS));
+ HVX_Vector v_sub = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(v_x), Q6_Vh_vsplat_R(NVFP4_UE4M3_F16_SUB_SCALE)));
+
+ HVX_VectorPred q_sub = Q6_Q_vcmp_gt_VhVh(Q6_Vh_vsplat_R(NVFP4_UE4M3_MIN_NORMAL), v_x);
+ HVX_VectorPred q_nan = Q6_Q_vcmp_eq_VhVh(v_x, v_mask);
+
+ return Q6_V_vmux_QVV(q_nan, Q6_V_vzero(), Q6_V_vmux_QVV(q_sub, v_sub, v_norm));
+}
+
+// NVFP4 tiles are MXFP4 quants with one UE4M3 scale per 16 k (sub-block 0: k 0..15, sub-block 1: k 16..31)
+static void dequantize_tiled_weight_to_fp16_task_nvfp4(
+ const tiled_dequantize_state_t *state,
+ uint32_t start_tile, uint32_t end_tile) {
+
+ const HVX_Vector mask_h4 = Q6_Vb_vsplat_R(0x0F);
+ const HVX_Vector vlut_cvt = hvx_vmem(mxfp4_to_fp16_lut);
+
+ for (uint32_t t = start_tile; t < end_tile; t++) {
+ const uint8_t * tile_src = state->src + t * state->aligned_tile_size;
+ __fp16 * dst_ptr = state->dst + t * HTP_MM_HMX_TILE_N_ELMS;
+
+ HVX_Vector v_d = nvfp4_ue4m3_to_f16(Q6_V_lo_W(Q6_Wuh_vunpack_Vub(hvx_vmem(tile_src + NVFP4_SCALE_OFF))));
+ HVX_VectorPair vp_d = Q6_W_vshuff_VVR(v_d, v_d, -2);
+ HVX_Vector v_scale_k0 = Q6_V_lo_W(vp_d);
+ HVX_Vector v_scale_k16 = Q6_V_hi_W(vp_d);
+
+ #pragma unroll
+ for (int g = 0; g < 4; g++) {
+ const HVX_Vector v_scale = (g < 2) ? v_scale_k0 : v_scale_k16;
+
+ HVX_Vector vq = hvx_vmem(tile_src + g * 128);
+ HVX_Vector v_lo = Q6_V_vand_VV(vq, mask_h4);
+ HVX_Vector v_hi = Q6_Vub_vlsr_VubR(vq, 4);
+
+ HVX_VectorPair vp_shuf = Q6_W_vshuff_VVR(v_hi, v_lo, -1);
+ HVX_VectorPair vp_lo = Q6_Wh_vlut16_VbVhR(Q6_Vb_vshuff_Vb(Q6_V_lo_W(vp_shuf)), vlut_cvt, 0);
+ HVX_VectorPair vp_hi = Q6_Wh_vlut16_VbVhR(Q6_Vb_vshuff_Vb(Q6_V_hi_W(vp_shuf)), vlut_cvt, 0);
+
+ hvx_vmem(dst_ptr + (4 * g + 0) * 64) = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_V_lo_W(vp_lo), v_scale));
+ hvx_vmem(dst_ptr + (4 * g + 1) * 64) = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_V_hi_W(vp_lo), v_scale));
+ hvx_vmem(dst_ptr + (4 * g + 2) * 64) = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_V_lo_W(vp_hi), v_scale));
+ hvx_vmem(dst_ptr + (4 * g + 3) * 64) = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_V_hi_W(vp_hi), v_scale));
+ }
+ }
+}
+
+// IQ2_XXS: int8 grid weights from the VTCM LUT, scaled by d / 8 * (2 * scale + 1) per row
+static void dequantize_tiled_weight_to_fp16_task_iq2_xxs(
+ const tiled_dequantize_state_t *state,
+ uint32_t start_tile, uint32_t end_tile) {
+
+ HVX_Vector v_d = Q6_V_vzero();
+ const HVX_Vector v_eighth = Q6_Vh_vsplat_R(IQ2XXS_D_SCALE_F16);
+
+ HVX_Vector v_sgn = Q6_V_vzero();
+ if (start_tile < end_tile) {
+ v_sgn = iq2xxs_gather_issue(state->src + start_tile * state->aligned_tile_size);
+ }
+
+ for (uint32_t t = start_tile; t < end_tile; t++) {
+ const uint8_t * tile = state->src + t * state->aligned_tile_size;
+ __fp16 * dst_ptr = state->dst + t * HTP_MM_HMX_TILE_N_ELMS;
+
+ // rows hold all k tiles and n_k_tiles is a multiple of 8, so t % 8 is the tile within its super-block
+ if (t == start_tile || (t & (IQ2XXS_TILES_PER_SB - 1)) == 0) {
+ v_d = iq2xxs_gather_d(tile - (t & (IQ2XXS_TILES_PER_SB - 1)) * state->aligned_tile_size);
+ }
+
+ HVX_Vector v_c[2], v_w[2 * IQ2XXS_GROUPS], v_ls;
+ iq2xxs_gather_wait(tile, v_c);
+ const HVX_Vector v_sgn_cur = v_sgn;
+ if (t + 1 < end_tile) {
+ v_sgn = iq2xxs_gather_issue(tile + state->aligned_tile_size);
+ }
+ iq2xxs_decode_expand(v_c, v_sgn_cur, v_w, &v_ls);
+
+ HVX_Vector v_ls_h = Q6_Vhf_equals_Vh(Q6_Vh_vpacke_VwVw(v_ls, v_ls));
+ HVX_Vector v_sc = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(v_d, v_ls_h));
+ HVX_Vector v_scale = Q6_V_lo_W(Q6_W_vshuff_VVR(v_sc, v_sc, -2));
+
+ #pragma unroll
+ for (int g = 0; g < 2 * IQ2XXS_GROUPS; g++) {
+ HVX_VectorPair vp16 = Q6_Wh_vunpack_Vb(v_w[g]);
+ HVX_VectorPair vp_k = Q6_W_vdeal_VVR(Q6_V_hi_W(vp16), Q6_V_lo_W(vp16), -4);
+
+ // d / 8 is subnormal in fp16 for real weights, so apply the 1/8 after the multiply by d * ls
+ HVX_Vector v_lo = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_k)), v_scale));
+ HVX_Vector v_hi = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_k)), v_scale));
+ hvx_vmem(dst_ptr + (2 * g + 0) * 64) = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(v_lo, v_eighth));
+ hvx_vmem(dst_ptr + (2 * g + 1) * 64) = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(v_hi, v_eighth));
+ }
+ }
+}
+
// 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,
diff --git a/ggml/src/ggml-hexagon/htp/htp-ops.h b/ggml/src/ggml-hexagon/htp/htp-ops.h
index d0b53c015..4f24eda1e 100644
--- a/ggml/src/ggml-hexagon/htp/htp-ops.h
+++ b/ggml/src/ggml-hexagon/htp/htp-ops.h
@@ -29,11 +29,13 @@ enum htp_data_type {
HTP_TYPE_Q4_K = 12,
HTP_TYPE_Q5_K = 13,
HTP_TYPE_Q6_K = 14,
+ HTP_TYPE_IQ2_XXS = 16,
HTP_TYPE_IQ4_NL = 20,
HTP_TYPE_I32 = 26,
HTP_TYPE_I64 = 27,
HTP_TYPE_BF16 = 30,
HTP_TYPE_MXFP4 = 39,
+ HTP_TYPE_NVFP4 = 40,
HTP_TYPE_Q1_0 = 41,
HTP_TYPE_Q2_0 = 42,
diff --git a/ggml/src/ggml-hexagon/htp/htp-tensor.h b/ggml/src/ggml-hexagon/htp/htp-tensor.h
index beacc7e4d..2f9a4b129 100644
--- a/ggml/src/ggml-hexagon/htp/htp-tensor.h
+++ b/ggml/src/ggml-hexagon/htp/htp-tensor.h
@@ -22,6 +22,12 @@
#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_NVFP4 576
+#define HTP_TENSOR_TILE_SIZE_IQ2_XXS 264
+#define HTP_TENSOR_TILE_SCRATCH_IQ2_XXS 256
+#define HTP_IQ2XXS_NIBBLE_BYTES 128 // 16 words for vlut16, padded to a vector
+#define HTP_IQ2XXS_GRID_BYTES 512 // 256 entries x 8 elements x 2-bit magnitude code
+#define HTP_IQ2XXS_LUT_BYTES (HTP_IQ2XXS_NIBBLE_BYTES + HTP_IQ2XXS_GRID_BYTES)
#define HTP_TENSOR_TILE_SIZE_Q4_K 640
#define HTP_TENSOR_TILE_SIZE_Q5_K 768
#define HTP_TENSOR_TILE_SIZE_Q6_K 896
@@ -32,7 +38,7 @@
#define HTP_TENSOR_TILE_K_Q2_0 64
#define HTP_TENSOR_TILE_SIZE_Q2_0 576
-// Tiled weight tensor layouts (Q4_0, Q5_0, Q5_1, Q8_0, IQ4_NL, MXFP4, Q1_0, Q2_0, K-quants):
+// Tiled weight tensor layouts (Q4_0, Q5_0, Q5_1, Q8_0, IQ4_NL, MXFP4, NVFP4, IQ2_XXS, Q1_0, Q2_0, K-quants):
// See docs/backend/snapdragon/developer.md ("Tiled Weight Formats").
struct htp_tensor_mdev_range {
@@ -156,6 +162,8 @@ static inline uint32_t htp_tensor_get_tile_size(int type) {
case HTP_TYPE_Q1_0: return HTP_TENSOR_TILE_SIZE_Q1_0;
case HTP_TYPE_Q2_0: return HTP_TENSOR_TILE_SIZE_Q2_0;
case HTP_TYPE_MXFP4: return HTP_TENSOR_TILE_SIZE_MXFP4;
+ case HTP_TYPE_NVFP4: return HTP_TENSOR_TILE_SIZE_NVFP4;
+ case HTP_TYPE_IQ2_XXS: return HTP_TENSOR_TILE_SIZE_IQ2_XXS;
case HTP_TYPE_F16: return HTP_TENSOR_TILE_SIZE_F16;
default: return 0;
}
@@ -170,7 +178,9 @@ static inline uint32_t htp_tensor_get_tile_k(int type) {
}
static inline uint32_t htp_tensor_get_tile_padded_size(int type) {
- return (htp_tensor_get_tile_size(type) + 127u) & ~127u;
+ // IQ2_XXS: the last 256 B of the padded tile are the vgather scratch of its vec-dot
+ const uint32_t scratch = (type == HTP_TYPE_IQ2_XXS) ? HTP_TENSOR_TILE_SCRATCH_IQ2_XXS : 0u;
+ return (htp_tensor_get_tile_size(type) + scratch + 127u) & ~127u;
}
static inline bool htp_tensor_is_tiled(const struct htp_tensor * t) {
@@ -189,6 +199,8 @@ static inline uint32_t htp_tensor_get_row_size(int type, uint32_t ne00) {
case HTP_TYPE_Q5_0: return (ne00 / 32) * 22;
case HTP_TYPE_Q5_1: return (ne00 / 32) * 24;
case HTP_TYPE_MXFP4: return (ne00 / 32) * 17;
+ case HTP_TYPE_NVFP4: return (ne00 / 64) * 36;
+ case HTP_TYPE_IQ2_XXS: return (ne00 / 256) * 66;
case HTP_TYPE_Q2_K: return (ne00 / 256) * 84;
case HTP_TYPE_Q2_0: return (ne00 / 64) * 18;
case HTP_TYPE_Q3_K: return (ne00 / 256) * 110;
diff --git a/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h b/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h
index 4a6b2b5ea..e39ca336b 100644
--- a/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h
+++ b/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h
@@ -1850,6 +1850,212 @@ static void tiled_vec_dot_iq4nl_32x2(const uint32_t n, float * restrict s0, floa
}
}
+#define IQ2XXS_GROUPS 4 // 8-element grid entries per tile row
+#define IQ2XXS_TILES_PER_SB 8 // 256 / 32
+#define IQ2XXS_CODE_SHIFT 1 // log2(2 B per grid code)
+#define IQ2XXS_WORD_BYTES 4
+#define IQ2XXS_IDX_BITS 8
+#define IQ2XXS_IDX_MASK 0xFF
+#define IQ2XXS_SIGN_BITS 7
+#define IQ2XXS_SIGN_MASK 0x7F
+#define IQ2XXS_PARITY_SHIFT 7
+#define IQ2XXS_SCALE_SHIFT 28
+#define IQ2XXS_D_SLICE_OFF 256 // 4 fp16 d of the tile, after the 32 rows x 8 B
+#define IQ2XXS_D_SCALE 0.125f // db = d * (2 * scale + 1) / 8
+#define IQ2XXS_D_SLICE_BYTES 8
+#define IQ2XXS_SIGN_PRED_LO 0x08040201 // sign bit of bytes 0..3 of an 8-element group
+#define IQ2XXS_SIGN_PRED_HI 0x80402010 // sign bit of bytes 4..7
+#define IQ2XXS_IDX_LO_MASK 0x1FE // index byte << 1, halfword 0 of a lane
+#define IQ2XXS_IDX_HI_MASK 0x1FE0000 // index byte << 1, halfword 1 of a lane
+#define IQ2XXS_NIB_LO_MASK 0x00000F0F // nibbles of the code bytes, lane bytes 0..1
+#define IQ2XXS_NIB_HI_MASK 0x0F0F0000 // nibbles of the code bytes, lane bytes 2..3
+
+static inline HVX_Vector * iq2xxs_tile_scratch(const uint8_t * tile) {
+ return (HVX_Vector *) (tile + htp_tensor_get_tile_padded_size(HTP_TYPE_IQ2_XXS) - HTP_TENSOR_TILE_SCRATCH_IQ2_XXS);
+}
+
+// IQ2_XXS: start the 2 halfword grid-code gathers of a tile into its scratch, returns the sign/scale words
+static inline HVX_Vector iq2xxs_gather_issue(const uint8_t * restrict tile) {
+ HVX_Vector * restrict v_scr = iq2xxs_tile_scratch(tile);
+ const uint8_t * grid = iq2xxs_vtcm_lut + HTP_IQ2XXS_NIBBLE_BYTES;
+
+ // lo: 4 grid indices per row, hi: 4 sign indices + scale per row
+ HVX_VectorPair p = Q6_W_vdeal_VVR(hvx_vmem(tile + VLEN), hvx_vmem(tile), -IQ2XXS_WORD_BYTES);
+ HVX_Vector v_idx = Q6_V_lo_W(p);
+
+ // one halfword gather per 2 groups: lane i gets the codes of groups 2h and 2h+1 in halfwords 0 and 1
+ HVX_Vector v_off0 = Q6_V_vor_VV(Q6_V_vand_VV(Q6_Vw_vasl_VwR(v_idx, IQ2XXS_CODE_SHIFT), Q6_V_vsplat_R(IQ2XXS_IDX_LO_MASK)),
+ Q6_V_vand_VV(Q6_Vw_vasl_VwR(v_idx, IQ2XXS_CODE_SHIFT + IQ2XXS_IDX_BITS), Q6_V_vsplat_R(IQ2XXS_IDX_HI_MASK)));
+ HVX_Vector v_off1 = Q6_V_vor_VV(Q6_V_vand_VV(Q6_Vuw_vlsr_VuwR(v_idx, 2 * IQ2XXS_IDX_BITS - IQ2XXS_CODE_SHIFT), Q6_V_vsplat_R(IQ2XXS_IDX_LO_MASK)),
+ Q6_V_vand_VV(Q6_Vuw_vlsr_VuwR(v_idx, IQ2XXS_IDX_BITS - IQ2XXS_CODE_SHIFT), Q6_V_vsplat_R(IQ2XXS_IDX_HI_MASK)));
+ Q6_vgather_ARMVh(&v_scr[0], (size_t) grid, HTP_IQ2XXS_GRID_BYTES - 1, v_off0);
+ Q6_vgather_ARMVh(&v_scr[1], (size_t) grid, HTP_IQ2XXS_GRID_BYTES - 1, v_off1);
+
+ return Q6_V_hi_W(p);
+}
+
+// IQ2_XXS: wait for the gathers of iq2xxs_gather_issue and read the grid codes
+static inline void iq2xxs_gather_wait(const uint8_t * restrict tile, HVX_Vector * restrict v_c) {
+ HVX_Vector * restrict v_scr = iq2xxs_tile_scratch(tile);
+ hvx_gather_sync(&v_scr[1]);
+ v_c[0] = v_scr[0];
+ v_c[1] = v_scr[1];
+}
+
+// IQ2_XXS: 8 int8 weight words (k = 4g..4g+3 per row lane) and the odd group scale (2 * scale + 1) of a 32x32 tile
+static inline void iq2xxs_decode_expand(const HVX_Vector * restrict v_c, HVX_Vector v_sgn, HVX_Vector * restrict v_w, HVX_Vector * restrict v_ls) {
+ const HVX_Vector v_zero = Q6_V_vzero();
+ const HVX_Vector v_one = Q6_V_vsplat_R(1);
+ const HVX_Vector v_sgn_mask = Q6_V_vsplat_R(IQ2XXS_SIGN_MASK);
+ const HVX_Vector v_nib_lut = hvx_vmem(iq2xxs_vtcm_lut);
+
+ #pragma unroll
+ for (int l = 0; l < IQ2XXS_GROUPS; l++) {
+ // 7 stored sign bits + even-parity bit
+ HVX_Vector v_s7 = Q6_V_vand_VV(Q6_Vuw_vlsr_VuwR(v_sgn, IQ2XXS_SIGN_BITS * l), v_sgn_mask);
+ HVX_Vector v_t = Q6_Vh_vpopcount_Vh(v_s7);
+ HVX_Vector v_s = Q6_V_vor_VV(v_s7, Q6_Vw_vasl_VwR(Q6_V_vand_VV(v_t, v_one), IQ2XXS_PARITY_SHIFT));
+ HVX_Vector v_s16 = Q6_V_vor_VV(v_s, Q6_Vw_vasl_VwR(v_s, 8));
+ HVX_Vector v_rep = Q6_V_vor_VV(v_s16, Q6_Vw_vasl_VwR(v_s16, 16));
+ HVX_VectorPred q_lo = Q6_Q_vand_VR(v_rep, IQ2XXS_SIGN_PRED_LO);
+ HVX_VectorPred q_hi = Q6_Q_vand_VR(v_rep, IQ2XXS_SIGN_PRED_HI);
+
+ HVX_Vector v_c2 = v_c[l >> 1];
+ // code halfword duplicated in both halves of each lane
+ HVX_Vector v_d = (l & 1) ? Q6_Vh_vshuffo_VhVh(v_c2, v_c2) : Q6_Vh_vshuffe_VhVh(v_c2, v_c2);
+ // lane bytes: nibbles 0, 2, 1, 3 of the code, so vlut16 returns elements 0..3 / 4..7 as words
+ HVX_Vector v_n = Q6_V_vor_VV(Q6_V_vand_VV(v_d, Q6_V_vsplat_R(IQ2XXS_NIB_LO_MASK)),
+ Q6_V_vand_VV(Q6_Vuw_vlsr_VuwR(v_d, 4), Q6_V_vsplat_R(IQ2XXS_NIB_HI_MASK)));
+ HVX_VectorPair p_g = Q6_Wh_vlut16_VbVhR(v_n, v_nib_lut, 0);
+ HVX_Vector v_g0 = Q6_V_lo_W(p_g);
+ HVX_Vector v_g1 = Q6_V_hi_W(p_g);
+
+ v_w[2 * l + 0] = Q6_V_vmux_QVV(q_lo, Q6_Vb_vsub_VbVb(v_zero, v_g0), v_g0);
+ v_w[2 * l + 1] = Q6_V_vmux_QVV(q_hi, Q6_Vb_vsub_VbVb(v_zero, v_g1), v_g1);
+ }
+
+ *v_ls = Q6_Vw_vadd_VwVw(Q6_Vw_vasl_VwR(Q6_Vuw_vlsr_VuwR(v_sgn, IQ2XXS_SCALE_SHIFT), 1), v_one);
+}
+
+// IQ2_XXS: fp16 d of the 32 rows of a super-block (8 tiles) in halfword lanes 0..31
+static inline HVX_Vector iq2xxs_gather_d(const uint8_t * restrict sb) {
+ const uint32_t padded = htp_tensor_get_tile_padded_size(HTP_TYPE_IQ2_XXS);
+ HVX_Vector v_dd = Q6_V_vzero();
+
+ #pragma unroll
+ for (int j = 0; j < IQ2XXS_TILES_PER_SB; j++) {
+ HVX_Vector v = hvx_vmem(sb + j * padded + IQ2XXS_D_SLICE_OFF);
+ v_dd = Q6_V_vmux_QVV(Q6_Q_vsetq_R(IQ2XXS_D_SLICE_BYTES * j), v_dd,
+ Q6_V_vror_VR(v, (VLEN - IQ2XXS_D_SLICE_BYTES * j) & (VLEN - 1)));
+ }
+
+ return v_dd;
+}
+
+// IQ2_XXS: fp32 d / 8 of the 32 rows of a super-block (8 tiles) in lanes 0..31
+static inline HVX_Vector iq2xxs_load_d(const uint8_t * restrict sb) {
+ // db = d * (0.5 + scale) * 0.25 = d * (2 * scale + 1) / 8
+ return hvx_vec_mul_f32_f32(Q6_V_lo_W(hvx_vec_f16_to_f32(iq2xxs_gather_d(sb))), hvx_vec_splat_f32(IQ2XXS_D_SCALE));
+}
+
+static void tiled_vec_dot_iq2xxs_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 uint8_t * restrict tile_ptr = vx;
+ const uint8_t * restrict y_q = vy;
+ const uint32_t padded = htp_tensor_get_tile_padded_size(HTP_TYPE_IQ2_XXS);
+
+ HVX_Vector v_sum_float = Q6_V_vzero();
+
+ uint32_t n_k_tiles = n / 32;
+ HVX_Vector v_sgn = iq2xxs_gather_issue(tile_ptr);
+ for (uint32_t kb = 0; kb < n_k_tiles; kb += IQ2XXS_TILES_PER_SB) {
+ HVX_Vector v_d = iq2xxs_load_d(tile_ptr + kb * padded);
+
+ for (uint32_t kt = kb; kt < kb + IQ2XXS_TILES_PER_SB; kt++) {
+ const HVX_Vector * restrict v_act = (const HVX_Vector *) (y_q + kt * 1152);
+
+ HVX_Vector v_c[2], v_w[8], v_ls;
+ iq2xxs_gather_wait(tile_ptr + kt * padded, v_c);
+ HVX_Vector v_sgn_cur = v_sgn;
+ if (kt + 1 < n_k_tiles) {
+ v_sgn = iq2xxs_gather_issue(tile_ptr + (kt + 1) * padded);
+ }
+ iq2xxs_decode_expand(v_c, v_sgn_cur, v_w, &v_ls);
+
+ HVX_Vector v_sum = Q6_V_vzero();
+ #pragma unroll
+ for (int g = 0; g < 8; g++) {
+ v_sum = Q6_Vw_vrmpyacc_VwVbVb(v_sum, v_w[g], v_act[g]);
+ }
+
+ HVX_Vector v_da = hvx_vec_mul_f32_f32(v_d, Q6_V_lo_W(hvx_vec_f16_to_f32(v_act[8])));
+ HVX_Vector v_t = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(v_sum), Q6_Vsf_equals_Vw(v_ls));
+ v_sum_float = hvx_vec_add_f32_f32(v_sum_float, hvx_vec_mul_f32_f32(v_t, v_da));
+ }
+ }
+
+ if (sz) {
+ hvx_vec_store_u(s, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum_float, hvx_vmemu(sz)));
+ } else {
+ hvx_vec_store_u(s, valid_rows * sizeof(float), v_sum_float);
+ }
+}
+
+static void tiled_vec_dot_iq2xxs_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 uint8_t * restrict tile_ptr = vx;
+ const uint8_t * restrict y0_q = vy0;
+ const uint8_t * restrict y1_q = vy1;
+ const uint32_t padded = htp_tensor_get_tile_padded_size(HTP_TYPE_IQ2_XXS);
+
+ HVX_Vector v_sum_float_c0 = Q6_V_vzero();
+ HVX_Vector v_sum_float_c1 = Q6_V_vzero();
+
+ uint32_t n_k_tiles = n / 32;
+ HVX_Vector v_sgn = iq2xxs_gather_issue(tile_ptr);
+ for (uint32_t kb = 0; kb < n_k_tiles; kb += IQ2XXS_TILES_PER_SB) {
+ HVX_Vector v_d = iq2xxs_load_d(tile_ptr + kb * padded);
+
+ for (uint32_t kt = kb; kt < kb + IQ2XXS_TILES_PER_SB; kt++) {
+ const HVX_Vector * restrict v_act0 = (const HVX_Vector *) (y0_q + kt * 1152);
+ const HVX_Vector * restrict v_act1 = (const HVX_Vector *) (y1_q + kt * 1152);
+
+ HVX_Vector v_c[2], v_w[8], v_ls;
+ iq2xxs_gather_wait(tile_ptr + kt * padded, v_c);
+ HVX_Vector v_sgn_cur = v_sgn;
+ if (kt + 1 < n_k_tiles) {
+ v_sgn = iq2xxs_gather_issue(tile_ptr + (kt + 1) * padded);
+ }
+ iq2xxs_decode_expand(v_c, v_sgn_cur, v_w, &v_ls);
+
+ HVX_Vector v_sum0 = Q6_V_vzero();
+ HVX_Vector v_sum1 = Q6_V_vzero();
+ #pragma unroll
+ for (int g = 0; g < 8; g++) {
+ v_sum0 = Q6_Vw_vrmpyacc_VwVbVb(v_sum0, v_w[g], v_act0[g]);
+ v_sum1 = Q6_Vw_vrmpyacc_VwVbVb(v_sum1, v_w[g], v_act1[g]);
+ }
+
+ HVX_Vector v_ls_sf = Q6_Vsf_equals_Vw(v_ls);
+ HVX_Vector v_da0 = hvx_vec_mul_f32_f32(v_d, Q6_V_lo_W(hvx_vec_f16_to_f32(v_act0[8])));
+ HVX_Vector v_da1 = hvx_vec_mul_f32_f32(v_d, Q6_V_lo_W(hvx_vec_f16_to_f32(v_act1[8])));
+ HVX_Vector v_t0 = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(v_sum0), v_ls_sf);
+ HVX_Vector v_t1 = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(v_sum1), v_ls_sf);
+ v_sum_float_c0 = hvx_vec_add_f32_f32(v_sum_float_c0, hvx_vec_mul_f32_f32(v_t0, v_da0));
+ v_sum_float_c1 = hvx_vec_add_f32_f32(v_sum_float_c1, hvx_vec_mul_f32_f32(v_t1, v_da1));
+ }
+ }
+
+ if (sz0) {
+ hvx_vec_store_u(s0, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum_float_c0, hvx_vmemu(sz0)));
+ } else {
+ hvx_vec_store_u(s0, valid_rows * sizeof(float), v_sum_float_c0);
+ }
+ if (sz1) {
+ hvx_vec_store_u(s1, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum_float_c1, hvx_vmemu(sz1)));
+ } else {
+ hvx_vec_store_u(s1, valid_rows * sizeof(float), v_sum_float_c1);
+ }
+}
+
static void tiled_vec_dot_mxfp4_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 uint8_t * restrict tile_ptr = vx;
const uint8_t * restrict y_q = vy;
@@ -1956,6 +2162,163 @@ static void tiled_vec_dot_mxfp4_32x2(const uint32_t n, float * restrict s0, floa
}
}
+#define NVFP4_SCALE_OFF 512 // UE4M3 scales after the 32 rows x 16 B of quants
+#define NVFP4_UE4M3_MASK 0x7F // code bits; 0x7F is the NaN code, decoded as 0
+#define NVFP4_UE4M3_MIN_NORMAL 0x08 // exponent field 1
+#define NVFP4_UE4M3_MAN_SHIFT 20 // 23 - 3 mantissa bits
+#define NVFP4_UE4M3_REBIAS 0x3C000000 // (127 - 7) << 23
+#define NVFP4_UE4M3_SUB_SCALE 0x1p-9f // subnormal: mantissa * 2^-9
+#define NVFP4_VECS_PER_GROUP 2 // quant vectors per scale group
+#define NVFP4_ACT_PER_GROUP 4 // activation vectors per scale group
+
+// UE4M3 codes in the low byte of each word -> fp32 (without the 0.5 of kvalues_mxfp4)
+static inline HVX_Vector nvfp4_ue4m3_to_f32(HVX_Vector v_code) {
+ HVX_Vector v_mask = Q6_V_vsplat_R(NVFP4_UE4M3_MASK);
+ HVX_Vector v_x = Q6_V_vand_VV(v_code, v_mask);
+
+ HVX_Vector v_norm = Q6_Vw_vadd_VwVw(Q6_Vw_vasl_VwR(v_x, NVFP4_UE4M3_MAN_SHIFT), Q6_V_vsplat_R(NVFP4_UE4M3_REBIAS));
+ HVX_Vector v_sub = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(v_x), hvx_vec_splat_f32(NVFP4_UE4M3_SUB_SCALE));
+
+ HVX_VectorPred q_sub = Q6_Q_vcmp_gt_VwVw(Q6_V_vsplat_R(NVFP4_UE4M3_MIN_NORMAL), v_x);
+ HVX_VectorPred q_nan = Q6_Q_vcmp_eq_VwVw(v_x, v_mask);
+
+ return Q6_V_vmux_QVV(q_nan, Q6_V_vzero(), Q6_V_vmux_QVV(q_sub, v_sub, v_norm));
+}
+
+// fp32 scales of the two 16-element sub-blocks of a tile, one lane per row
+static inline void nvfp4_load_scales(const uint8_t * restrict tile, HVX_Vector * restrict v_d0, HVX_Vector * restrict v_d1) {
+ HVX_VectorPair v_code = Q6_Wuw_vunpack_Vuh(Q6_V_lo_W(Q6_Wuh_vunpack_Vub(hvx_vmem(tile + NVFP4_SCALE_OFF))));
+
+ *v_d0 = nvfp4_ue4m3_to_f32(Q6_V_lo_W(v_code));
+ *v_d1 = nvfp4_ue4m3_to_f32(Q6_V_hi_W(v_code));
+}
+
+// int8 dot of one 16-element scale group (2 quant vectors, 4 activation vectors)
+static inline HVX_Vector accum_nvfp4_group_32x1(const HVX_Vector * restrict vptr, const HVX_Vector * restrict v_act, HVX_Vector mask_h4, HVX_Vector lut) {
+ HVX_Vector v_sum0 = Q6_V_vzero();
+ HVX_Vector v_sum1 = Q6_V_vzero();
+
+ #pragma unroll
+ for (int i = 0; i < NVFP4_VECS_PER_GROUP; i++) {
+ HVX_VectorPair v_W_pair = unpack_and_interleave_4bit_x2(vptr[i], mask_h4);
+ HVX_Vector v_W0 = Q6_Vb_vlut32_VbVbI(Q6_V_lo_W(v_W_pair), lut, 0);
+ HVX_Vector v_W1 = Q6_Vb_vlut32_VbVbI(Q6_V_hi_W(v_W_pair), lut, 0);
+ v_sum0 = Q6_Vw_vrmpyacc_VwVbVb(v_sum0, v_W0, v_act[i * 2 + 0]);
+ v_sum1 = Q6_Vw_vrmpyacc_VwVbVb(v_sum1, v_W1, v_act[i * 2 + 1]);
+ }
+
+ return Q6_Vw_vadd_VwVw(v_sum0, v_sum1);
+}
+
+static inline HVX_VectorPair accum_nvfp4_group_32x2(const HVX_Vector * restrict vptr, const HVX_Vector * restrict v_act0, const HVX_Vector * restrict v_act1, HVX_Vector mask_h4, HVX_Vector lut) {
+ HVX_Vector v_sum0 = Q6_V_vzero();
+ HVX_Vector v_sum1 = Q6_V_vzero();
+
+ #pragma unroll
+ for (int i = 0; i < NVFP4_VECS_PER_GROUP; i++) {
+ HVX_VectorPair v_W_pair = unpack_and_interleave_4bit_x2(vptr[i], mask_h4);
+ HVX_Vector v_W0 = Q6_Vb_vlut32_VbVbI(Q6_V_lo_W(v_W_pair), lut, 0);
+ HVX_Vector v_W1 = Q6_Vb_vlut32_VbVbI(Q6_V_hi_W(v_W_pair), lut, 0);
+
+ v_sum0 = Q6_Vw_vrmpyacc_VwVbVb(v_sum0, v_W0, v_act0[i * 2 + 0]);
+ v_sum0 = Q6_Vw_vrmpyacc_VwVbVb(v_sum0, v_W1, v_act0[i * 2 + 1]);
+
+ v_sum1 = Q6_Vw_vrmpyacc_VwVbVb(v_sum1, v_W0, v_act1[i * 2 + 0]);
+ v_sum1 = Q6_Vw_vrmpyacc_VwVbVb(v_sum1, v_W1, v_act1[i * 2 + 1]);
+ }
+
+ return Q6_W_vcombine_VV(v_sum1, v_sum0);
+}
+
+static void tiled_vec_dot_nvfp4_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 uint8_t * restrict tile_ptr = vx;
+ const uint8_t * restrict y_q = vy;
+ const uint32_t padded = htp_tensor_get_tile_padded_size(HTP_TYPE_NVFP4);
+
+ HVX_Vector v_sum_float = Q6_V_vzero();
+ HVX_Vector mask_h4 = Q6_Vb_vsplat_R(0x0F);
+ HVX_Vector lut = *(const HVX_Vector *) kvalues_mxfp4_lut;
+
+ uint32_t n_k_tiles = n / 32;
+ for (uint32_t kt = 0; kt < n_k_tiles; kt++) {
+ const uint8_t * restrict tile = tile_ptr + kt * padded;
+ const HVX_Vector * restrict vptr = (const HVX_Vector *) tile;
+ const HVX_Vector * restrict v_act = (const HVX_Vector *) (y_q + kt * 1152);
+
+ HVX_Vector v_sum0 = accum_nvfp4_group_32x1(vptr, v_act, mask_h4, lut);
+ HVX_Vector v_sum1 = accum_nvfp4_group_32x1(vptr + NVFP4_VECS_PER_GROUP, v_act + NVFP4_ACT_PER_GROUP, mask_h4, lut);
+
+ HVX_Vector v_d0, v_d1;
+ nvfp4_load_scales(tile, &v_d0, &v_d1);
+
+ HVX_Vector v_scale_a = Q6_V_lo_W(hvx_vec_f16_to_f32_shuff(v_act[8]));
+
+ HVX_Vector v_t = hvx_vec_add_f32_f32(hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(v_sum0), v_d0),
+ hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(v_sum1), v_d1));
+
+ v_sum_float = hvx_vec_add_f32_f32(v_sum_float, hvx_vec_mul_f32_f32(v_t, v_scale_a));
+ }
+
+ v_sum_float = hvx_vec_mul_f32_f32(v_sum_float, hvx_vec_splat_f32(0.5f));
+
+ if (sz) {
+ hvx_vec_store_u(s, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum_float, hvx_vmemu(sz)));
+ } else {
+ hvx_vec_store_u(s, valid_rows * sizeof(float), v_sum_float);
+ }
+}
+
+static void tiled_vec_dot_nvfp4_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 uint8_t * restrict tile_ptr = vx;
+ const uint8_t * restrict y0_q = vy0;
+ const uint8_t * restrict y1_q = vy1;
+ const uint32_t padded = htp_tensor_get_tile_padded_size(HTP_TYPE_NVFP4);
+
+ HVX_Vector v_sum_float_c0 = Q6_V_vzero();
+ HVX_Vector v_sum_float_c1 = Q6_V_vzero();
+ HVX_Vector mask_h4 = Q6_Vb_vsplat_R(0x0F);
+ HVX_Vector lut = *(const HVX_Vector *) kvalues_mxfp4_lut;
+
+ uint32_t n_k_tiles = n / 32;
+ for (uint32_t kt = 0; kt < n_k_tiles; kt++) {
+ const uint8_t * restrict tile = tile_ptr + kt * padded;
+ const HVX_Vector * restrict vptr = (const HVX_Vector *) tile;
+ const HVX_Vector * restrict v_act0 = (const HVX_Vector *) (y0_q + kt * 1152);
+ const HVX_Vector * restrict v_act1 = (const HVX_Vector *) (y1_q + kt * 1152);
+
+ HVX_Vector v_d0, v_d1;
+ nvfp4_load_scales(tile, &v_d0, &v_d1);
+
+ HVX_VectorPair v_g0 = accum_nvfp4_group_32x2(vptr, v_act0, v_act1, mask_h4, lut);
+ HVX_Vector v_t_c0 = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(v_g0)), v_d0);
+ HVX_Vector v_t_c1 = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_hi_W(v_g0)), v_d0);
+
+ HVX_VectorPair v_g1 = accum_nvfp4_group_32x2(vptr + NVFP4_VECS_PER_GROUP, v_act0 + NVFP4_ACT_PER_GROUP, v_act1 + NVFP4_ACT_PER_GROUP, mask_h4, lut);
+ v_t_c0 = hvx_vec_add_f32_f32(v_t_c0, hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(v_g1)), v_d1));
+ v_t_c1 = hvx_vec_add_f32_f32(v_t_c1, hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_hi_W(v_g1)), v_d1));
+
+ HVX_Vector v_scale_a_c0 = Q6_V_lo_W(hvx_vec_f16_to_f32_shuff(v_act0[8]));
+ HVX_Vector v_scale_a_c1 = Q6_V_lo_W(hvx_vec_f16_to_f32_shuff(v_act1[8]));
+
+ v_sum_float_c0 = hvx_vec_add_f32_f32(v_sum_float_c0, hvx_vec_mul_f32_f32(v_t_c0, v_scale_a_c0));
+ v_sum_float_c1 = hvx_vec_add_f32_f32(v_sum_float_c1, hvx_vec_mul_f32_f32(v_t_c1, v_scale_a_c1));
+ }
+
+ v_sum_float_c0 = hvx_vec_mul_f32_f32(v_sum_float_c0, hvx_vec_splat_f32(0.5f));
+ v_sum_float_c1 = hvx_vec_mul_f32_f32(v_sum_float_c1, hvx_vec_splat_f32(0.5f));
+
+ if (sz0) {
+ hvx_vec_store_u(s0, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum_float_c0, hvx_vmemu(sz0)));
+ } else {
+ hvx_vec_store_u(s0, valid_rows * sizeof(float), v_sum_float_c0);
+ }
+ if (sz1) {
+ hvx_vec_store_u(s1, valid_rows * sizeof(float), hvx_vec_add_f32_f32(v_sum_float_c1, hvx_vmemu(sz1)));
+ } else {
+ hvx_vec_store_u(s1, valid_rows * sizeof(float), v_sum_float_c1);
+ }
+}
+
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;
diff --git a/ggml/src/ggml-hexagon/htp/iq2xxs-tables.h b/ggml/src/ggml-hexagon/htp/iq2xxs-tables.h
new file mode 100644
index 000000000..5a3dbe26e
--- /dev/null
+++ b/ggml/src/ggml-hexagon/htp/iq2xxs-tables.h
@@ -0,0 +1,54 @@
+#ifndef HTP_IQ2XXS_TABLES_H
+#define HTP_IQ2XXS_TABLES_H
+
+// IQ2_XXS lookup tables, copied to VTCM per op (see ggml-common.h: iq2xxs_grid)
+// nibble table: 16 words (low halfword used by vlut16), 2 grid magnitudes (2-bit codes -> 8, 25, 43) per 4-bit index, then
+// 256 halfword grid codes (2 bits per element, element 0 in the low bits)
+
+#include <stdint.h>
+#include "htp-tensor.h"
+
+static const uint8_t __attribute__((aligned(128))) htp_iq2xxs_lut[HTP_IQ2XXS_LUT_BYTES] = {
+ 0x08, 0x08, 0x00, 0x00, 0x19, 0x08, 0x00, 0x00, 0x2b, 0x08, 0x00, 0x00, 0x2b, 0x08, 0x00, 0x00,
+ 0x08, 0x19, 0x00, 0x00, 0x19, 0x19, 0x00, 0x00, 0x2b, 0x19, 0x00, 0x00, 0x2b, 0x19, 0x00, 0x00,
+ 0x08, 0x2b, 0x00, 0x00, 0x19, 0x2b, 0x00, 0x00, 0x2b, 0x2b, 0x00, 0x00, 0x2b, 0x2b, 0x00, 0x00,
+ 0x08, 0x2b, 0x00, 0x00, 0x19, 0x2b, 0x00, 0x00, 0x2b, 0x2b, 0x00, 0x00, 0x2b, 0x2b, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
+ 0x00, 0x00, 0x02, 0x00, 0x05, 0x00, 0x08, 0x00, 0x0a, 0x00, 0x11, 0x00, 0x14, 0x00, 0x20, 0x00,
+ 0x22, 0x00, 0x28, 0x00, 0x2a, 0x00, 0x41, 0x00, 0x44, 0x00, 0x50, 0x00, 0x58, 0x00, 0x61, 0x00,
+ 0x64, 0x00, 0x80, 0x00, 0x82, 0x00, 0x8a, 0x00, 0xa2, 0x00, 0x01, 0x01, 0x04, 0x01, 0x10, 0x01,
+ 0x15, 0x01, 0x40, 0x01, 0x84, 0x01, 0x98, 0x01, 0x00, 0x02, 0x02, 0x02, 0x22, 0x02, 0x82, 0x02,
+ 0x01, 0x04, 0x04, 0x04, 0x10, 0x04, 0x21, 0x04, 0x24, 0x04, 0x40, 0x04, 0x42, 0x04, 0x48, 0x04,
+ 0x60, 0x04, 0x81, 0x04, 0x84, 0x04, 0x90, 0x04, 0xa4, 0x04, 0x00, 0x05, 0x02, 0x05, 0x08, 0x05,
+ 0x20, 0x05, 0x46, 0x05, 0x69, 0x05, 0x80, 0x05, 0x91, 0x05, 0x09, 0x06, 0x10, 0x06, 0x40, 0x06,
+ 0x84, 0x06, 0xa4, 0x06, 0x00, 0x08, 0x05, 0x08, 0x08, 0x08, 0x14, 0x08, 0x28, 0x08, 0x41, 0x08,
+ 0x44, 0x08, 0x50, 0x08, 0x52, 0x08, 0x88, 0x08, 0x04, 0x09, 0x40, 0x09, 0x02, 0x0a, 0x14, 0x0a,
+ 0x01, 0x10, 0x04, 0x10, 0x10, 0x10, 0x21, 0x10, 0x40, 0x10, 0x60, 0x10, 0x84, 0x10, 0x90, 0x10,
+ 0x95, 0x10, 0x00, 0x11, 0x08, 0x11, 0x20, 0x11, 0x50, 0x11, 0x5a, 0x11, 0x80, 0x11, 0x24, 0x12,
+ 0x45, 0x12, 0x00, 0x14, 0x08, 0x14, 0x20, 0x14, 0x25, 0x14, 0x49, 0x14, 0x80, 0x14, 0x18, 0x15,
+ 0x62, 0x15, 0x00, 0x16, 0x16, 0x16, 0x01, 0x18, 0x04, 0x18, 0x10, 0x18, 0x40, 0x18, 0x81, 0x18,
+ 0x00, 0x19, 0x05, 0x19, 0xa0, 0x19, 0x51, 0x1a, 0x00, 0x20, 0x02, 0x20, 0x0a, 0x20, 0x44, 0x20,
+ 0x61, 0x20, 0x80, 0x20, 0x82, 0x20, 0x29, 0x21, 0x48, 0x21, 0x00, 0x22, 0x02, 0x22, 0x01, 0x24,
+ 0x04, 0x24, 0x10, 0x24, 0x40, 0x24, 0x56, 0x24, 0x00, 0x25, 0x41, 0x25, 0x64, 0x25, 0x90, 0x26,
+ 0x08, 0x28, 0x20, 0x28, 0x94, 0x28, 0x44, 0x2a, 0x01, 0x40, 0x04, 0x40, 0x10, 0x40, 0x18, 0x40,
+ 0x21, 0x40, 0x24, 0x40, 0x40, 0x40, 0x48, 0x40, 0x56, 0x40, 0x60, 0x40, 0x81, 0x40, 0x84, 0x40,
+ 0x90, 0x40, 0x00, 0x41, 0x20, 0x41, 0x61, 0x41, 0x80, 0x41, 0x85, 0x41, 0x01, 0x42, 0x10, 0x42,
+ 0x48, 0x42, 0x56, 0x42, 0x68, 0x42, 0x00, 0x44, 0x08, 0x44, 0x20, 0x44, 0x80, 0x44, 0x99, 0x44,
+ 0x12, 0x45, 0x24, 0x45, 0x00, 0x46, 0x01, 0x48, 0x04, 0x48, 0x10, 0x48, 0x40, 0x48, 0x45, 0x48,
+ 0x00, 0x49, 0x58, 0x49, 0x61, 0x49, 0x82, 0x49, 0x45, 0x4a, 0x90, 0x4a, 0x00, 0x50, 0x08, 0x50,
+ 0x11, 0x50, 0x19, 0x50, 0x20, 0x50, 0x80, 0x50, 0x88, 0x50, 0x04, 0x51, 0x42, 0x51, 0xa4, 0x51,
+ 0x91, 0x52, 0x90, 0x54, 0x92, 0x54, 0x0a, 0x55, 0x01, 0x56, 0x54, 0x56, 0x00, 0x58, 0x11, 0x58,
+ 0x19, 0x58, 0x64, 0x58, 0x40, 0x59, 0x08, 0x5a, 0x04, 0x60, 0x10, 0x60, 0x40, 0x60, 0x68, 0x60,
+ 0x00, 0x61, 0x55, 0x61, 0x18, 0x62, 0x60, 0x62, 0x00, 0x64, 0x05, 0x64, 0x10, 0x65, 0x12, 0x65,
+ 0x84, 0x65, 0x42, 0x68, 0x00, 0x80, 0x02, 0x80, 0x0a, 0x80, 0x41, 0x80, 0x82, 0x80, 0x04, 0x81,
+ 0x18, 0x81, 0x40, 0x81, 0x11, 0x82, 0x01, 0x84, 0x04, 0x84, 0x10, 0x84, 0x15, 0x84, 0x40, 0x84,
+ 0x60, 0x84, 0x00, 0x85, 0x46, 0x85, 0x94, 0x85, 0x09, 0x86, 0x40, 0x86, 0x60, 0x86, 0x02, 0x88,
+ 0x04, 0x89, 0x11, 0x8a, 0x04, 0x90, 0x10, 0x90, 0x24, 0x90, 0x40, 0x90, 0xa1, 0x90, 0x16, 0x91,
+ 0x80, 0x91, 0x45, 0x92, 0x00, 0x94, 0x22, 0x94, 0x44, 0x94, 0x51, 0x95, 0x81, 0x98, 0x20, 0x99,
+ 0x02, 0xa0, 0x50, 0xa0, 0x85, 0xa0, 0x09, 0xa1, 0x00, 0xa2, 0x18, 0xa4, 0x50, 0xa8, 0x04, 0xa9,
+};
+
+#endif // HTP_IQ2XXS_TABLES_H
diff --git a/ggml/src/ggml-hexagon/htp/matmul-ops.c b/ggml/src/ggml-hexagon/htp/matmul-ops.c
index b2dbe861e..ea05937a5 100644
--- a/ggml/src/ggml-hexagon/htp/matmul-ops.c
+++ b/ggml/src/ggml-hexagon/htp/matmul-ops.c
@@ -191,6 +191,25 @@ static const uint8_t __attribute__((aligned(VLEN))) kvalues_mxfp4_lut[] = {
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
};
+#include "iq2xxs-tables.h"
+
+// VTCM copy of htp_iq2xxs_lut for the current op (set by hvx_mm_iq2xxs_lut_init)
+static const uint8_t * iq2xxs_vtcm_lut;
+
+static inline void iq2xxs_lut_copy(uint8_t * restrict lut) {
+ for (uint32_t i = 0; i < HTP_IQ2XXS_LUT_BYTES / VLEN; i++) {
+ hvx_vmem(lut + i * VLEN) = hvx_vmem(htp_iq2xxs_lut + i * VLEN);
+ }
+ iq2xxs_vtcm_lut = lut;
+}
+
+static inline void hvx_mm_iq2xxs_lut_init(uint8_t * restrict base, const struct htp_mm_hvx_vtcm_layout * restrict L) {
+ if (L->lut_bytes == 0) {
+ return;
+ }
+ iq2xxs_lut_copy(base + L->off_lut);
+}
+
#define htp_matmul_tensors_preamble \
const struct htp_tensor * restrict src0 = octx->src[0]; \
const struct htp_tensor * restrict src1 = octx->src[1]; \
@@ -1700,11 +1719,21 @@ static int hvx_mm_init_tiled_vec_dot(struct htp_mm_context * mmctx, enum htp_dat
mmctx->vec_dot_32x1 = tiled_vec_dot_iq4nl_32x1;
mmctx->vec_dot_32x2 = tiled_vec_dot_iq4nl_32x2;
return 0;
+ case HTP_TYPE_IQ2_XXS:
+ mmctx->type = "iq2xxs_tiled-f32";
+ mmctx->vec_dot_32x1 = tiled_vec_dot_iq2xxs_32x1;
+ mmctx->vec_dot_32x2 = tiled_vec_dot_iq2xxs_32x2;
+ return 0;
case HTP_TYPE_MXFP4:
mmctx->type = "mxfp4_tiled-f32";
mmctx->vec_dot_32x1 = tiled_vec_dot_mxfp4_32x1;
mmctx->vec_dot_32x2 = tiled_vec_dot_mxfp4_32x2;
return 0;
+ case HTP_TYPE_NVFP4:
+ mmctx->type = "nvfp4_tiled-f32";
+ mmctx->vec_dot_32x1 = tiled_vec_dot_nvfp4_32x1;
+ mmctx->vec_dot_32x2 = tiled_vec_dot_nvfp4_32x2;
+ return 0;
case HTP_TYPE_Q1_0:
mmctx->type = "q1_0_tiled-f32";
mmctx->vec_dot_32x1 = tiled_vec_dot_q1_0_32x1;
@@ -1889,6 +1918,7 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) {
mmctx->vtcm_bias = VTCM_LAYOUT_PTR(uint8_t, base, L.off_bias);
mmctx->vtcm_dst = VTCM_LAYOUT_PTR(uint8_t, base, L.off_dst);
mmctx->vtcm_act_raw = VTCM_LAYOUT_PTR(uint8_t, base, L.off_act_raw);
+ hvx_mm_iq2xxs_lut_init(base, &L);
mmctx->vtcm_src0_stride = src0_row_size_padded;
mmctx->vtcm_act_stride = act_row_size;
@@ -2321,6 +2351,8 @@ static inline tiled_dequant_task_fn_t hmx_get_dequant_worker_fn(int weight_type,
case HTP_TYPE_Q4_1:
case HTP_TYPE_Q4_K: return dequantize_tiled_weight_to_fp16_task_q4_1;
case HTP_TYPE_MXFP4: return dequantize_tiled_weight_to_fp16_task_mxfp4;
+ case HTP_TYPE_NVFP4: return dequantize_tiled_weight_to_fp16_task_nvfp4;
+ case HTP_TYPE_IQ2_XXS: return dequantize_tiled_weight_to_fp16_task_iq2_xxs;
case HTP_TYPE_Q8_0: return dequantize_tiled_weight_to_fp16_task_q8_0;
case HTP_TYPE_Q5_0: return dequantize_tiled_weight_to_fp16_task_q5_0;
case HTP_TYPE_Q5_1:
@@ -2688,6 +2720,9 @@ static int hmx_mm_2d_f32(struct htp_context *ctx,
__fp16 *vtcm_scales = VTCM_LAYOUT_PTR(__fp16, base, L.off_scales);
hmx_init_column_scales(vtcm_scales, Q6_V_vsplat_R(0x3c00)); // scale: 1.0, bias: 0.0 in FP16
+ if (L.lut_bytes > 0) {
+ iq2xxs_lut_copy(base + L.off_lut);
+ }
const bool has_bias = (src2_bytes > 0 && src2_addr != 0);
float *vtcm_bias = VTCM_LAYOUT_PTR_OPTIONAL(float, base, L.off_bias, has_bias);
@@ -2957,6 +2992,9 @@ static int hmx_mm_nx_2d_f32(struct htp_ops_context * octx, const struct htp_mm_k
__fp16 *vtcm_scales = VTCM_LAYOUT_PTR(__fp16, base, L.off_scales);
hmx_init_column_scales(vtcm_scales, Q6_V_vsplat_R(0x3c00)); // scale: 1.0, bias: 0.0 in FP16
+ if (L.lut_bytes > 0) {
+ iq2xxs_lut_copy(base + L.off_lut);
+ }
int m_start = 0;
int m_rows = m;
@@ -3539,7 +3577,7 @@ static int hmx_mm_id_2d_f32(struct htp_context *ctx,
htp_mm_hmx_get_2d_chunk_costs(weight_type, k, /*pipeline=*/false, aligned_tile_size,
&size_per_n, &size_per_m, &size_per_mn);
- const size_t overhead = htp_mm_hmx_get_2d_overhead(/*pipeline=*/false, /*is_matmul_id=*/true);
+ const size_t overhead = htp_mm_hmx_get_2d_overhead(weight_type, /*pipeline=*/false, /*is_matmul_id=*/true);
size_t m_chunk_n_rows = 0, n_chunk_n_cols = 0;
if (htp_mm_hmx_compute_chunks(vtcm_budget, overhead, size_per_n, size_per_m, size_per_mn,
m_core_padded, n,
@@ -3561,6 +3599,8 @@ static int hmx_mm_id_2d_f32(struct htp_context *ctx,
__fp16 *vtcm_output = (__fp16 *) vtcm_seq_alloc(&vtcm_ptr, output_area_size);
void *vtcm_scratch0 = vtcm_seq_alloc(&vtcm_ptr, scratch0_size);
__fp16 *vtcm_scales = (__fp16 *) vtcm_seq_alloc(&vtcm_ptr, 256);
+ const size_t lut_size = htp_mm_hmx_lut_bytes(weight_type);
+ uint8_t *vtcm_lut = lut_size ? (uint8_t *) vtcm_seq_alloc(&vtcm_ptr, lut_size) : NULL;
vtcm_used = vtcm_ptr - (uint8_t *) ctx->vtcm_base;
if (vtcm_used > vtcm_budget) {
@@ -3569,6 +3609,9 @@ static int hmx_mm_id_2d_f32(struct htp_context *ctx,
}
hmx_init_column_scales(vtcm_scales, Q6_V_vsplat_R(0x3c00));
+ if (lut_size) {
+ iq2xxs_lut_copy(vtcm_lut);
+ }
htp_trace_event_stop(tr, HTP_TRACE_EVT_INIT, 0);
@@ -3894,6 +3937,7 @@ static int hvx_mm_matmul_id(
mmctx->vtcm_bias = NULL;
mmctx->vtcm_dst = VTCM_LAYOUT_PTR(uint8_t, base, L.off_dst);
mmctx->vtcm_act_raw = VTCM_LAYOUT_PTR(uint8_t, base, L.off_act_raw);
+ hvx_mm_iq2xxs_lut_init(base, &L);
mmctx->vtcm_src0_stride = src0_row_size_padded;
mmctx->vtcm_act_stride = act_row_size;
@@ -4052,6 +4096,7 @@ static int hvx_mm_matmul_id_nx(
mmctx->vtcm_act = VTCM_LAYOUT_PTR(uint8_t, base, L.off_act);
mmctx->vtcm_dst = VTCM_LAYOUT_PTR(uint8_t, base, L.off_dst);
mmctx->vtcm_act_raw = VTCM_LAYOUT_PTR(uint8_t, base, L.off_act_raw);
+ hvx_mm_iq2xxs_lut_init(base, &L);
mmctx->vtcm_src0_stride = 0;
mmctx->vtcm_act_stride = act_row_size;
@@ -4537,6 +4582,7 @@ static int op_matmul_nx_impl(struct htp_ops_context * octx) {
mmctx->vtcm_act = VTCM_LAYOUT_PTR(uint8_t, base, L.off_act);
mmctx->vtcm_dst = VTCM_LAYOUT_PTR(uint8_t, base, L.off_dst);
mmctx->vtcm_act_raw = VTCM_LAYOUT_PTR(uint8_t, base, L.off_act_raw);
+ hvx_mm_iq2xxs_lut_init(base, &L);
mmctx->vtcm_src0_stride = 0;
mmctx->vtcm_act_stride = act_row_size;
diff --git a/ggml/src/ggml-hexagon/htp/matmul-ops.h b/ggml/src/ggml-hexagon/htp/matmul-ops.h
index 7a85624cf..94441cf8c 100644
--- a/ggml/src/ggml-hexagon/htp/matmul-ops.h
+++ b/ggml/src/ggml-hexagon/htp/matmul-ops.h
@@ -252,9 +252,13 @@ static inline void htp_mm_hmx_get_batched_chunk_costs(
*size_per_mn_out = sizeof(uint16_t);
}
-static inline size_t htp_mm_hmx_get_2d_overhead(bool pipeline, bool is_matmul_id) {
+static inline size_t htp_mm_hmx_lut_bytes(int wtype) {
+ return (wtype == HTP_TYPE_IQ2_XXS) ? hex_align_up(HTP_IQ2XXS_LUT_BYTES, HTP_MM_HMX_TILE_SIZE) : 0;
+}
+
+static inline size_t htp_mm_hmx_get_2d_overhead(int wtype, bool pipeline, bool is_matmul_id) {
size_t num_regions = pipeline ? 7 : (is_matmul_id ? 4 : 5);
- return num_regions * HTP_MM_HMX_TILE_SIZE + 256;
+ return num_regions * HTP_MM_HMX_TILE_SIZE + 256 + htp_mm_hmx_lut_bytes(wtype);
}
static inline size_t htp_mm_hmx_get_batched_overhead(void) {
@@ -270,6 +274,7 @@ struct htp_mm_hmx_vtcm_layout {
size_t off_scratch[2]; // dequantization scratch pads
size_t off_scales; // HMX scales (256 bytes)
size_t off_bias; // bias in VTCM
+ size_t off_lut; // IQ2_XXS grid + sign tables (2D only)
// Cached sizes of regions for HMX kernel use
size_t weight_area_bytes;
@@ -279,6 +284,7 @@ struct htp_mm_hmx_vtcm_layout {
size_t scratch_bytes[2];
size_t act_head_stride;
size_t bias_bytes;
+ size_t lut_bytes;
size_t total_bytes;
};
@@ -290,6 +296,7 @@ struct htp_mm_hvx_vtcm_layout {
size_t off_bias; // vtcm_bias (bias / fused add only)
size_t off_dst; // vtcm_dst (output scratch)
size_t off_act_raw; // vtcm_act_raw (raw activation DMA staging)
+ size_t off_lut; // vtcm_lut (IQ2_XXS grid + sign tables)
// Cached sizes
size_t src0_bytes;
@@ -297,6 +304,7 @@ struct htp_mm_hvx_vtcm_layout {
size_t bias_bytes;
size_t dst_bytes;
size_t act_raw_bytes;
+ size_t lut_bytes;
size_t total_bytes;
};
@@ -315,6 +323,8 @@ static inline void htp_mm_hmx_vtcm_layout_build(
size_t bias_size
) {
size_t off = 0;
+ L->lut_bytes = 0;
+ L->off_lut = 0;
if (kernel_type == HTP_MM_KERNEL_HMX_BATCHED) {
const size_t vec_dot_size = k * sizeof(uint16_t);
@@ -412,6 +422,9 @@ static inline void htp_mm_hmx_vtcm_layout_build(
L->bias_bytes = bias_size;
off = off_group_a + hex_smax(group_b_size, group_c_size);
+
+ L->lut_bytes = htp_mm_hmx_lut_bytes(wtype);
+ VTCM_LAYOUT_ALLOC_OPTIONAL(off, off_lut, L->lut_bytes, L->lut_bytes > 0);
}
L->total_bytes = off;
@@ -438,6 +451,7 @@ static inline void htp_mm_hvx_vtcm_layout_build(
size_t bias_sz = bias_row_size > 0 ? htp_mm_round_up(bias_row_size, 128) : 0;
size_t dst_sz = 0;
size_t act_raw_sz = 0;
+ const size_t lut_sz = (wtype == HTP_TYPE_IQ2_XXS) ? HTP_IQ2XXS_LUT_BYTES : 0;
if (is_fused_nx) {
uint32_t aligned_tile_size = htp_tensor_get_tile_padded_size(wtype);
@@ -517,6 +531,7 @@ static inline void htp_mm_hvx_vtcm_layout_build(
size_t off_group_a = 0;
VTCM_LAYOUT_ALLOC(off_group_a, off_act, act_sz);
VTCM_LAYOUT_ALLOC_OPTIONAL(off_group_a, off_bias, bias_sz, bias_sz > 0);
+ VTCM_LAYOUT_ALLOC_OPTIONAL(off_group_a, off_lut, lut_sz, lut_sz > 0);
// Group B: Compute-only buffers (starts at off_group_a)
size_t off_group_b = off_group_a;
@@ -534,6 +549,7 @@ static inline void htp_mm_hvx_vtcm_layout_build(
L->bias_bytes = bias_sz;
L->dst_bytes = dst_sz;
L->act_raw_bytes = act_raw_sz;
+ L->lut_bytes = lut_sz;
L->total_bytes = off_group_a + hex_smax(group_b_size, group_c_size);
}
@@ -564,7 +580,7 @@ static inline bool htp_mm_hvx_solve_vtcm_params(
return true;
}
- const size_t fixed_bytes = L.src0_bytes + L.bias_bytes + L.dst_bytes;
+ const size_t fixed_bytes = L.src0_bytes + L.bias_bytes + L.dst_bytes + L.lut_bytes;
if (vtcm_budget <= fixed_bytes) {
return false;
}
@@ -727,7 +743,7 @@ static inline bool htp_mm_hmx_solve_2d_params(
int act_hvx = n_hvx;
while (act_hvx >= 1) {
- size_t simple_2d_overhead = htp_mm_hmx_get_2d_overhead(pipeline, is_matmul_id) + (bias_size > 0 ? hex_align_up(bias_size, HTP_MM_HMX_TILE_SIZE) : 0);
+ size_t simple_2d_overhead = htp_mm_hmx_get_2d_overhead(wtype, pipeline, is_matmul_id) + (bias_size > 0 ? hex_align_up(bias_size, HTP_MM_HMX_TILE_SIZE) : 0);
size_t simple_2d_size_per_n, simple_2d_size_per_m, simple_2d_size_per_mn;
htp_mm_hmx_get_2d_chunk_costs(wtype, k, pipeline, aligned_tile_size, &simple_2d_size_per_n, &simple_2d_size_per_m, &simple_2d_size_per_mn);
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index d0a585a22..c328072bc 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -12122,6 +12122,13 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
GGML_TYPE_IQ3_S, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M, GGML_TYPE_IQ4_XS}) {
test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 16, 10, 256, {1, 1}, {1, 1}));
}
+
+ // Tile-aligned shapes for NVFP4 and IQ2_XXS across GEMV (n=1, 2) and HMX (n=16, 64) paths
+ for (ggml_type type_a : {GGML_TYPE_NVFP4, GGML_TYPE_IQ2_XXS}) {
+ for (int n : {1, 2, 16, 64}) {
+ test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 64, n, 256, {1, 1}, {1, 1}));
+ }
+ }
#else
// m = a rows
// n = b rows