Commit 2a4daf0e0 for llama.cpp
commit 2a4daf0e08d6a9b6993ee0d4a9a3cfc35a6f405c
Author: Aparna M P <aparmp@qti.qualcomm.com>
Date: Thu Nov 12 22:03:16 2026 +0530
hexagon: add Q5_0 and Q5_1 support (#32094)
* hexagon: add q5_0 and q5_1 support
* hex-mm: dedicated q5_0 kernels
diff --git a/docs/backend/snapdragon/developer.md b/docs/backend/snapdragon/developer.md
index 95b6127a4..69a8d5f7d 100644
--- a/docs/backend/snapdragon/developer.md
+++ b/docs/backend/snapdragon/developer.md
@@ -60,7 +60,7 @@ Implementation: [`htp-repack.cpp`](../../../ggml/src/ggml-hexagon/htp-repack.cpp
### Tile Dimensions and Grid Layout
Each tile represents a submatrix of `N_rows = 32` rows and `K_cols` columns:
-- Standard types (`Q4_0`, `Q4_1`, `Q8_0`, `IQ4_NL`, `MXFP4`, and K-quants): `K_cols = 32` (`HTP_TENSOR_TILE_N_COLS = 32`).
+- Standard types (`Q4_0`, `Q4_1`, `Q5_0`, `Q5_1`, `Q8_0`, `IQ4_NL`, `MXFP4`, and K-quants): `K_cols = 32` (`HTP_TENSOR_TILE_N_COLS = 32`).
- `Q2_0`: `K_cols = 64` (`HTP_TENSOR_TILE_K_Q2_0 = 64`).
- `Q1_0`: `K_cols = 128` (`HTP_TENSOR_TILE_K_Q1_0 = 128`).
@@ -123,6 +123,17 @@ Block-32 formats use a tile size of 32 rows x 32 columns (`K_cols = 32`):
- **Scales** (bytes 512..639, 128 bytes):
32 pairs of FP16 values (`ggml_half`): `scale[2 * row + 0] = d` (scale) and `scale[2 * row + 1] = m` (min) for each row `row = 0..31`.
+- **Q5_0** (704 bytes per tile, `HTP_TENSOR_TILE_SIZE_Q5_0`):
+ - **Quants** (bytes 0..511): same 4-bit nibble transposition and storage layout as `Q4_0`, holding the low 4 bits.
+ - **5th-bit plane** (bytes 512..639): byte `32*c + row` (`c = 0..3`) holds a flag for `k = 2*(4i + c)` at bit `2i` and for `k = 2*(4i + c) + 1` at bit `2i + 1`, for `i = 0..3`. The flag is set when the 5th bit of the quant is clear.
+ - **Scales** (bytes 640..703, 64 bytes): 32 FP16 values, storing scale `d` for each row `row = 0..31`.
+ - **Dequantization**: `0xF0` is OR-ed into the nibble where the flag is set, giving the signed value `q - 16` (-16..15), scaled by `d`.
+
+- **Q5_1** (768 bytes per tile, `HTP_TENSOR_TILE_SIZE_Q5_1`, same layout as the Q5_K tile):
+ - **Quants** (bytes 0..511) and **scales** (bytes 512..639): same as `Q4_1`, holding the low 4 bits and `(d, m)` pairs.
+ - **5th-bit plane** (bytes 640..767): same as `Q5_K`; a bit is set when the 5th bit of the quant is set.
+ - Uses the `Q5_K` HVX/HMX kernels (Q8_1 activations with block sums).
+
### Q1_0 Format (1-Bit Quantization)
`Q1_0` packs 128 columns per tile (`K_cols = 128`, `HTP_TENSOR_TILE_K_Q1_0 = 128`):
diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
index c38d0005c..03da443c8 100644
--- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp
+++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
@@ -3428,6 +3428,10 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) {
"please update hexagon_type to match ggml_type");
static_assert((unsigned int) HTP_TYPE_Q4_1 == (unsigned int) GGML_TYPE_Q4_1,
"please update hexagon_type to match ggml_type");
+ static_assert((unsigned int) HTP_TYPE_Q5_0 == (unsigned int) GGML_TYPE_Q5_0,
+ "please update hexagon_type to match ggml_type");
+ static_assert((unsigned int) HTP_TYPE_Q5_1 == (unsigned int) GGML_TYPE_Q5_1,
+ "please update hexagon_type to match ggml_type");
static_assert((unsigned int) HTP_TYPE_Q8_0 == (unsigned int) GGML_TYPE_Q8_0,
"please update hexagon_type to match ggml_type");
static_assert((unsigned int) HTP_TYPE_MXFP4 == (unsigned int) GGML_TYPE_MXFP4,
diff --git a/ggml/src/ggml-hexagon/htp-kparams.cpp b/ggml/src/ggml-hexagon/htp-kparams.cpp
index 2fbd419ac..8c75be289 100644
--- a/ggml/src/ggml-hexagon/htp-kparams.cpp
+++ b/ggml/src/ggml-hexagon/htp-kparams.cpp
@@ -953,6 +953,8 @@ static bool htp_matmul_check_types(
case GGML_TYPE_Q2_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
+ case GGML_TYPE_Q5_0:
+ case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_MXFP4:
@@ -995,6 +997,8 @@ static bool htp_matmul_check_types(
case GGML_TYPE_Q2_0:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
+ case GGML_TYPE_Q5_0:
+ case GGML_TYPE_Q5_1:
case GGML_TYPE_Q8_0:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_MXFP4:
diff --git a/ggml/src/ggml-hexagon/htp-repack.cpp b/ggml/src/ggml-hexagon/htp-repack.cpp
index d2e3df2cf..b5413791c 100644
--- a/ggml/src/ggml-hexagon/htp-repack.cpp
+++ b/ggml/src/ggml-hexagon/htp-repack.cpp
@@ -15,7 +15,7 @@
# include <arm_neon.h>
#endif
-// --- 4-bit block-32 nibble helpers (shared by Q4_0, IQ4_NL, Q4_1, MXFP4) ---
+// --- 4-bit block-32 nibble helpers (shared by Q4_0, IQ4_NL, Q4_1, Q5_0, Q5_1, MXFP4) ---
static inline void unpack_nibbles_32(uint8_t * dst, const uint8_t * qs) {
for (unsigned int i = 0; i < 16; ++i) {
@@ -48,6 +48,28 @@ static inline void transpose_tile_to_nibbles(uint8_t tile_quants[32][32], const
}
}
+// 5th-bit plane of one row (qh bit k = quant k), same bit layout as the Q5_K plane
+static inline void pack_q5_plane(uint8_t * plane, int row, uint32_t qh) {
+ for (int c = 0; c < 4; c++) {
+ uint32_t v = 0;
+ for (int i = 0; i < 4; i++) {
+ v |= ((qh >> (8 * i + 2 * c)) & 3) << (2 * i);
+ }
+ plane[c * 32 + row] = (uint8_t) v;
+ }
+}
+
+static inline uint32_t unpack_q5_plane(const uint8_t * plane, int row) {
+ uint32_t qh = 0;
+ for (int c = 0; c < 4; c++) {
+ const uint32_t v = plane[c * 32 + row];
+ for (int i = 0; i < 4; i++) {
+ qh |= ((v >> (2 * i)) & 3) << (8 * i + 2 * c);
+ }
+ }
+ return qh;
+}
+
#if defined(__ARM_NEON) || defined(__aarch64__)
static void repack_neon_transpose8x8(const uint8x8_t rows[8], uint8x8_t columns[8]) {
const uint8x8x2_t t0 = vtrn_u8(rows[0], rows[1]);
@@ -433,7 +455,7 @@ static void unpack_from_tiled_grid_32(void * data, const ggml_tensor * t, size_t
}
}
-// --- Block-32 Quantized Types (Q4_0, IQ4_NL, Q4_1, Q8_0, MXFP4) ---
+// --- Block-32 Quantized Types (Q4_0, IQ4_NL, Q4_1, Q5_0, Q5_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,
@@ -546,6 +568,141 @@ static void repack_tiled_q4_1(void * data, const ggml_tensor * t, size_t offset,
});
}
+static void repack_q5_0_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+ repack_to_tiled_grid_32<block_q5_0, 1>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q5_0,
+ [](uint8_t * tile_dst, const block_q5_0 * src_slice, int64_t ne0, int64_t ne1, int64_t row_stride, int ct, int kt) {
+ uint8_t * plane_dst = tile_dst + 512;
+ ggml_half * scale_dst = (ggml_half *) (tile_dst + 640);
+#if defined(__ARM_NEON) || defined(__aarch64__)
+ if (ct * 32 + 32 <= ne1 && kt < ne0 / 32) {
+ for (int row0 = 0; row0 < 32; row0 += 8) {
+ const uint8_t * qs[8];
+ for (int rr = 0; rr < 8; ++rr) {
+ qs[rr] = src_slice[(ct * 32 + row0 + rr) * row_stride + kt].qs;
+ }
+ repack_neon_4bit_rows(qs, tile_dst, row0);
+ }
+
+ for (int row = 0; row < 32; ++row) {
+ const block_q5_0 & src_block = src_slice[(ct * 32 + row) * row_stride + kt];
+ uint32_t qh;
+ memcpy(&qh, src_block.qh, sizeof(qh));
+ pack_q5_plane(plane_dst, row, ~qh);
+ scale_dst[row] = src_block.d;
+ }
+ return;
+ }
+#endif
+ uint8_t tile_quants[32][32];
+ for (int row = 0; row < 32; row++) {
+ const int64_t r = ct * 32 + row;
+ if (r < ne1 && kt < ne0 / 32) {
+ const block_q5_0 & b = src_slice[r * row_stride + kt];
+ unpack_nibbles_32(tile_quants[row], b.qs);
+ uint32_t qh;
+ memcpy(&qh, b.qh, sizeof(qh));
+ pack_q5_plane(plane_dst, row, ~qh);
+ scale_dst[row] = b.d;
+ } else {
+ memset(tile_quants[row], 0, 32);
+ pack_q5_plane(plane_dst, row, 0);
+ scale_dst[row] = 0;
+ }
+ }
+ transpose_nibbles_to_tile(tile_dst, tile_quants);
+ });
+}
+
+static void repack_tiled_q5_0(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+ unpack_from_tiled_grid_32<block_q5_0, 1>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q5_0,
+ [](block_q5_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 uint8_t * plane_src = tile_src + 512;
+ const ggml_half * scale_src = (const ggml_half *) (tile_src + 640);
+ 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_q5_0 & b = dst_slice[(r - start_row) * row_stride + kt];
+ pack_nibbles_32(b.qs, tile_quants[row]);
+ const uint32_t qh = ~unpack_q5_plane(plane_src, row);
+ memcpy(b.qh, &qh, sizeof(qh));
+ b.d = scale_src[row];
+ }
+ }
+ });
+}
+
+// Q5_1 uses the Q5_K tile layout
+static void repack_q5_1_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) {
+ repack_to_tiled_grid_32<block_q5_1, 1>(t, data, offset, size, HTP_TENSOR_TILE_SIZE_Q5_1,
+ [](uint8_t * tile_dst, const block_q5_1 * src_slice, int64_t ne0, int64_t ne1, int64_t row_stride, int ct, int kt) {
+ ggml_half * scale_dst = (ggml_half *) (tile_dst + 512);
+ uint8_t * plane_dst = tile_dst + 640;
+#if defined(__ARM_NEON) || defined(__aarch64__)
+ if (ct * 32 + 32 <= ne1 && kt < ne0 / 32) {
+ for (int row0 = 0; row0 < 32; row0 += 8) {
+ const uint8_t * qs[8];
+ for (int rr = 0; rr < 8; ++rr) {
+ qs[rr] = src_slice[(ct * 32 + row0 + rr) * row_stride + kt].qs;
+ }
+ repack_neon_4bit_rows(qs, tile_dst, row0);
+ }
+
+ for (int row = 0; row < 32; ++row) {
+ const block_q5_1 & src_block = src_slice[(ct * 32 + row) * row_stride + kt];
+ uint32_t qh;
+ memcpy(&qh, src_block.qh, sizeof(qh));
+ pack_q5_plane(plane_dst, row, qh);
+ scale_dst[2 * row + 0] = src_block.d;
+ scale_dst[2 * row + 1] = src_block.m;
+ }
+ return;
+ }
+#endif
+ uint8_t tile_quants[32][32];
+ for (int row = 0; row < 32; row++) {
+ const int64_t r = ct * 32 + row;
+ if (r < ne1 && kt < ne0 / 32) {
+ const block_q5_1 & b = src_slice[r * row_stride + kt];
+ unpack_nibbles_32(tile_quants[row], b.qs);
+ uint32_t qh;
+ memcpy(&qh, b.qh, sizeof(qh));
+ pack_q5_plane(plane_dst, row, qh);
+ scale_dst[2 * row + 0] = b.d;
+ scale_dst[2 * row + 1] = b.m;
+ } else {
+ memset(tile_quants[row], 0, 32);
+ pack_q5_plane(plane_dst, row, 0);
+ scale_dst[2 * row + 0] = 0;
+ scale_dst[2 * row + 1] = 0;
+ }
+ }
+ transpose_nibbles_to_tile(tile_dst, tile_quants);
+ });
+}
+
+static void repack_tiled_q5_1(void * data, const ggml_tensor * t, size_t offset, size_t size) {
+ unpack_from_tiled_grid_32<block_q5_1, 1>(data, t, offset, size, HTP_TENSOR_TILE_SIZE_Q5_1,
+ [](block_q5_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);
+ const uint8_t * plane_src = tile_src + 640;
+ 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_q5_1 & b = dst_slice[(r - start_row) * row_stride + kt];
+ pack_nibbles_32(b.qs, tile_quants[row]);
+ const uint32_t qh = unpack_q5_plane(plane_src, row);
+ memcpy(b.qh, &qh, sizeof(qh));
+ 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) {
@@ -1538,6 +1695,12 @@ void htp_repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t siz
case GGML_TYPE_Q4_1:
repack_q4_1_tiled(tensor, data, 0, size);
break;
+ case GGML_TYPE_Q5_0:
+ repack_q5_0_tiled(tensor, data, 0, size);
+ break;
+ case GGML_TYPE_Q5_1:
+ repack_q5_1_tiled(tensor, data, 0, size);
+ break;
case GGML_TYPE_Q8_0:
repack_q8_0_tiled(tensor, data, 0, size);
break;
@@ -1593,6 +1756,12 @@ void htp_unpack_tensor_tiled(void * data, const ggml_tensor * tensor, size_t off
case GGML_TYPE_Q4_1:
repack_tiled_q4_1(data, tensor, offset, size);
break;
+ case GGML_TYPE_Q5_0:
+ repack_tiled_q5_0(data, tensor, offset, size);
+ break;
+ case GGML_TYPE_Q5_1:
+ repack_tiled_q5_1(data, tensor, offset, size);
+ break;
case GGML_TYPE_Q8_0:
repack_tiled_q8_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 f3206bc87..5b71f9c5a 100644
--- a/ggml/src/ggml-hexagon/htp-repack.h
+++ b/ggml/src/ggml-hexagon/htp-repack.h
@@ -10,6 +10,7 @@
static inline bool htp_is_tiled_weight_type(enum ggml_type type) {
return type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q2_0 ||
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_Q8_0 || type == GGML_TYPE_Q6_K ||
type == GGML_TYPE_Q2_K || type == GGML_TYPE_Q3_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 590167492..a828012dd 100644
--- a/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h
+++ b/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h
@@ -647,6 +647,106 @@ static void dequantize_tiled_weight_to_fp16_task_q5_k(
}
}
+static void dequantize_tiled_weight_to_fp16_task_q5_0(
+ const tiled_dequantize_state_t *state,
+ uint32_t start_tile, uint32_t end_tile) {
+
+ const HVX_Vector mask_h4 = Q6_Vb_vsplat_R(0x0F);
+
+ 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_sc = hvx_vmem(tile_src + 640);
+ HVX_Vector v_scale_duplicated = Q6_V_lo_W(Q6_W_vshuff_VVR(v_sc, v_sc, -2));
+
+ // Load all 4 groups in parallel
+ HVX_Vector vq0 = hvx_vmem(tile_src + 0 * 128);
+ HVX_Vector vq1 = hvx_vmem(tile_src + 1 * 128);
+ HVX_Vector vq2 = hvx_vmem(tile_src + 2 * 128);
+ HVX_Vector vq3 = hvx_vmem(tile_src + 3 * 128);
+
+ // Nibble extraction
+ HVX_Vector v_lo0 = Q6_V_vand_VV(vq0, mask_h4);
+ HVX_Vector v_hi0 = Q6_Vub_vlsr_VubR(vq0, 4);
+ HVX_Vector v_lo1 = Q6_V_vand_VV(vq1, mask_h4);
+ HVX_Vector v_hi1 = Q6_Vub_vlsr_VubR(vq1, 4);
+ HVX_Vector v_lo2 = Q6_V_vand_VV(vq2, mask_h4);
+ HVX_Vector v_hi2 = Q6_Vub_vlsr_VubR(vq2, 4);
+ HVX_Vector v_lo3 = Q6_V_vand_VV(vq3, mask_h4);
+ HVX_Vector v_hi3 = Q6_Vub_vlsr_VubR(vq3, 4);
+
+ // Q5_0: OR 0xF0 where the plane flags a clear 5th bit (q - 16)
+ HVX_Vector v_plane = hvx_vmem(tile_src + 512);
+ v_lo0 = hvx_q5_0_or_neg(v_lo0, v_plane, 0);
+ v_hi0 = hvx_q5_0_or_neg(v_hi0, v_plane, 1);
+ v_lo1 = hvx_q5_0_or_neg(v_lo1, v_plane, 2);
+ v_hi1 = hvx_q5_0_or_neg(v_hi1, v_plane, 3);
+ v_lo2 = hvx_q5_0_or_neg(v_lo2, v_plane, 4);
+ v_hi2 = hvx_q5_0_or_neg(v_hi2, v_plane, 5);
+ v_lo3 = hvx_q5_0_or_neg(v_lo3, v_plane, 6);
+ v_hi3 = hvx_q5_0_or_neg(v_hi3, v_plane, 7);
+
+ // Shuffling
+ HVX_VectorPair vp_shuf0 = Q6_W_vshuff_VVR(v_hi0, v_lo0, -1);
+ HVX_VectorPair vp_shuf1 = Q6_W_vshuff_VVR(v_hi1, v_lo1, -1);
+ HVX_VectorPair vp_shuf2 = Q6_W_vshuff_VVR(v_hi2, v_lo2, -1);
+ HVX_VectorPair vp_shuf3 = Q6_W_vshuff_VVR(v_hi3, v_lo3, -1);
+
+ // Unpack to 16-bit
+ HVX_VectorPair vp_int16_lo0 = Q6_Wh_vunpack_Vb(Q6_V_lo_W(vp_shuf0));
+ HVX_VectorPair vp_int16_hi0 = Q6_Wh_vunpack_Vb(Q6_V_hi_W(vp_shuf0));
+ HVX_VectorPair vp_int16_lo1 = Q6_Wh_vunpack_Vb(Q6_V_lo_W(vp_shuf1));
+ HVX_VectorPair vp_int16_hi1 = Q6_Wh_vunpack_Vb(Q6_V_hi_W(vp_shuf1));
+ HVX_VectorPair vp_int16_lo2 = Q6_Wh_vunpack_Vb(Q6_V_lo_W(vp_shuf2));
+ HVX_VectorPair vp_int16_hi2 = Q6_Wh_vunpack_Vb(Q6_V_hi_W(vp_shuf2));
+ HVX_VectorPair vp_int16_lo3 = Q6_Wh_vunpack_Vb(Q6_V_lo_W(vp_shuf3));
+ HVX_VectorPair vp_int16_hi3 = Q6_Wh_vunpack_Vb(Q6_V_hi_W(vp_shuf3));
+
+ // Convert and scale multiplication
+ HVX_Vector v_grp0_0 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_lo0)), v_scale_duplicated));
+ HVX_Vector v_grp0_1 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_lo0)), v_scale_duplicated));
+ HVX_Vector v_grp0_2 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_hi0)), v_scale_duplicated));
+ HVX_Vector v_grp0_3 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_hi0)), v_scale_duplicated));
+
+ HVX_Vector v_grp1_0 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_lo1)), v_scale_duplicated));
+ HVX_Vector v_grp1_1 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_lo1)), v_scale_duplicated));
+ HVX_Vector v_grp1_2 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_hi1)), v_scale_duplicated));
+ HVX_Vector v_grp1_3 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_hi1)), v_scale_duplicated));
+
+ HVX_Vector v_grp2_0 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_lo2)), v_scale_duplicated));
+ HVX_Vector v_grp2_1 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_lo2)), v_scale_duplicated));
+ HVX_Vector v_grp2_2 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_hi2)), v_scale_duplicated));
+ HVX_Vector v_grp2_3 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_hi2)), v_scale_duplicated));
+
+ HVX_Vector v_grp3_0 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_lo3)), v_scale_duplicated));
+ HVX_Vector v_grp3_1 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_lo3)), v_scale_duplicated));
+ HVX_Vector v_grp3_2 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_hi3)), v_scale_duplicated));
+ HVX_Vector v_grp3_3 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_hi3)), v_scale_duplicated));
+
+ // Parallel Stores
+ hvx_vmem(dst_ptr + 0 * 64) = v_grp0_0;
+ hvx_vmem(dst_ptr + 1 * 64) = v_grp0_1;
+ hvx_vmem(dst_ptr + 2 * 64) = v_grp0_2;
+ hvx_vmem(dst_ptr + 3 * 64) = v_grp0_3;
+
+ hvx_vmem(dst_ptr + 4 * 64) = v_grp1_0;
+ hvx_vmem(dst_ptr + 5 * 64) = v_grp1_1;
+ hvx_vmem(dst_ptr + 6 * 64) = v_grp1_2;
+ hvx_vmem(dst_ptr + 7 * 64) = v_grp1_3;
+
+ hvx_vmem(dst_ptr + 8 * 64) = v_grp2_0;
+ hvx_vmem(dst_ptr + 9 * 64) = v_grp2_1;
+ hvx_vmem(dst_ptr + 10 * 64) = v_grp2_2;
+ hvx_vmem(dst_ptr + 11 * 64) = v_grp2_3;
+
+ hvx_vmem(dst_ptr + 12 * 64) = v_grp3_0;
+ hvx_vmem(dst_ptr + 13 * 64) = v_grp3_1;
+ hvx_vmem(dst_ptr + 14 * 64) = v_grp3_2;
+ hvx_vmem(dst_ptr + 15 * 64) = v_grp3_3;
+ }
+}
+
// 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(
diff --git a/ggml/src/ggml-hexagon/htp/htp-ops.h b/ggml/src/ggml-hexagon/htp/htp-ops.h
index 7c4b97307..c28b8d61c 100644
--- a/ggml/src/ggml-hexagon/htp/htp-ops.h
+++ b/ggml/src/ggml-hexagon/htp/htp-ops.h
@@ -21,6 +21,8 @@ enum htp_data_type {
HTP_TYPE_F16 = 1,
HTP_TYPE_Q4_0 = 2,
HTP_TYPE_Q4_1 = 3,
+ HTP_TYPE_Q5_0 = 6,
+ HTP_TYPE_Q5_1 = 7,
HTP_TYPE_Q8_0 = 8,
HTP_TYPE_Q2_K = 10,
HTP_TYPE_Q3_K = 11,
diff --git a/ggml/src/ggml-hexagon/htp/htp-tensor.h b/ggml/src/ggml-hexagon/htp/htp-tensor.h
index a5ea0d533..beacc7e4d 100644
--- a/ggml/src/ggml-hexagon/htp/htp-tensor.h
+++ b/ggml/src/ggml-hexagon/htp/htp-tensor.h
@@ -17,6 +17,8 @@
#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_Q5_0 704
+#define HTP_TENSOR_TILE_SIZE_Q5_1 768
#define HTP_TENSOR_TILE_SIZE_Q8_0 1088
#define HTP_TENSOR_TILE_SIZE_IQ4_NL 576
#define HTP_TENSOR_TILE_SIZE_MXFP4 544
@@ -30,7 +32,7 @@
#define HTP_TENSOR_TILE_K_Q2_0 64
#define HTP_TENSOR_TILE_SIZE_Q2_0 576
-// Tiled weight tensor layouts (Q4_0, 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, Q1_0, Q2_0, K-quants):
// See docs/backend/snapdragon/developer.md ("Tiled Weight Formats").
struct htp_tensor_mdev_range {
@@ -143,6 +145,8 @@ static inline uint32_t htp_tensor_get_tile_size(int type) {
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_Q5_0: return HTP_TENSOR_TILE_SIZE_Q5_0;
+ case HTP_TYPE_Q5_1: return HTP_TENSOR_TILE_SIZE_Q5_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;
@@ -182,6 +186,8 @@ static inline uint32_t htp_tensor_get_row_size(int type, uint32_t ne00) {
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_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_Q2_K: return (ne00 / 256) * 84;
case HTP_TYPE_Q2_0: return (ne00 / 64) * 18;
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 01c671679..4a6b2b5ea 100644
--- a/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h
+++ b/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h
@@ -639,6 +639,74 @@ static inline HVX_VectorPair accum_5bit_32x2(
return Q6_W_vcombine_VV(v_sum1, v_sum0);
}
+// OR 0xF0 into every lane of v whose flag j is set in the plane, i.e. 4-bit value | 0xF0 = q - 16 where the 5th bit is clear (Q5_0)
+static inline HVX_Vector hvx_q5_0_or_neg(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, 0xF0F0F0F0);
+}
+
+// Q5_0: signed bytes (q - 16) straight from the nibbles and the plane
+static inline HVX_VectorPair unpack_and_interleave_q5_0_x2(HVX_Vector v_src, HVX_Vector v_plane, int i, HVX_Vector mask_h4) {
+ HVX_Vector v_lo = hvx_q5_0_or_neg(Q6_V_vand_VV(v_src, mask_h4), v_plane, 2 * i);
+ HVX_Vector v_hi = hvx_q5_0_or_neg(Q6_Vub_vlsr_VubR(v_src, 4), v_plane, 2 * i + 1);
+ HVX_VectorPair v01_pair = Q6_W_vshuff_VVR(v_hi, v_lo, -1);
+ HVX_Vector v01_lo = Q6_V_lo_W(v01_pair);
+ HVX_Vector v01_hi = Q6_V_hi_W(v01_pair);
+
+ HVX_Vector v23_lo = Q6_V_valign_VVR(v01_hi, v01_lo, 64);
+ HVX_Vector v_W0 = Q6_V_lo_W(Q6_W_vshuff_VVR(v23_lo, v01_lo, -2));
+
+ HVX_Vector v67_lo = Q6_V_valign_VVR(v01_lo, v01_hi, 64);
+ HVX_Vector v_W1 = Q6_V_lo_W(Q6_W_vshuff_VVR(v67_lo, v01_hi, -2));
+
+ return Q6_W_vcombine_VV(v_W1, v_W0);
+}
+
+static inline HVX_Vector accum_q5_0_32x1(
+ const HVX_Vector * restrict vptr,
+ const HVX_Vector * restrict v_act
+) {
+ HVX_Vector v_sum0 = Q6_V_vzero();
+ HVX_Vector v_sum1 = Q6_V_vzero();
+ HVX_Vector mask_h4 = Q6_Vb_vsplat_R(0x0F);
+ HVX_Vector v_plane = vptr[4];
+
+ #pragma unroll
+ for (int i = 0; i < 4; i++) {
+ HVX_VectorPair v_W_pair = unpack_and_interleave_q5_0_x2(vptr[i], v_plane, i, mask_h4);
+ v_sum0 = Q6_Vw_vrmpyacc_VwVbVb(v_sum0, Q6_V_lo_W(v_W_pair), v_act[i * 2 + 0]);
+ v_sum1 = Q6_Vw_vrmpyacc_VwVbVb(v_sum1, Q6_V_hi_W(v_W_pair), v_act[i * 2 + 1]);
+ }
+
+ return Q6_Vw_vadd_VwVw(v_sum0, v_sum1);
+}
+
+static inline HVX_VectorPair accum_q5_0_32x2(
+ const HVX_Vector * restrict vptr,
+ const HVX_Vector * restrict v_act0,
+ const HVX_Vector * restrict v_act1
+) {
+ HVX_Vector v_sum0 = Q6_V_vzero();
+ HVX_Vector v_sum1 = Q6_V_vzero();
+ HVX_Vector mask_h4 = Q6_Vb_vsplat_R(0x0F);
+ HVX_Vector v_plane = vptr[4];
+
+ #pragma unroll
+ for (int i = 0; i < 4; i++) {
+ HVX_VectorPair v_W_pair = unpack_and_interleave_q5_0_x2(vptr[i], v_plane, i, mask_h4);
+ HVX_Vector v_W0 = Q6_V_lo_W(v_W_pair);
+ HVX_Vector v_W1 = Q6_V_hi_W(v_W_pair);
+
+ 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);
+}
+
// 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);
@@ -1251,6 +1319,82 @@ static void tiled_vec_dot_q8_0_32x2(const uint32_t n, float * restrict s0, float
}
}
+static void tiled_vec_dot_q5_0_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;
+
+ HVX_Vector v_sum_float = Q6_V_vzero();
+
+ uint32_t n_k_tiles = n / 32;
+ for (uint32_t kt = 0; kt < n_k_tiles; kt++) {
+ const HVX_Vector * restrict vptr = (const HVX_Vector *) (tile_ptr + kt * 768);
+ const HVX_Vector * restrict v_act = (const HVX_Vector *) (y_q + kt * 1152);
+
+ HVX_Vector v_sum = accum_q5_0_32x1(vptr, v_act);
+ HVX_Vector v_sum_sf = Q6_Vsf_equals_Vw(v_sum);
+
+ HVX_Vector v_scale_w = vptr[5];
+ HVX_Vector v_scale_a = v_act[8];
+ HVX_Vector v_scale_comb = hvx_vec_mul_f16_f16_to_f32_lower32(v_scale_w, v_scale_a);
+ HVX_Vector v_sum_scaled = hvx_vec_mul_f32_f32(v_sum_sf, v_scale_comb);
+
+ v_sum_float = hvx_vec_add_f32_f32(v_sum_float, v_sum_scaled);
+ }
+
+ 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_q5_0_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;
+
+ 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;
+ for (uint32_t kt = 0; kt < n_k_tiles; kt++) {
+ const HVX_Vector * restrict vptr = (const HVX_Vector *) (tile_ptr + kt * 768);
+ 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_VectorPair v_sums = accum_q5_0_32x2(vptr, v_act0, v_act1);
+ HVX_Vector v_sum_c0 = Q6_V_lo_W(v_sums);
+ HVX_Vector v_sum_c1 = Q6_V_hi_W(v_sums);
+
+ HVX_Vector v_sum_sf_c0 = Q6_Vsf_equals_Vw(v_sum_c0);
+ HVX_Vector v_sum_sf_c1 = Q6_Vsf_equals_Vw(v_sum_c1);
+
+ HVX_Vector v_scale_w = vptr[5];
+ HVX_Vector v_scale_a_c0 = v_act0[8];
+ HVX_Vector v_scale_a_c1 = v_act1[8];
+
+ HVX_Vector v_scale_comb_c0 = hvx_vec_mul_f16_f16_to_f32_lower32(v_scale_w, v_scale_a_c0);
+ HVX_Vector v_scale_comb_c1 = hvx_vec_mul_f16_f16_to_f32_lower32(v_scale_w, v_scale_a_c1);
+
+ HVX_Vector v_sum_scaled_c0 = hvx_vec_mul_f32_f32(v_sum_sf_c0, v_scale_comb_c0);
+ HVX_Vector v_sum_scaled_c1 = hvx_vec_mul_f32_f32(v_sum_sf_c1, v_scale_comb_c1);
+
+ v_sum_float_c0 = hvx_vec_add_f32_f32(v_sum_float_c0, v_sum_scaled_c0);
+ v_sum_float_c1 = hvx_vec_add_f32_f32(v_sum_float_c1, v_sum_scaled_c1);
+ }
+
+ 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_q5_k_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;
diff --git a/ggml/src/ggml-hexagon/htp/matmul-ops.c b/ggml/src/ggml-hexagon/htp/matmul-ops.c
index fb15310cf..2d0750954 100644
--- a/ggml/src/ggml-hexagon/htp/matmul-ops.c
+++ b/ggml/src/ggml-hexagon/htp/matmul-ops.c
@@ -1664,6 +1664,12 @@ static int hvx_mm_init_tiled_vec_dot(struct htp_mm_context * mmctx, enum htp_dat
mmctx->vec_dot_32x1 = tiled_vec_dot_q8_0_32x1;
mmctx->vec_dot_32x2 = tiled_vec_dot_q8_0_32x2;
return 0;
+ case HTP_TYPE_Q5_0:
+ mmctx->type = "q5_0_tiled-f32";
+ mmctx->vec_dot_32x1 = tiled_vec_dot_q5_0_32x1;
+ mmctx->vec_dot_32x2 = tiled_vec_dot_q5_0_32x2;
+ return 0;
+ case HTP_TYPE_Q5_1:
case HTP_TYPE_Q5_K:
mmctx->type = "q5_k_tiled-f32";
mmctx->vec_dot_32x1 = tiled_vec_dot_q5_k_32x1;
@@ -2316,6 +2322,8 @@ static inline tiled_dequant_task_fn_t hmx_get_dequant_worker_fn(int weight_type,
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_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:
case HTP_TYPE_Q5_K: return dequantize_tiled_weight_to_fp16_task_q5_k;
case HTP_TYPE_Q6_K: return dequantize_tiled_weight_to_fp16_task_q6_k;
case HTP_TYPE_Q3_K: return dequantize_tiled_weight_to_fp16_task_q3_k;
diff --git a/ggml/src/ggml-hexagon/htp/matmul-ops.h b/ggml/src/ggml-hexagon/htp/matmul-ops.h
index 0c16fe7b8..7a85624cf 100644
--- a/ggml/src/ggml-hexagon/htp/matmul-ops.h
+++ b/ggml/src/ggml-hexagon/htp/matmul-ops.h
@@ -185,7 +185,7 @@ next_nc:
// (Q2_K: per-16 k sums, q8_1_s16)
static inline bool htp_mm_weight_has_offset(int weight_type) {
return weight_type == HTP_TYPE_Q4_1 || weight_type == HTP_TYPE_Q4_K || weight_type == HTP_TYPE_Q5_K ||
- weight_type == HTP_TYPE_Q2_K || weight_type == HTP_TYPE_Q2_0;
+ weight_type == HTP_TYPE_Q2_K || weight_type == HTP_TYPE_Q2_0 || weight_type == HTP_TYPE_Q5_1;
}
// --- Activation/Row Size Helpers ---