From 0295ef6846d9dc7f485b1e6cda549a197670ef7b Mon Sep 17 00:00:00 2001 From: Jhen-Jie Hong Date: Fri, 2 Oct 2026 05:12:25 +0800 Subject: [PATCH] hexagon: add q2_k and q3_k quant type support (llama/29717) * hexagon: add q2_k and q3_k quant type support * hex-qk: consistent allocation of src1_row_size --------- Co-authored-by: Max Krasnyansky --- ggml/src/ggml-hexagon/ggml-hexagon.cpp | 452 +++++++++++++++++- .../ggml-hexagon/htp/hmx-mm-kernels-tiled.h | 69 +++ ggml/src/ggml-hexagon/htp/htp-ops.h | 2 + .../ggml-hexagon/htp/hvx-mm-kernels-tiled.h | 342 ++++++++++++- ggml/src/ggml-hexagon/htp/matmul-ops.c | 100 +++- ggml/src/ggml-hexagon/htp/matmul-ops.h | 29 +- 6 files changed, 967 insertions(+), 27 deletions(-) diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp index ad0905f6f..f59c49c46 100644 --- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp +++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp @@ -269,10 +269,11 @@ static inline bool ggml_hexagon_is_repack_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_Q4_K || type == GGML_TYPE_Q5_K || + type == GGML_TYPE_Q3_K || type == GGML_TYPE_Q2_K; } -// Size of one repacked row in the DSP tiled layout. The Q6_K, Q5_K and Q4_K tiles store uncompressed scales/mins, +// Size of one repacked row in the DSP tiled layout. The K-quant tiles store uncompressed scales/mins, // so they are larger than the ggml blocks. For the other repack types the tile has the same size as the ggml blocks. static inline size_t ggml_hexagon_tiled_row_size(enum ggml_type type, int64_t ne0) { if (type == GGML_TYPE_Q6_K) { @@ -284,6 +285,12 @@ static inline size_t ggml_hexagon_tiled_row_size(enum ggml_type type, int64_t ne if (type == GGML_TYPE_Q5_K) { return (size_t) (ne0 / 32) * (HTP_MM_WEIGHT_TILE_SIZE_Q5_K / 32); } + if (type == GGML_TYPE_Q3_K) { + return (size_t) (ne0 / 32) * (HTP_MM_WEIGHT_TILE_SIZE_Q3_K / 32); + } + if (type == GGML_TYPE_Q2_K) { + return (size_t) (ne0 / 32) * (HTP_MM_WEIGHT_TILE_SIZE_Q2_K / 32); + } return ggml_row_size(type, ne0); } @@ -1571,6 +1578,397 @@ static void repack_tiled_q6_K(void * data, const ggml_tensor * t, size_t offset, GGML_UNUSED(size); } +// 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; +} + +// tile layout: see HTP_MM_WEIGHT_TILE_SIZE_Q3_K in htp/matmul-ops.h +static void repack_q3_K_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) { + GGML_ASSERT(offset == 0); + + const block_q3_K * src_matrix = (const block_q3_K *) data; + int64_t ne0 = t->ne[0]; + int64_t ne1 = t->ne[1]; + int64_t ne2 = t->ne[2]; + int64_t ne3 = t->ne[3]; + int64_t ne0_padded = hex_round_up(ne0, 32); + int64_t ne1_padded = hex_round_up(ne1, 32); + + GGML_ASSERT(ne0 % QK_K == 0); + + const int n_col_tiles = ne1_padded / 32; + const int n_k_tiles = ne0_padded / 32; + const size_t tile_size = HTP_MM_WEIGHT_TILE_SIZE_Q3_K; + 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 (int i3 = 0; i3 < ne3; i3++) { + for (int i2 = 0; i2 < ne2; i2++) { + const block_q3_K * src_slice = src_matrix + (i3 * ne2 + i2) * (ne1 * sb_per_row); + uint8_t * matrix_dst = (uint8_t *) t->data + (i3 * ne2 + i2) * matrix_size; + + memset(matrix_dst, 0, matrix_size); // padding rows and the OR-ed bits below need zeroed tiles + + for (int64_t r = 0; r < ne1; r++) { + const int ct = (int) (r / 32); + const int row = (int) (r % 32); + const block_q3_K * src_row = src_slice + r * sb_per_row; + + for (int kt = 0; kt < n_k_tiles; kt++) { + const int kt_local = kt % 8; // k-tile within the super-block + const block_q3_K * b = &src_row[kt / 8]; + const float d = GGML_FP16_TO_FP32(b->d); + + uint8_t * tile = matrix_dst + ((size_t) ct * n_k_tiles + kt) * tile_size; + 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)); + } + } + } + } + } + + GGML_UNUSED(size); +} + +// Reverse of repack_q3_K_tiled. Unpacks quants losslessly and normalizes sub-block scales. Read-back only. +static void repack_tiled_q3_K(void * data, const ggml_tensor * t, size_t offset, size_t size) { + GGML_ASSERT(offset == 0); + + block_q3_K * dst_matrix = (block_q3_K *) data; + int64_t ne0 = t->ne[0]; + int64_t ne1 = t->ne[1]; + int64_t ne2 = t->ne[2]; + int64_t ne3 = t->ne[3]; + int64_t ne0_padded = hex_round_up(ne0, 32); + int64_t ne1_padded = hex_round_up(ne1, 32); + + GGML_ASSERT(ne0 % QK_K == 0); + + const int n_col_tiles = ne1_padded / 32; + const int n_k_tiles = ne0_padded / 32; + const size_t tile_size = HTP_MM_WEIGHT_TILE_SIZE_Q3_K; + 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 (int i3 = 0; i3 < ne3; i3++) { + for (int i2 = 0; i2 < ne2; i2++) { + block_q3_K * dst_slice = dst_matrix + (i3 * ne2 + i2) * (ne1 * sb_per_row); + const uint8_t * matrix_src = (const uint8_t *) t->data + (i3 * ne2 + i2) * matrix_size; + + for (int64_t r = 0; r < ne1; r++) { + const int ct = (int) (r / 32); + const int row = (int) (r % 32); + block_q3_K * dst_row = dst_slice + r * sb_per_row; + + for (int64_t sb = 0; sb < sb_per_row; sb++) { + block_q3_K * b = &dst_row[sb]; + memset(b, 0, sizeof(block_q3_K)); + + ggml_half sub_scales[16]; + for (int kt_local = 0; kt_local < 8; kt_local++) { + const int kt = sb * 8 + kt_local; + const uint8_t * tile = matrix_src + ((size_t) ct * n_k_tiles + kt) * 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)) { + // no exact match: same scale choice as quantize_row_q3_K_ref + 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))); + } + } + } + } + } + + GGML_UNUSED(size); +} + +// tile layout: see HTP_MM_WEIGHT_TILE_SIZE_Q2_K in htp/matmul-ops.h +static void repack_q2_K_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) { + GGML_ASSERT(offset == 0); + + const block_q2_K * src_matrix = (const block_q2_K *) data; + int64_t ne0 = t->ne[0]; + int64_t ne1 = t->ne[1]; + int64_t ne2 = t->ne[2]; + int64_t ne3 = t->ne[3]; + int64_t ne0_padded = hex_round_up(ne0, 32); + int64_t ne1_padded = hex_round_up(ne1, 32); + + GGML_ASSERT(ne0 % QK_K == 0); + + const int n_col_tiles = ne1_padded / 32; + const int n_k_tiles = ne0_padded / 32; + const size_t tile_size = HTP_MM_WEIGHT_TILE_SIZE_Q2_K; + 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 (int i3 = 0; i3 < ne3; i3++) { + for (int i2 = 0; i2 < ne2; i2++) { + const block_q2_K * src_slice = src_matrix + (i3 * ne2 + i2) * (ne1 * sb_per_row); + uint8_t * matrix_dst = (uint8_t *) t->data + (i3 * ne2 + i2) * matrix_size; + + memset(matrix_dst, 0, matrix_size); // padding rows and the OR-ed bits below need zeroed tiles + + for (int64_t r = 0; r < ne1; r++) { + const int ct = (int) (r / 32); + const int row = (int) (r % 32); + const block_q2_K * src_row = src_slice + r * sb_per_row; + + for (int kt = 0; kt < n_k_tiles; kt++) { + const int kt_local = kt % 8; // k-tile within the super-block + const block_q2_K * b = &src_row[kt / 8]; + const float d = GGML_FP16_TO_FP32(b->d); + const float dmin = GGML_FP16_TO_FP32(b->dmin); + + uint8_t * tile = matrix_dst + ((size_t) ct * n_k_tiles + kt) * tile_size; + uint8_t * lo_pl = tile; + ggml_half * sc_pl = (ggml_half *) (tile + 256); + ggml_half * m_pl = (ggml_half *) (tile + 384); + + 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++) { + const uint8_t sc = b->scales[kt_local * 2 + sub]; + sc_pl[sub * 32 + row] = GGML_FP32_TO_FP16( d * (float) (sc & 0xF)); + m_pl [sub * 32 + row] = GGML_FP32_TO_FP16(-dmin * (float) (sc >> 4)); + } + } + } + } + } + + GGML_UNUSED(size); +} + +// Reverse of repack_q2_K_tiled. Unpacks quants losslessly and normalizes scales/mins. Read-back only. +static void repack_tiled_q2_K(void * data, const ggml_tensor * t, size_t offset, size_t size) { + GGML_ASSERT(offset == 0); + + block_q2_K * dst_matrix = (block_q2_K *) data; + int64_t ne0 = t->ne[0]; + int64_t ne1 = t->ne[1]; + int64_t ne2 = t->ne[2]; + int64_t ne3 = t->ne[3]; + int64_t ne0_padded = hex_round_up(ne0, 32); + int64_t ne1_padded = hex_round_up(ne1, 32); + + GGML_ASSERT(ne0 % QK_K == 0); + + const int n_col_tiles = ne1_padded / 32; + const int n_k_tiles = ne0_padded / 32; + const size_t tile_size = HTP_MM_WEIGHT_TILE_SIZE_Q2_K; + 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 (int i3 = 0; i3 < ne3; i3++) { + for (int i2 = 0; i2 < ne2; i2++) { + block_q2_K * dst_slice = dst_matrix + (i3 * ne2 + i2) * (ne1 * sb_per_row); + const uint8_t * matrix_src = (const uint8_t *) t->data + (i3 * ne2 + i2) * matrix_size; + + for (int64_t r = 0; r < ne1; r++) { + const int ct = (int) (r / 32); + const int row = (int) (r % 32); + block_q2_K * dst_row = dst_slice + r * sb_per_row; + + for (int64_t sb = 0; sb < sb_per_row; sb++) { + block_q2_K * b = &dst_row[sb]; + memset(b, 0, sizeof(block_q2_K)); + + ggml_half sub_scales[16]; + ggml_half sub_mins[16]; + for (int kt_local = 0; kt_local < 8; kt_local++) { + const int kt = sb * 8 + kt_local; + const uint8_t * tile = matrix_src + ((size_t) ct * n_k_tiles + kt) * tile_size; + const uint8_t * lo_pl = tile; + const ggml_half * sc_pl = (const ggml_half *) (tile + 256); + const ggml_half * m_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; + b->qs[(e / 128) * 32 + e % 32] |= (uint8_t) (lo << (2 * ((e % 128) / 32))); + } + + for (int sub = 0; sub < 2; sub++) { + const float D = GGML_FP16_TO_FP32(sc_pl[sub * 32 + row]); + const float M = GGML_FP16_TO_FP32(m_pl[sub * 32 + row]); + sub_scales[kt_local * 2 + sub] = GGML_FP32_TO_FP16((D > 0.0f) ? D : 0.0f); + sub_mins[kt_local * 2 + sub] = GGML_FP32_TO_FP16((-M > 0.0f) ? -M : 0.0f); + } + } + + int ls[16]; + int lm[16]; + ggml_half * const dd[2] = { &b->d, &b->dmin }; + const ggml_half * const prod[2] = { sub_scales, sub_mins }; + int * const ll[2] = { ls, lm }; + for (int w = 0; w < 2; w++) { + if (hexagon_recover_k_scales(prod[w], 16, 0, 15, dd[w], ll[w])) { + continue; + } + float max_val = 0.0f; + for (int j = 0; j < 16; j++) { + max_val = (std::max)(max_val, GGML_FP16_TO_FP32(prod[w][j])); + } + *dd[w] = GGML_FP32_TO_FP16(max_val / 15.0f); + const float d_actual = GGML_FP16_TO_FP32(*dd[w]); + const float inv_d = (d_actual > 0.0f) ? (1.0f / d_actual) : 0.0f; + for (int j = 0; j < 16; j++) { + ll[w][j] = (std::min)(15, (int) roundf(inv_d * GGML_FP16_TO_FP32(prod[w][j]))); + } + } + + for (int j = 0; j < 16; j++) { + b->scales[j] = (uint8_t) (ls[j] | (lm[j] << 4)); + } + } + } + } + } + + GGML_UNUSED(size); +} + 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; @@ -1985,6 +2383,14 @@ static void repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t repack_q6_K_tiled(tensor, data, 0, size); break; + case GGML_TYPE_Q3_K: + repack_q3_K_tiled(tensor, data, 0, size); + break; + + case GGML_TYPE_Q2_K: + repack_q2_K_tiled(tensor, data, 0, size); + break; + default: break; } @@ -2099,6 +2505,18 @@ static void ggml_backend_hexagon_buffer_get_tensor(ggml_backend_buffer_t buffer, repack_tiled_q6_K(data, tensor, offset, size); break; + case GGML_TYPE_Q3_K: + GGML_ASSERT(offset == 0); + GGML_ASSERT(offset + size <= ggml_nbytes(tensor)); + repack_tiled_q3_K(data, tensor, offset, size); + break; + + case GGML_TYPE_Q2_K: + GGML_ASSERT(offset == 0); + GGML_ASSERT(offset + size <= ggml_nbytes(tensor)); + repack_tiled_q2_K(data, tensor, offset, size); + break; + default: memcpy(data, (const char *) tensor->data + offset, size); break; @@ -2228,6 +2646,14 @@ static void ggml_backend_hexagon_buffer_get_tensor_2d(ggml_backend_buffer_t buff repack_tiled_q6_K(temp_buf.data(), tensor, offset, temp_size); break; + case GGML_TYPE_Q3_K: + repack_tiled_q3_K(temp_buf.data(), tensor, offset, temp_size); + break; + + case GGML_TYPE_Q2_K: + repack_tiled_q2_K(temp_buf.data(), tensor, offset, temp_size); + break; + default: memcpy(temp_buf.data(), (const uint8_t *) tensor->data + offset, temp_size); break; @@ -4773,7 +5199,7 @@ static bool ggml_hexagon_precompute_hmx_mm_params( kparams->n_act_threads = act_threads_selected; kparams->tile_size = htp_mm_get_weight_tile_size(wtype); kparams->aligned_tile_size = aligned_tile_size; - kparams->src1_row_size = (wtype == GGML_TYPE_Q4_1 || wtype == GGML_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); + kparams->src1_row_size = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); kparams->vtcm_size = vtcm_size; kparams->vtcm_src0_size = 0; kparams->div_n_act_threads = init_fastdiv_values(act_threads_selected); @@ -4829,7 +5255,7 @@ static void ggml_hexagon_precompute_hvx_mm_params( if (is_matmul_id) { kparams->kernel_type = (src1_nrows < (int) sess->n_threads) ? HTP_MM_KERNEL_HVX_QUANT_BLOCK : HTP_MM_KERNEL_HVX_QUANT_ROW; - kparams->src1_row_size = (wtype == GGML_TYPE_Q4_1 || wtype == GGML_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); + kparams->src1_row_size = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); struct htp_mm_hvx_vtcm_layout L; uint32_t max_prefetch = (src1_nrows > HTP_MM_HMX_MIN_NROWS) ? 2 : 16; @@ -4857,7 +5283,7 @@ static void ggml_hexagon_precompute_hvx_mm_params( } else { bool try_tiled = (k_align && opt_mm_select >= 1); if (try_tiled) { - kparams->src1_row_size = (wtype == GGML_TYPE_Q4_1 || wtype == GGML_TYPE_Q4_K) + kparams->src1_row_size = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); if (src1_nrows < (int) sess->n_threads) { @@ -5677,7 +6103,7 @@ static void ggml_hexagon_precompute_fused_mmnx_params( { const int src1_nrows = ne11 * ne12 * ne13; - const size_t src1_row_size = (wtype == GGML_TYPE_Q4_1 || wtype == GGML_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); + const size_t src1_row_size = 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 src0_row_size = src0->nb[1]; uint32_t best_n_prefetch = 16; @@ -5772,11 +6198,14 @@ static bool ggml_hexagon_supported_mul_mat(const struct ggml_hexagon_session * s case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: case GGML_TYPE_Q6_K: + case GGML_TYPE_Q3_K: + case GGML_TYPE_Q2_K: if (!ggml_is_contiguous(src0) || ggml_is_permuted(src0)) { return false; } - if (src0->ne[0] % ((src0->type == GGML_TYPE_Q6_K || src0->type == GGML_TYPE_Q5_K || src0->type == GGML_TYPE_Q4_K) ? QK_K : 32)) { + if (src0->ne[0] % ((src0->type == GGML_TYPE_Q6_K || src0->type == GGML_TYPE_Q5_K || src0->type == GGML_TYPE_Q4_K || + src0->type == GGML_TYPE_Q3_K || src0->type == GGML_TYPE_Q2_K) ? QK_K : 32)) { return false; } @@ -5856,11 +6285,14 @@ static bool ggml_hexagon_supported_mul_mat_id(const struct ggml_hexagon_session case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: case GGML_TYPE_Q6_K: + case GGML_TYPE_Q3_K: + case GGML_TYPE_Q2_K: if (!ggml_is_contiguous(src0) || ggml_is_permuted(src0)) { return false; } - if (src0->ne[0] % ((src0->type == GGML_TYPE_Q6_K || src0->type == GGML_TYPE_Q5_K || src0->type == GGML_TYPE_Q4_K) ? QK_K : 32)) { + if (src0->ne[0] % ((src0->type == GGML_TYPE_Q6_K || src0->type == GGML_TYPE_Q5_K || src0->type == GGML_TYPE_Q4_K || + src0->type == GGML_TYPE_Q3_K || src0->type == GGML_TYPE_Q2_K) ? QK_K : 32)) { return false; } @@ -8230,6 +8662,10 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) { "please update hexagon_type to match ggml_type"); static_assert((unsigned int) HTP_TYPE_Q6_K == (unsigned int) GGML_TYPE_Q6_K, "please update hexagon_type to match ggml_type"); + static_assert((unsigned int) HTP_TYPE_Q3_K == (unsigned int) GGML_TYPE_Q3_K, + "please update hexagon_type to match ggml_type"); + static_assert((unsigned int) HTP_TYPE_Q2_K == (unsigned int) GGML_TYPE_Q2_K, + "please update hexagon_type to match ggml_type"); const char * str_verbose = getenv("GGML_HEXAGON_VERBOSE"); const char * str_opbatch = getenv("GGML_HEXAGON_OPBATCH"); 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 5b7f34023..7d7455d04 100644 --- a/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h +++ b/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h @@ -646,6 +646,75 @@ 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. +static void dequantize_tiled_weight_to_fp16_task_q3_k( + const tiled_dequantize_state_t *state, + uint32_t start_tile, uint32_t end_tile) { + + const HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03); + + for (uint32_t t = start_tile; t < end_tile; t++) { + const HVX_Vector * vptr = (const HVX_Vector *) (state->src + t * state->aligned_tile_size); + __fp16 * dst_ptr = state->dst + t * HTP_MM_HMX_TILE_N_ELMS; + + HVX_Vector v_sc = vptr[3]; + HVX_Vector v_sc_k16 = Q6_V_vror_VR(v_sc, 64); + HVX_Vector v_scale_k0 = Q6_V_lo_W(Q6_W_vshuff_VVR(v_sc, v_sc, -2)); + HVX_Vector v_scale_k16 = Q6_V_lo_W(Q6_W_vshuff_VVR(v_sc_k16, v_sc_k16, -2)); + + #pragma unroll + for (int g = 0; g < 8; g++) { + const HVX_Vector v_scale = (g < 4) ? v_scale_k0 : v_scale_k16; + + HVX_Vector v_q = unpack_q3_k_group(vptr, g, mask_03); + HVX_VectorPair vp16 = Q6_Wh_vunpack_Vb(v_q); + HVX_VectorPair vp_k = Q6_W_vdeal_VVR(Q6_V_hi_W(vp16), Q6_V_lo_W(vp16), -4); + + hvx_vmem(dst_ptr + (2 * g + 0) * 64) = + Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_k)), v_scale)); + hvx_vmem(dst_ptr + (2 * g + 1) * 64) = + Q6_Vhf_equals_Vqf16(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_k)), v_scale)); + } + } +} + +// Q2_K stores 2-bit weights and one fp16 scale and offset per 16 k, see HTP_MM_WEIGHT_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) { + + const HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03); + + for (uint32_t t = start_tile; t < end_tile; t++) { + const HVX_Vector * vptr = (const HVX_Vector *) (state->src + t * state->aligned_tile_size); + __fp16 * dst_ptr = state->dst + t * HTP_MM_HMX_TILE_N_ELMS; + + HVX_Vector v_sc = vptr[2]; + HVX_Vector v_sc_k16 = Q6_V_vror_VR(v_sc, 64); + HVX_Vector v_m = vptr[3]; + HVX_Vector v_m_k16 = Q6_V_vror_VR(v_m, 64); + HVX_Vector v_scale_k0 = Q6_V_lo_W(Q6_W_vshuff_VVR(v_sc, v_sc, -2)); + HVX_Vector v_scale_k16 = Q6_V_lo_W(Q6_W_vshuff_VVR(v_sc_k16, v_sc_k16, -2)); + HVX_Vector v_offset_k0 = Q6_V_lo_W(Q6_W_vshuff_VVR(v_m, v_m, -2)); + HVX_Vector v_offset_k16 = Q6_V_lo_W(Q6_W_vshuff_VVR(v_m_k16, v_m_k16, -2)); + + #pragma unroll + for (int g = 0; g < 8; g++) { + const HVX_Vector v_scale = (g < 4) ? v_scale_k0 : v_scale_k16; + const HVX_Vector v_offset = (g < 4) ? v_offset_k0 : v_offset_k16; + + HVX_Vector v_q = unpack_q3_k_low2(vptr, g, mask_03); + HVX_VectorPair vp16 = Q6_Wh_vunpack_Vb(v_q); + HVX_VectorPair vp_k = Q6_W_vdeal_VVR(Q6_V_hi_W(vp16), Q6_V_lo_W(vp16), -4); + + hvx_vmem(dst_ptr + (2 * g + 0) * 64) = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf( + Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_k)), v_scale), v_offset)); + hvx_vmem(dst_ptr + (2 * g + 1) * 64) = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf( + Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_k)), v_scale), v_offset)); + } + } +} + static __attribute__((noinline)) void convert_f16_weight_to_fp16_tiles_task( 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 3dbde5adf..a6a3bb85a 100644 --- a/ggml/src/ggml-hexagon/htp/htp-ops.h +++ b/ggml/src/ggml-hexagon/htp/htp-ops.h @@ -22,6 +22,8 @@ enum htp_data_type { HTP_TYPE_Q4_0 = 2, HTP_TYPE_Q4_1 = 3, HTP_TYPE_Q8_0 = 8, + HTP_TYPE_Q2_K = 10, + HTP_TYPE_Q3_K = 11, HTP_TYPE_Q4_K = 12, HTP_TYPE_Q5_K = 13, HTP_TYPE_Q6_K = 14, 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 4564b1524..580bae4ac 100644 --- a/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h +++ b/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h @@ -1,6 +1,9 @@ // Dynamic quantizers that produce tiled activations -static inline void quantize_block_f32_q8_1_tiled(float * restrict x, uint8_t * restrict y_block) { +// vector 9: d * sum(q) of the 32 k, replicated (q8_1), or per 16 k for Q2_K (q8_1_s16): k 0..15 in lanes 0..31, +// k 16..31 in lanes 32..63, with d the fp16 scale of vector 8 +__attribute__((always_inline)) +static inline void quantize_block_f32_q8_1_tiled_impl(float * restrict x, uint8_t * restrict y_block, const bool sums16) { assert((unsigned long) x % 128 == 0); assert((unsigned long) y_block % 128 == 0); @@ -46,7 +49,10 @@ static inline void quantize_block_f32_q8_1_tiled(float * restrict x, uint8_t * r HVX_Vector v_sums = Q6_Vw_vrmpy_VbVb(vx_i8, ones); v_sums = Q6_Vw_vadd_VwVw(v_sums, Q6_V_vror_VR(v_sums, 4)); v_sums = Q6_Vw_vadd_VwVw(v_sums, Q6_V_vror_VR(v_sums, 8)); - v_sums = Q6_Vw_vadd_VwVw(v_sums, Q6_V_vror_VR(v_sums, 16)); + if (!sums16) { + v_sums = Q6_Vw_vadd_VwVw(v_sums, Q6_V_vror_VR(v_sums, 16)); + } + // word 8b: sum of block b (sums16: k 0..15, and word 8b + 4: k 16..31) const HVX_Vector v_inv127 = hvx_vec_splat_f32(1.0f / 127.0f); HVX_Vector vd0_sf = hvx_vec_mul_f32_f32(vmax0_sf, v_inv127); @@ -55,6 +61,15 @@ static inline void quantize_block_f32_q8_1_tiled(float * restrict x, uint8_t * r HVX_Vector vd3_sf = hvx_vec_mul_f32_f32(vmax3_sf, v_inv127); HVX_Vector v_sums_sf = Q6_Vsf_equals_Vw(v_sums); + if (sums16) { + // the fp16 d of vector 8 + HVX_VectorPair vd01_sf = hvx_vec_f16_to_f32(vd01_hf); + HVX_VectorPair vd23_sf = hvx_vec_f16_to_f32(vd23_hf); + vd0_sf = Q6_V_lo_W(vd01_sf); + vd1_sf = Q6_V_hi_W(vd01_sf); + vd2_sf = Q6_V_lo_W(vd23_sf); + vd3_sf = Q6_V_hi_W(vd23_sf); + } HVX_Vector voff0_sf = hvx_vec_mul_f32_f32(vd0_sf, v_sums_sf); HVX_Vector voff1_sf = hvx_vec_mul_f32_f32(vd1_sf, Q6_V_vror_VR(v_sums_sf, 32)); HVX_Vector voff2_sf = hvx_vec_mul_f32_f32(vd2_sf, Q6_V_vror_VR(v_sums_sf, 64)); @@ -75,6 +90,14 @@ static inline void quantize_block_f32_q8_1_tiled(float * restrict x, uint8_t * r hvx_vec_repl_f16(voff23_hf), hvx_vec_repl_f16(Q6_V_vror_VR(voff23_hf, 64)), }; + if (sums16) { + // the k 16..31 sums sit 4 halfwords after the k 0..15 sums + const HVX_VectorPred q_lo = Q6_Q_vsetq_R(64); + r_offset[0] = Q6_V_vmux_QVV(q_lo, r_offset[0], hvx_vec_repl_f16(Q6_V_vror_VR(voff01_hf, 8))); + r_offset[1] = Q6_V_vmux_QVV(q_lo, r_offset[1], hvx_vec_repl_f16(Q6_V_vror_VR(voff01_hf, 72))); + r_offset[2] = Q6_V_vmux_QVV(q_lo, r_offset[2], hvx_vec_repl_f16(Q6_V_vror_VR(voff23_hf, 8))); + r_offset[3] = Q6_V_vmux_QVV(q_lo, r_offset[3], hvx_vec_repl_f16(Q6_V_vror_VR(voff23_hf, 72))); + } 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, @@ -114,6 +137,14 @@ static inline void quantize_block_f32_q8_1_tiled(float * restrict x, uint8_t * r } } +static inline void quantize_block_f32_q8_1_tiled(float * restrict x, uint8_t * restrict y_block) { + quantize_block_f32_q8_1_tiled_impl(x, y_block, false); +} + +static inline void quantize_block_f32_q8_1_s16_tiled(float * restrict x, uint8_t * restrict y_block) { + quantize_block_f32_q8_1_tiled_impl(x, y_block, true); +} + static inline void quantize_block_f32_q8_0_tiled(float * restrict x, uint8_t * restrict y_block) { assert((unsigned long) x % 128 == 0); assert((unsigned long) y_block % 128 == 0); @@ -229,6 +260,17 @@ static void quantize_row_f32_q8_1_tiled(float * restrict x, uint8_t * restrict y } } +static void quantize_row_f32_q8_1_s16_tiled(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 * 1280; + quantize_block_f32_q8_1_s16_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) { @@ -243,6 +285,19 @@ static inline HVX_Vector hvx_vec_mul_f16_f16_to_f32_lower32(HVX_Vector v1, HVX_V #endif } +// both halves of hvx_vec_mul_f16_f16_to_f32_lower32: lo = products of lanes 0..31, hi = lanes 32..63 +static inline HVX_VectorPair hvx_vec_mul_f16_f16_to_f32_pair(HVX_Vector v1, HVX_Vector v2) { +#if __HVX_ARCH__ >= 79 + HVX_VectorPair p = Q6_Wsf_vmpy_VhfVhf(v1, v2); + return Q6_W_vshuff_VVR(Q6_V_hi_W(p), Q6_V_lo_W(p), -4); +#else + HVX_VectorPair p = Q6_Wqf32_vmpy_VhfVhf(v1, v2); + HVX_Vector hi = Q6_Vsf_equals_Vqf32(Q6_V_hi_W(p)); + HVX_Vector lo = Q6_Vsf_equals_Vqf32(Q6_V_lo_W(p)); + return Q6_W_vshuff_VVR(hi, lo, -4); +#endif +} + static inline HVX_Vector unpack_and_interleave_4bit(HVX_Vector v_a, HVX_Vector v_b, HVX_Vector mask_h4) { HVX_Vector v_W0 = Q6_V_vand_VV(v_a, mask_h4); HVX_Vector v_W1 = Q6_Vub_vlsr_VubR(v_a, 4); @@ -529,13 +584,131 @@ static inline void accum_q6_k_32x2( // scale the two half sums with the per-row tile scales (v_scale_w = vptr[6]) and the activation scale static inline HVX_Vector scale_q6_k_32x1(HVX_VectorPair v_sums, HVX_Vector v_scale_w, HVX_Vector v_scale_a) { - HVX_Vector v_scale_lo = hvx_vec_mul_f16_f16_to_f32_lower32(v_scale_w, v_scale_a); - HVX_Vector v_scale_hi = hvx_vec_mul_f16_f16_to_f32_lower32(Q6_V_vror_VR(v_scale_w, 64), v_scale_a); - HVX_Vector v_lo = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(v_sums)), v_scale_lo); - HVX_Vector v_hi = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_hi_W(v_sums)), v_scale_hi); + HVX_VectorPair v_scale = hvx_vec_mul_f16_f16_to_f32_pair(v_scale_w, v_scale_a); + HVX_Vector v_lo = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_lo_W(v_sums)), Q6_V_lo_W(v_scale)); + HVX_Vector v_hi = hvx_vec_mul_f32_f32(Q6_Vsf_equals_Vw(Q6_V_hi_W(v_sums)), Q6_V_hi_W(v_scale)); 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 +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) { + return Q6_Vub_vlsr_VubR(v, 6); + } + if (g & 3) { + v = Q6_Vub_vlsr_VubR(v, 2 * (g & 3)); + } + return Q6_V_vand_VV(v, mask_03); +} + +// Q3_K k-group g as signed bytes: low2 | 0xFC (= low2 - 4) where bit g of vector 2 is set +static inline HVX_Vector unpack_q3_k_group(const HVX_Vector * restrict vptr, int g, HVX_Vector mask_03) { + HVX_VectorPred q_neg = Q6_Q_vand_VR(vptr[2], 0x01010101u << g); + return Q6_V_vandor_VQR(unpack_q3_k_low2(vptr, g, mask_03), q_neg, 0xFCFCFCFC); +} + +// same half split as accum_q6_k_32x1 +static inline HVX_VectorPair accum_q3_k_32x1( + const HVX_Vector * restrict vptr, + const HVX_Vector * restrict v_act +) { + HVX_Vector v_sum_lo = Q6_V_vzero(); + HVX_Vector v_sum_hi = Q6_V_vzero(); + HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03); + + #pragma unroll + for (int g = 0; g < 4; g++) { + HVX_Vector v_W_lo = unpack_q3_k_group(vptr, g, mask_03); + HVX_Vector v_W_hi = unpack_q3_k_group(vptr, g + 4, mask_03); + v_sum_lo = Q6_Vw_vrmpyacc_VwVbVb(v_sum_lo, v_W_lo, v_act[g]); + v_sum_hi = Q6_Vw_vrmpyacc_VwVbVb(v_sum_hi, v_W_hi, v_act[g + 4]); + } + + return Q6_W_vcombine_VV(v_sum_hi, v_sum_lo); +} + +static inline void accum_q3_k_32x2( + const HVX_Vector * restrict vptr, + const HVX_Vector * restrict v_act0, + const HVX_Vector * restrict v_act1, + HVX_VectorPair * v_sums0, + HVX_VectorPair * v_sums1 +) { + HVX_Vector v_sum0_lo = Q6_V_vzero(); + HVX_Vector v_sum0_hi = Q6_V_vzero(); + HVX_Vector v_sum1_lo = Q6_V_vzero(); + HVX_Vector v_sum1_hi = Q6_V_vzero(); + HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03); + + #pragma unroll + for (int g = 0; g < 4; g++) { + HVX_Vector v_W_lo = unpack_q3_k_group(vptr, g, mask_03); + HVX_Vector v_W_hi = unpack_q3_k_group(vptr, g + 4, mask_03); + v_sum0_lo = Q6_Vw_vrmpyacc_VwVbVb(v_sum0_lo, v_W_lo, v_act0[g]); + v_sum0_hi = Q6_Vw_vrmpyacc_VwVbVb(v_sum0_hi, v_W_hi, v_act0[g + 4]); + v_sum1_lo = Q6_Vw_vrmpyacc_VwVbVb(v_sum1_lo, v_W_lo, v_act1[g]); + v_sum1_hi = Q6_Vw_vrmpyacc_VwVbVb(v_sum1_hi, v_W_hi, v_act1[g + 4]); + } + + *v_sums0 = Q6_W_vcombine_VV(v_sum0_hi, v_sum0_lo); + *v_sums1 = Q6_W_vcombine_VV(v_sum1_hi, v_sum1_lo); +} + +// Q2_K (x = D * q + M): the Q3_K dot products without the -4 flags, M uses the q8_1_s16 sums +static inline HVX_VectorPair accum_q2_k_32x1( + const HVX_Vector * restrict vptr, + const HVX_Vector * restrict v_act +) { + HVX_Vector v_sum_lo = Q6_V_vzero(); + HVX_Vector v_sum_hi = Q6_V_vzero(); + HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03); + + #pragma unroll + for (int g = 0; g < 4; g++) { + HVX_Vector v_W_lo = unpack_q3_k_low2(vptr, g, mask_03); + HVX_Vector v_W_hi = unpack_q3_k_low2(vptr, g + 4, mask_03); + v_sum_lo = Q6_Vw_vrmpyacc_VwVbVb(v_sum_lo, v_W_lo, v_act[g]); + v_sum_hi = Q6_Vw_vrmpyacc_VwVbVb(v_sum_hi, v_W_hi, v_act[g + 4]); + } + + return Q6_W_vcombine_VV(v_sum_hi, v_sum_lo); +} + +static inline void accum_q2_k_32x2( + const HVX_Vector * restrict vptr, + const HVX_Vector * restrict v_act0, + const HVX_Vector * restrict v_act1, + HVX_VectorPair * v_sums0, + HVX_VectorPair * v_sums1 +) { + HVX_Vector v_sum0_lo = Q6_V_vzero(); + HVX_Vector v_sum0_hi = Q6_V_vzero(); + HVX_Vector v_sum1_lo = Q6_V_vzero(); + HVX_Vector v_sum1_hi = Q6_V_vzero(); + HVX_Vector mask_03 = Q6_Vb_vsplat_R(0x03); + + #pragma unroll + for (int g = 0; g < 4; g++) { + HVX_Vector v_W_lo = unpack_q3_k_low2(vptr, g, mask_03); + HVX_Vector v_W_hi = unpack_q3_k_low2(vptr, g + 4, mask_03); + v_sum0_lo = Q6_Vw_vrmpyacc_VwVbVb(v_sum0_lo, v_W_lo, v_act0[g]); + v_sum0_hi = Q6_Vw_vrmpyacc_VwVbVb(v_sum0_hi, v_W_hi, v_act0[g + 4]); + v_sum1_lo = Q6_Vw_vrmpyacc_VwVbVb(v_sum1_lo, v_W_lo, v_act1[g]); + v_sum1_hi = Q6_Vw_vrmpyacc_VwVbVb(v_sum1_hi, v_W_hi, v_act1[g + 4]); + } + + *v_sums0 = Q6_W_vcombine_VV(v_sum0_hi, v_sum0_lo); + *v_sums1 = Q6_W_vcombine_VV(v_sum1_hi, v_sum1_lo); +} + +// D as the Q6_K scales, plus M (vector 3) times the per-16 activation sums in v_act[9] +static inline HVX_Vector scale_q2_k_32x1(HVX_VectorPair v_sums, const HVX_Vector * restrict vptr, const HVX_Vector * restrict v_act) { + HVX_VectorPair v_m = hvx_vec_mul_f16_f16_to_f32_pair(vptr[3], v_act[9]); + HVX_Vector v_d = scale_q6_k_32x1(v_sums, vptr[2], v_act[8]); + return hvx_vec_add_f32_f32(v_d, hvx_vec_add_f32_f32(Q6_V_lo_W(v_m), Q6_V_hi_W(v_m))); +} + static void tiled_vec_dot_q4_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; @@ -939,6 +1112,116 @@ static void tiled_vec_dot_q6_k_32x2(const uint32_t n, float * restrict s0, float } } +static void tiled_vec_dot_q3_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; + + 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 * 512); + const HVX_Vector * restrict v_act = (const HVX_Vector *) (y_q + kt * 1152); + + HVX_VectorPair v_sums = accum_q3_k_32x1(vptr, v_act); + v_sum_float = hvx_vec_add_f32_f32(v_sum_float, scale_q6_k_32x1(v_sums, vptr[3], v_act[8])); + } + + 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_q3_k_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 * 512); + 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_sums0, v_sums1; + accum_q3_k_32x2(vptr, v_act0, v_act1, &v_sums0, &v_sums1); + + v_sum_float_c0 = hvx_vec_add_f32_f32(v_sum_float_c0, scale_q6_k_32x1(v_sums0, vptr[3], v_act0[8])); + v_sum_float_c1 = hvx_vec_add_f32_f32(v_sum_float_c1, scale_q6_k_32x1(v_sums1, vptr[3], v_act1[8])); + } + + 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_q2_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; + + 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 * 512); + const HVX_Vector * restrict v_act = (const HVX_Vector *) (y_q + kt * 1280); + + HVX_VectorPair v_sums = accum_q2_k_32x1(vptr, v_act); + v_sum_float = hvx_vec_add_f32_f32(v_sum_float, scale_q2_k_32x1(v_sums, vptr, v_act)); + } + + 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_q2_k_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 * 512); + const HVX_Vector * restrict v_act0 = (const HVX_Vector *) (y0_q + kt * 1280); + const HVX_Vector * restrict v_act1 = (const HVX_Vector *) (y1_q + kt * 1280); + + HVX_VectorPair v_sums0, v_sums1; + accum_q2_k_32x2(vptr, v_act0, v_act1, &v_sums0, &v_sums1); + + v_sum_float_c0 = hvx_vec_add_f32_f32(v_sum_float_c0, scale_q2_k_32x1(v_sums0, vptr, v_act0)); + v_sum_float_c1 = hvx_vec_add_f32_f32(v_sum_float_c1, scale_q2_k_32x1(v_sums1, vptr, v_act1)); + } + + 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_iq4nl_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; @@ -1159,6 +1442,23 @@ static inline void quantize_f32_q8_1_tiled_kernel( } } +static inline void quantize_f32_q8_1_s16_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_q8_1_s16_tiled((float *) 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, @@ -1218,3 +1518,33 @@ static inline void quantize_f32_q8_1_tiled_block_kernel( } } } + +static inline void quantize_f32_q8_1_s16_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 * 1280; + + quantize_block_f32_q8_1_s16_tiled((float *) src_ptr, dst_ptr); + + c++; + if (c == nb) { + c = 0; + r++; + } + } +} diff --git a/ggml/src/ggml-hexagon/htp/matmul-ops.c b/ggml/src/ggml-hexagon/htp/matmul-ops.c index 727dc281f..9dfd35649 100644 --- a/ggml/src/ggml-hexagon/htp/matmul-ops.c +++ b/ggml/src/ggml-hexagon/htp/matmul-ops.c @@ -545,6 +545,8 @@ MATMUL_2D_REPACKED_IMPL(q4_1, 640, tiled_vec_dot_q4_1_32x2, tiled_vec_do MATMUL_2D_REPACKED_IMPL(q8_0, 1088, tiled_vec_dot_q8_0_32x2, tiled_vec_dot_q8_0_32x1) MATMUL_2D_REPACKED_IMPL(q6_k, 896, tiled_vec_dot_q6_k_32x2, tiled_vec_dot_q6_k_32x1) MATMUL_2D_REPACKED_IMPL(q5_k, 768, tiled_vec_dot_q5_k_32x2, tiled_vec_dot_q5_k_32x1) +MATMUL_2D_REPACKED_IMPL(q3_k, 512, tiled_vec_dot_q3_k_32x2, tiled_vec_dot_q3_k_32x1) +MATMUL_2D_REPACKED_IMPL(q2_k, 512, tiled_vec_dot_q2_k_32x2, tiled_vec_dot_q2_k_32x1) MATMUL_2D_REPACKED_IMPL(iq4nl, 576, tiled_vec_dot_iq4nl_32x2, tiled_vec_dot_iq4nl_32x1) MATMUL_2D_REPACKED_IMPL(mxfp4, 544, tiled_vec_dot_mxfp4_32x2, tiled_vec_dot_mxfp4_32x1) @@ -652,6 +654,7 @@ static void name(unsigned int nth, unsigned int ith, void * data) { QUANTIZE_IMPL(quantize_f32_q8_0_tiled, "quantize-f32-q8_0_tiled", quantize_f32_q8_0_tiled_kernel, htp_mm_q8_0_tiled_row_size(ne0)) QUANTIZE_IMPL(quantize_f32_q8_1_tiled, "quantize-f32-q8_1_tiled", quantize_f32_q8_1_tiled_kernel, htp_mm_q8_1_tiled_row_size(ne0)) +QUANTIZE_IMPL(quantize_f32_q8_1_s16_tiled, "quantize-f32-q8_1_s16_tiled", quantize_f32_q8_1_s16_tiled_kernel, htp_mm_q8_1_tiled_row_size(ne0)) QUANTIZE_IMPL(quantize_f32_f32, "quantize-f32-f32", quantize_f32_f32_kernel, mmctx->vtcm_src1_stride) QUANTIZE_IMPL(quantize_f32_f16, "quantize-f32-f16", quantize_f32_f16_kernel, mmctx->vtcm_src1_stride) QUANTIZE_IMPL(quantize_f16_f16, "quantize-f16-f16", quantize_f16_f16_kernel, mmctx->vtcm_src1_stride) @@ -712,11 +715,56 @@ static void quantize_f32_q8_1_tiled_block(unsigned int nth, unsigned int ith, vo htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_A_QUANT, mmctx->quant_ib_first[ith]); } +static void quantize_f32_q8_1_s16_tiled_block(unsigned int nth, unsigned int ith, void * data) { + (void) nth; + struct htp_mm_context * mmctx = data; + if (mmctx->quant_ib_first[ith] >= mmctx->quant_ib_last[ith]) { + return; + } + struct htp_ops_context * octx = mmctx->octx; + struct htp_thread_trace * tr = &octx->ctx->trace[ith]; + htp_trace_event_start(tr, HTP_TRACE_EVT_HVX_A_QUANT, mmctx->quant_ib_first[ith]); + + const struct htp_tensor * src = mmctx->act; + + quantize_f32_q8_1_s16_tiled_block_kernel( + (const float *) mmctx->vtcm_act_raw, + mmctx->vtcm_src1, + NULL, + src->ne[0], + mmctx->quant_ib_first[ith], + mmctx->quant_ib_last[ith], + mmctx->vtcm_act_raw_stride, + htp_mm_q8_1_tiled_row_size(src->ne[0]), + mmctx->quant_r[ith], + mmctx->quant_c[ith] + ); + + htp_trace_event_stop(tr, HTP_TRACE_EVT_HVX_A_QUANT, mmctx->quant_ib_first[ith]); +} + +// q8_1 for weight types with offsets (q8_1_s16 for Q2_K), otherwise q8_0 +static inline worker_callback_t htp_mm_act_quant_row_func(int weight_type) { + if (weight_type == HTP_TYPE_Q2_K) { + return quantize_f32_q8_1_s16_tiled; + } + return htp_mm_weight_has_offset(weight_type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; +} + +static inline worker_callback_t htp_mm_act_quant_block_func(int weight_type) { + if (weight_type == HTP_TYPE_Q2_K) { + return quantize_f32_q8_1_s16_tiled_block; + } + return htp_mm_weight_has_offset(weight_type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; +} + MATVEC_2D_REPACKED_IMPL(q4_0, 576, tiled_vec_dot_q4_0_32x1) MATVEC_2D_REPACKED_IMPL(q4_1, 640, tiled_vec_dot_q4_1_32x1) MATVEC_2D_REPACKED_IMPL(q8_0, 1088, tiled_vec_dot_q8_0_32x1) MATVEC_2D_REPACKED_IMPL(q5_k, 768, tiled_vec_dot_q5_k_32x1) MATVEC_2D_REPACKED_IMPL(q6_k, 896, tiled_vec_dot_q6_k_32x1) +MATVEC_2D_REPACKED_IMPL(q3_k, 512, tiled_vec_dot_q3_k_32x1) +MATVEC_2D_REPACKED_IMPL(q2_k, 512, tiled_vec_dot_q2_k_32x1) MATVEC_2D_REPACKED_IMPL(iq4nl, 576, tiled_vec_dot_iq4nl_32x1) MATVEC_2D_REPACKED_IMPL(mxfp4, 544, tiled_vec_dot_mxfp4_32x1) @@ -726,6 +774,8 @@ MATMUL_NX_2D_REPACKED_IMPL(q8_0, 1088, tiled_vec_dot_q8_0_32x2, tiled_vec_do MATMUL_NX_2D_REPACKED_IMPL(iq4nl, 576, tiled_vec_dot_iq4nl_32x2, tiled_vec_dot_iq4nl_32x1) MATMUL_NX_2D_REPACKED_IMPL(mxfp4, 544, tiled_vec_dot_mxfp4_32x2, tiled_vec_dot_mxfp4_32x1) MATMUL_NX_2D_REPACKED_IMPL(q5_k, 768, tiled_vec_dot_q5_k_32x2, tiled_vec_dot_q5_k_32x1) +MATMUL_NX_2D_REPACKED_IMPL(q3_k, 512, tiled_vec_dot_q3_k_32x2, tiled_vec_dot_q3_k_32x1) +MATMUL_NX_2D_REPACKED_IMPL(q2_k, 512, tiled_vec_dot_q2_k_32x2, tiled_vec_dot_q2_k_32x1) #define MATMUL_4D_REPACKED_IMPL(SUFFIX, TILE_SIZE, DOT_2X2, DOT_2X1) \ static void hvx_mm_4d_repacked_##SUFFIX(unsigned int nth, unsigned int ith, void * data) { \ @@ -859,6 +909,8 @@ MATMUL_4D_REPACKED_IMPL(q4_1, 640, tiled_vec_dot_q4_1_32x2, tiled_vec_do MATMUL_4D_REPACKED_IMPL(q8_0, 1088, tiled_vec_dot_q8_0_32x2, tiled_vec_dot_q8_0_32x1) MATMUL_4D_REPACKED_IMPL(q6_k, 896, tiled_vec_dot_q6_k_32x2, tiled_vec_dot_q6_k_32x1) MATMUL_4D_REPACKED_IMPL(q5_k, 768, tiled_vec_dot_q5_k_32x2, tiled_vec_dot_q5_k_32x1) +MATMUL_4D_REPACKED_IMPL(q3_k, 512, tiled_vec_dot_q3_k_32x2, tiled_vec_dot_q3_k_32x1) +MATMUL_4D_REPACKED_IMPL(q2_k, 512, tiled_vec_dot_q2_k_32x2, tiled_vec_dot_q2_k_32x1) MATMUL_4D_REPACKED_IMPL(iq4nl, 576, tiled_vec_dot_iq4nl_32x2, tiled_vec_dot_iq4nl_32x1) MATMUL_4D_REPACKED_IMPL(mxfp4, 544, tiled_vec_dot_mxfp4_32x2, tiled_vec_dot_mxfp4_32x1) @@ -1600,6 +1652,14 @@ static int hvx_mm_init_vec_dot(struct htp_mm_context * mmctx, enum htp_data_type mmctx->type = "q6_k_tiled-f32"; mmctx->vec_dot_32x1 = tiled_vec_dot_q6_k_32x1; return 0; + case HTP_TYPE_Q3_K: + mmctx->type = "q3_k_tiled-f32"; + mmctx->vec_dot_32x1 = tiled_vec_dot_q3_k_32x1; + return 0; + case HTP_TYPE_Q2_K: + mmctx->type = "q2_k_tiled-f32"; + mmctx->vec_dot_32x1 = tiled_vec_dot_q2_k_32x1; + return 0; case HTP_TYPE_IQ4_NL: mmctx->type = "iq4nl_tiled-f32"; mmctx->vec_dot_32x1 = tiled_vec_dot_iq4nl_32x1; @@ -1651,7 +1711,8 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { bool is_repacked = (src0->type == HTP_TYPE_Q4_0 || src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q8_0 || src0->type == HTP_TYPE_IQ4_NL || src0->type == HTP_TYPE_MXFP4 || src0->type == HTP_TYPE_Q6_K || - src0->type == HTP_TYPE_Q4_K || src0->type == HTP_TYPE_Q5_K); + src0->type == HTP_TYPE_Q4_K || src0->type == HTP_TYPE_Q5_K || + src0->type == HTP_TYPE_Q3_K || src0->type == HTP_TYPE_Q2_K); // Compute src0_nrows_per_thread mmctx->src0_nrows_per_thread = fastdiv(nrows + octx->n_threads - 1, &octx->n_threads_div); @@ -1681,6 +1742,8 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { case HTP_TYPE_Q8_0: matmul_job_func = hvx_mm_4d_repacked_q8_0; break; case HTP_TYPE_Q6_K: matmul_job_func = hvx_mm_4d_repacked_q6_k; break; case HTP_TYPE_Q5_K: matmul_job_func = hvx_mm_4d_repacked_q5_k; break; + case HTP_TYPE_Q3_K: matmul_job_func = hvx_mm_4d_repacked_q3_k; break; + case HTP_TYPE_Q2_K: matmul_job_func = hvx_mm_4d_repacked_q2_k; break; case HTP_TYPE_IQ4_NL: matmul_job_func = hvx_mm_4d_repacked_iq4nl; break; case HTP_TYPE_MXFP4: matmul_job_func = hvx_mm_4d_repacked_mxfp4; break; default: return HTP_STATUS_NO_SUPPORT; @@ -1697,6 +1760,8 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { case HTP_TYPE_Q8_0: matmul_job_func = hvx_mm_2d_repacked_q8_0; break; case HTP_TYPE_Q6_K: matmul_job_func = hvx_mm_2d_repacked_q6_k; break; case HTP_TYPE_Q5_K: matmul_job_func = hvx_mm_2d_repacked_q5_k; break; + case HTP_TYPE_Q3_K: matmul_job_func = hvx_mm_2d_repacked_q3_k; break; + case HTP_TYPE_Q2_K: matmul_job_func = hvx_mm_2d_repacked_q2_k; break; case HTP_TYPE_IQ4_NL: matmul_job_func = hvx_mm_2d_repacked_iq4nl; break; case HTP_TYPE_MXFP4: matmul_job_func = hvx_mm_2d_repacked_mxfp4; break; default: return HTP_STATUS_NO_SUPPORT; @@ -1713,6 +1778,8 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { case HTP_TYPE_Q8_0: matmul_job_func = hvx_mv_2d_repacked_q8_0; break; case HTP_TYPE_Q5_K: matmul_job_func = hvx_mv_2d_repacked_q5_k; break; case HTP_TYPE_Q6_K: matmul_job_func = hvx_mv_2d_repacked_q6_k; break; + case HTP_TYPE_Q3_K: matmul_job_func = hvx_mv_2d_repacked_q3_k; break; + case HTP_TYPE_Q2_K: matmul_job_func = hvx_mv_2d_repacked_q2_k; break; case HTP_TYPE_IQ4_NL: matmul_job_func = hvx_mv_2d_repacked_iq4nl; break; case HTP_TYPE_MXFP4: matmul_job_func = hvx_mv_2d_repacked_mxfp4; break; default: return HTP_STATUS_NO_SUPPORT; @@ -1758,7 +1825,7 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { if (src1_nrows < octx->n_threads && !is_batched) { n_quant_tasks = MIN(total_nb, octx->n_threads); - quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + quant_task_func = htp_mm_act_quant_block_func(src0->type); for (uint32_t ith = 0; ith < n_quant_tasks; ++ith) { uint32_t ib_first = (total_nb * ith) / n_quant_tasks; uint32_t ib_last = (total_nb * (ith + 1)) / n_quant_tasks; @@ -1769,7 +1836,7 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { } } else { n_quant_tasks = MIN(src1_nrows, octx->n_threads); - quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; + quant_task_func = htp_mm_act_quant_row_func(src0->type); } src1_row_size = htp_mm_weight_has_offset(src0->type) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); break; @@ -1849,7 +1916,7 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { work_queue_func_t q_func; if (cur_m_rows < octx->n_threads && (kparams->kernel_type == HTP_MM_KERNEL_HVX_QUANT_BLOCK || kparams->kernel_type == HTP_MM_KERNEL_HVX_QUANT_ROW)) { quant_tasks = MIN(total_nb, octx->n_threads); - q_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + q_func = htp_mm_act_quant_block_func(src0->type); for (uint32_t ith = 0; ith < quant_tasks; ++ith) { uint32_t ib_first = (total_nb * ith) / quant_tasks; uint32_t ib_last = (total_nb * (ith + 1)) / quant_tasks; @@ -2012,6 +2079,8 @@ DEQUANTIZE_WORKER_LOOP_IMPL(mxfp4) DEQUANTIZE_WORKER_LOOP_IMPL(q8_0) DEQUANTIZE_WORKER_LOOP_IMPL(q6_k) DEQUANTIZE_WORKER_LOOP_IMPL(q5_k) +DEQUANTIZE_WORKER_LOOP_IMPL(q3_k) +DEQUANTIZE_WORKER_LOOP_IMPL(q2_k) static void convert_f16_worker_loop(unsigned int n, unsigned int i, void *data) { tiled_dequantize_state_t *state = (tiled_dequantize_state_t *)data; @@ -2697,6 +2766,8 @@ static int hmx_mm_2d_f32(struct htp_context *ctx, case HTP_TYPE_Q8_0: dequant_worker_fn = dequantize_tiled_worker_loop_q8_0; break; case HTP_TYPE_Q5_K: dequant_worker_fn = dequantize_tiled_worker_loop_q5_k; break; case HTP_TYPE_Q6_K: dequant_worker_fn = dequantize_tiled_worker_loop_q6_k; break; + case HTP_TYPE_Q3_K: dequant_worker_fn = dequantize_tiled_worker_loop_q3_k; break; + case HTP_TYPE_Q2_K: dequant_worker_fn = dequantize_tiled_worker_loop_q2_k; break; case HTP_TYPE_F16: dequant_worker_fn = convert_f16_worker_loop; break; case HTP_TYPE_F32: dequant_worker_fn = quantize_f32_worker_loop; break; default: @@ -2963,6 +3034,8 @@ static int hmx_mm_nx_2d_f32(struct htp_ops_context * octx, const struct htp_mm_k case HTP_TYPE_Q8_0: dequant_worker_fn = dequantize_tiled_worker_loop_q8_0; break; case HTP_TYPE_Q5_K: dequant_worker_fn = dequantize_tiled_worker_loop_q5_k; break; case HTP_TYPE_Q6_K: dequant_worker_fn = dequantize_tiled_worker_loop_q6_k; break; + case HTP_TYPE_Q3_K: dequant_worker_fn = dequantize_tiled_worker_loop_q3_k; break; + case HTP_TYPE_Q2_K: dequant_worker_fn = dequantize_tiled_worker_loop_q2_k; break; case HTP_TYPE_F16: dequant_worker_fn = convert_f16_worker_loop; break; case HTP_TYPE_F32: dequant_worker_fn = quantize_f32_worker_loop; break; default: @@ -3559,6 +3632,8 @@ static int hmx_mm_id_2d_f32(struct htp_context *ctx, case HTP_TYPE_Q8_0: dequant_worker_fn = dequantize_tiled_worker_loop_q8_0; break; case HTP_TYPE_Q5_K: dequant_worker_fn = dequantize_tiled_worker_loop_q5_k; break; case HTP_TYPE_Q6_K: dequant_worker_fn = dequantize_tiled_worker_loop_q6_k; break; + case HTP_TYPE_Q3_K: dequant_worker_fn = dequantize_tiled_worker_loop_q3_k; break; + case HTP_TYPE_Q2_K: dequant_worker_fn = dequantize_tiled_worker_loop_q2_k; break; case HTP_TYPE_F16: dequant_worker_fn = convert_f16_worker_loop; break; case HTP_TYPE_F32: dequant_worker_fn = quantize_f32_worker_loop; break; default: @@ -3863,7 +3938,7 @@ static int hvx_mm_matmul_id( uint32_t n_quant_tasks = 1; if (act_nrows < octx->n_threads) { n_quant_tasks = MIN(total_nb, octx->n_threads); - quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + quant_task_func = htp_mm_act_quant_block_func(src0->type); for (uint32_t ith = 0; ith < n_quant_tasks; ++ith) { uint32_t ib_first = (total_nb * ith) / n_quant_tasks; uint32_t ib_last = (total_nb * (ith + 1)) / n_quant_tasks; @@ -3874,7 +3949,7 @@ static int hvx_mm_matmul_id( } } else { n_quant_tasks = MIN(act_nrows, octx->n_threads); - quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; + quant_task_func = htp_mm_act_quant_row_func(src0->type); } size_t src1_row_size = htp_mm_weight_has_offset(src0->type) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); @@ -4016,7 +4091,7 @@ static int hvx_mm_matmul_id_nx( uint32_t n_quant_tasks = 1; if (act_nrows < octx->n_threads) { n_quant_tasks = MIN(total_nb, octx->n_threads); - quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + quant_task_func = htp_mm_act_quant_block_func(src0->type); for (uint32_t ith = 0; ith < n_quant_tasks; ++ith) { uint32_t ib_first = (total_nb * ith) / n_quant_tasks; uint32_t ib_last = (total_nb * (ith + 1)) / n_quant_tasks; @@ -4027,7 +4102,7 @@ static int hvx_mm_matmul_id_nx( } } else { n_quant_tasks = MIN(act_nrows, octx->n_threads); - quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; + quant_task_func = htp_mm_act_quant_row_func(src0->type); } size_t src1_row_size = htp_mm_weight_has_offset(src0->type) ? htp_mm_q8_1_tiled_row_size(act->ne[0]) : htp_mm_q8_0_tiled_row_size(act->ne[0]); @@ -4394,7 +4469,8 @@ int op_matmul_nx(struct htp_ops_context * octx) { bool is_repacked = (src0->type == HTP_TYPE_Q4_0 || src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q8_0 || src0->type == HTP_TYPE_IQ4_NL || src0->type == HTP_TYPE_MXFP4 || src0->type == HTP_TYPE_Q4_K || - src0->type == HTP_TYPE_Q5_K); + src0->type == HTP_TYPE_Q5_K || src0->type == HTP_TYPE_Q3_K || + src0->type == HTP_TYPE_Q2_K); struct htp_mm_context mmctx_struct = {0}; struct htp_mm_context * mmctx = &mmctx_struct; @@ -4421,7 +4497,7 @@ int op_matmul_nx(struct htp_ops_context * octx) { uint32_t n_quant_tasks = 1; if (act_nrows < octx->n_threads) { n_quant_tasks = MIN(total_nb, octx->n_threads); - quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + quant_task_func = htp_mm_act_quant_block_func(src0->type); for (uint32_t ith = 0; ith < n_quant_tasks; ++ith) { uint32_t ib_first = (total_nb * ith) / n_quant_tasks; uint32_t ib_last = (total_nb * (ith + 1)) / n_quant_tasks; @@ -4432,7 +4508,7 @@ int op_matmul_nx(struct htp_ops_context * octx) { } } else { n_quant_tasks = MIN(act_nrows, octx->n_threads); - quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; + quant_task_func = htp_mm_act_quant_row_func(src0->type); } const size_t src1_row_size = htp_mm_weight_has_offset(src0->type) @@ -4480,6 +4556,8 @@ int op_matmul_nx(struct htp_ops_context * octx) { case HTP_TYPE_Q4_1: case HTP_TYPE_Q4_K: matmul_job_func = hvx_mm_nx_2d_repacked_q4_1; break; case HTP_TYPE_Q5_K: matmul_job_func = hvx_mm_nx_2d_repacked_q5_k; break; + case HTP_TYPE_Q3_K: matmul_job_func = hvx_mm_nx_2d_repacked_q3_k; break; + case HTP_TYPE_Q2_K: matmul_job_func = hvx_mm_nx_2d_repacked_q2_k; break; case HTP_TYPE_Q8_0: matmul_job_func = hvx_mm_nx_2d_repacked_q8_0; break; case HTP_TYPE_IQ4_NL: matmul_job_func = hvx_mm_nx_2d_repacked_iq4nl; break; case HTP_TYPE_MXFP4: matmul_job_func = hvx_mm_nx_2d_repacked_mxfp4; break; diff --git a/ggml/src/ggml-hexagon/htp/matmul-ops.h b/ggml/src/ggml-hexagon/htp/matmul-ops.h index cfb3bfbbd..386cb3049 100644 --- a/ggml/src/ggml-hexagon/htp/matmul-ops.h +++ b/ggml/src/ggml-hexagon/htp/matmul-ops.h @@ -33,6 +33,16 @@ extern "C" { // 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, vrmpy-ready like Q6_K +// vectors 0..1: unsigned 2-bit quants, vector m holds groups 4m..4m+3 at bit offsets 0,2,4,6 +// vector 2: fp16 scales per row, d * (scales[] & 0xF), same lanes as Q3_K vector 3 +// vector 3: fp16 offsets per row, -dmin * (scales[] >> 4), same lanes +#define HTP_MM_WEIGHT_TILE_SIZE_Q2_K 512 // --- Weight Repacked Aligned Tile Sizes --- #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q4_0 640 @@ -42,6 +52,8 @@ extern "C" { #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 512 // --- Activation Tiled Block Sizes (including padding) --- #define HTP_MM_ACT_TILE_SIZE_Q8_0 1152 @@ -207,6 +219,10 @@ static inline uint32_t htp_mm_get_weight_tile_size(int weight_type) { 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: @@ -228,6 +244,10 @@ static inline uint32_t htp_mm_get_weight_aligned_tile_size(int weight_type) { 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: @@ -236,8 +256,10 @@ static inline uint32_t htp_mm_get_weight_aligned_tile_size(int weight_type) { } // 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) { - return weight_type == HTP_TYPE_Q4_1 || weight_type == HTP_TYPE_Q4_K || weight_type == HTP_TYPE_Q5_K; + 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; } // --- Activation/Row Size Helpers --- @@ -263,6 +285,8 @@ static inline size_t htp_mm_get_tiled_row_stride(int weight_type, uint32_t 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: @@ -501,7 +525,8 @@ static inline void htp_mm_hvx_vtcm_layout_build( 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_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);