diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp index 9d3c97a998..581f07aa2e 100644 --- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp +++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp @@ -267,10 +267,10 @@ 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_Q4_K || type == GGML_TYPE_Q5_K; } -// Size of one repacked row in the DSP tiled layout. The Q6_K and Q4_K tiles store uncompressed scales/mins, +// Size of one repacked row in the DSP tiled layout. The Q6_K, Q5_K and Q4_K 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) { @@ -279,6 +279,9 @@ static inline size_t ggml_hexagon_tiled_row_size(enum ggml_type type, int64_t ne if (type == GGML_TYPE_Q4_K) { return (size_t) (ne0 / 32) * (HTP_MM_WEIGHT_TILE_SIZE_Q4_1 / 32); } + if (type == GGML_TYPE_Q5_K) { + return (size_t) (ne0 / 32) * (HTP_MM_WEIGHT_TILE_SIZE_Q5_K / 32); + } return ggml_row_size(type, ne0); } @@ -1742,6 +1745,202 @@ static void repack_tiled_q4_K(void * data, const ggml_tensor * t, size_t offset, GGML_UNUSED(size); } +// tile layout: see HTP_MM_WEIGHT_TILE_SIZE_Q5_K in htp/matmul-ops.h +static void repack_q5_K_tiled(ggml_tensor * t, const void * data, size_t offset, size_t size) { + GGML_ASSERT(offset == 0); + + const block_q5_K * src_matrix = (const block_q5_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_Q5_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_q5_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); + + for (int64_t r = 0; r < ne1; r++) { + const int ct = (int) (r / 32); + const int row = (int) (r % 32); + const block_q5_K * src_row = src_slice + r * sb_per_row; + + for (int kt = 0; kt < n_k_tiles; kt++) { + const int kt_local = kt % 8; + const block_q5_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_dst = matrix_dst + ((size_t) ct * n_k_tiles + kt) * tile_size; + uint8_t * plane = tile_dst + 640; + + uint8_t sc, m; + get_scale_min_k4(kt_local, b->scales, &sc, &m); + + const float D = d * (float) sc; + const float M = -dmin * (float) m; + + const uint8_t * qs_sub = b->qs + (kt_local / 2) * 32; + const int shift = (kt_local & 1) ? 4 : 0; + const uint8_t hbit = (uint8_t) (1 << kt_local); + + for (int cp = 0; cp < 16; cp++) { + const uint8_t q0 = (qs_sub[2 * cp + 0] >> shift) & 0x0F; + const uint8_t q1 = (qs_sub[2 * cp + 1] >> shift) & 0x0F; + tile_dst[cp * 32 + row] = (uint8_t) ((q1 << 4) | q0); + + const int i = cp / 4; + const int lane = (cp % 4) * 32 + row; + if (b->qh[2 * cp + 0] & hbit) { + plane[lane] |= (uint8_t) (1 << (2 * i)); + } + if (b->qh[2 * cp + 1] & hbit) { + plane[lane] |= (uint8_t) (1 << (2 * i + 1)); + } + } + + ggml_half * scale_dst = (ggml_half *) (tile_dst + 512); + scale_dst[2 * row + 0] = GGML_FP32_TO_FP16(D); + scale_dst[2 * row + 1] = GGML_FP32_TO_FP16(M); + } + } + } + } + + GGML_UNUSED(size); +} + +// Reverse of repack_q5_K_tiled. Unpacks quants losslessly and normalizes scales/mins. Read-back only. +static void repack_tiled_q5_K(void * data, const ggml_tensor * t, size_t offset, size_t size) { + GGML_ASSERT(offset == 0); + + block_q5_K * dst_matrix = (block_q5_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_Q5_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_q5_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_q5_K * dst_row = dst_slice + r * sb_per_row; + + for (int64_t sb = 0; sb < sb_per_row; sb++) { + block_q5_K * b = &dst_row[sb]; + memset(b, 0, sizeof(block_q5_K)); + + float sub_scales[8]; + float sub_mins[8]; + + for (int kt_local = 0; kt_local < 8; kt_local++) { + const int kt = sb * 8 + kt_local; + const uint8_t * tile_src = matrix_src + ((size_t) ct * n_k_tiles + kt) * tile_size; + const uint8_t * plane = tile_src + 640; + const ggml_half * scale_src = (const ggml_half *) (tile_src + 512); + + uint8_t * qs_sub = b->qs + (kt_local / 2) * 32; + const int shift = (kt_local & 1) ? 4 : 0; + const uint8_t hbit = (uint8_t) (1 << kt_local); + + for (int cp = 0; cp < 16; cp++) { + const uint8_t val = tile_src[cp * 32 + row]; + const uint8_t q0 = val & 0x0F; + const uint8_t q1 = val >> 4; + qs_sub[2 * cp + 0] |= (uint8_t) (q0 << shift); + qs_sub[2 * cp + 1] |= (uint8_t) (q1 << shift); + + const int i = cp / 4; + const int lane = (cp % 4) * 32 + row; + if (plane[lane] & (1 << (2 * i))) { + b->qh[2 * cp + 0] |= hbit; + } + if (plane[lane] & (1 << (2 * i + 1))) { + b->qh[2 * cp + 1] |= hbit; + } + } + + const float D = GGML_FP16_TO_FP32(scale_src[2 * row + 0]); + const float M = GGML_FP16_TO_FP32(scale_src[2 * row + 1]); + sub_scales[kt_local] = (D > 0.0f) ? D : 0.0f; + sub_mins[kt_local] = (-M > 0.0f) ? -M : 0.0f; + } + + float max_scale = 0.0f; + float max_min = 0.0f; + for (int j = 0; j < 8; j++) { + if (sub_scales[j] > max_scale) max_scale = sub_scales[j]; + if (sub_mins[j] > max_min) max_min = sub_mins[j]; + } + + float inv_scale = 0.0f; + if (max_scale > 0.0f) { + b->d = GGML_FP32_TO_FP16(max_scale / 63.0f); + const float d_actual = GGML_FP16_TO_FP32(b->d); + inv_scale = (d_actual > 0.0f) ? (1.0f / d_actual) : 0.0f; + } else { + b->d = GGML_FP32_TO_FP16(0.0f); + } + + float inv_min = 0.0f; + if (max_min > 0.0f) { + b->dmin = GGML_FP32_TO_FP16(max_min / 63.0f); + const float dmin_actual = GGML_FP16_TO_FP32(b->dmin); + inv_min = (dmin_actual > 0.0f) ? (1.0f / dmin_actual) : 0.0f; + } else { + b->dmin = GGML_FP32_TO_FP16(0.0f); + } + + for (int j = 0; j < 8; j++) { + uint8_t ls = (uint8_t) roundf(inv_scale * sub_scales[j]); + uint8_t lm = (uint8_t) roundf(inv_min * sub_mins[j]); + ls = (std::min)((uint8_t) 63, ls); + lm = (std::min)((uint8_t) 63, lm); + if (j < 4) { + b->scales[j] = ls; + b->scales[j + 4] = lm; + } else { + b->scales[j + 4] = (ls & 0xF) | ((lm & 0xF) << 4); + b->scales[j - 4] |= ((ls >> 4) << 6); + b->scales[j - 0] |= ((lm >> 4) << 6); + } + } + } + } + } + } + + GGML_UNUSED(size); +} + static void repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t size) { switch (tensor->type) { case GGML_TYPE_Q4_0: @@ -1768,6 +1967,10 @@ static void repack_tensor_tiled(ggml_tensor * tensor, const void * data, size_t repack_mxfp4_tiled(tensor, data, 0, size); break; + case GGML_TYPE_Q5_K: + repack_q5_K_tiled(tensor, data, 0, size); + break; + case GGML_TYPE_Q6_K: repack_q6_K_tiled(tensor, data, 0, size); break; @@ -1856,6 +2059,12 @@ static void ggml_backend_hexagon_buffer_get_tensor(ggml_backend_buffer_t buffer, repack_tiled_q4_K(data, tensor, offset, size); break; + case GGML_TYPE_Q5_K: + GGML_ASSERT(offset == 0); + GGML_ASSERT(offset + size <= ggml_nbytes(tensor)); + repack_tiled_q5_K(data, tensor, offset, size); + break; + case GGML_TYPE_Q8_0: GGML_ASSERT(offset == 0); GGML_ASSERT(offset + size <= ggml_nbytes(tensor)); @@ -1989,6 +2198,10 @@ static void ggml_backend_hexagon_buffer_get_tensor_2d(ggml_backend_buffer_t buff repack_tiled_q4_K(temp_buf.data(), tensor, offset, temp_size); break; + case GGML_TYPE_Q5_K: + repack_tiled_q5_K(temp_buf.data(), tensor, offset, temp_size); + break; + case GGML_TYPE_Q8_0: repack_tiled_q8_0(temp_buf.data(), tensor, offset, temp_size); break; @@ -5403,12 +5616,13 @@ static bool ggml_hexagon_supported_mul_mat(const struct ggml_hexagon_session * s case GGML_TYPE_IQ4_NL: case GGML_TYPE_MXFP4: case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: case GGML_TYPE_Q6_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_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) ? QK_K : 32)) { return false; } @@ -5486,12 +5700,13 @@ static bool ggml_hexagon_supported_mul_mat_id(const struct ggml_hexagon_session case GGML_TYPE_IQ4_NL: case GGML_TYPE_MXFP4: case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: case GGML_TYPE_Q6_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_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) ? QK_K : 32)) { return false; } @@ -7757,6 +7972,8 @@ static void ggml_hexagon_init(ggml_backend_reg * reg) { "please update hexagon_type to match ggml_type"); static_assert((unsigned int) HTP_TYPE_Q4_K == (unsigned int) GGML_TYPE_Q4_K, "please update hexagon_type to match ggml_type"); + static_assert((unsigned int) HTP_TYPE_Q5_K == (unsigned int) GGML_TYPE_Q5_K, + "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"); 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 b75f601ccd..5b7f340238 100644 --- a/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h +++ b/ggml/src/ggml-hexagon/htp/hmx-mm-kernels-tiled.h @@ -506,6 +506,111 @@ static void dequantize_tiled_weight_to_fp16_task_q8_0( } } +static void dequantize_tiled_weight_to_fp16_task_q5_k( + 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 vscale_offset = hvx_vmem(tile_src + 512); + HVX_VectorPair dm_deal = Q6_W_vdeal_VVR(vscale_offset, vscale_offset, -2); + HVX_Vector vd = Q6_V_lo_W(dm_deal); + HVX_Vector vm = Q6_V_hi_W(dm_deal); + + HVX_Vector v_scale_duplicated = Q6_V_lo_W(Q6_W_vshuff_VVR(vd, vd, -2)); + HVX_Vector v_offset_duplicated = Q6_V_lo_W(Q6_W_vshuff_VVR(vm, vm, -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_K: OR in the 5th bit from the plane + HVX_Vector v_plane = hvx_vmem(tile_src + 640); + v_lo0 = hvx_q5k_or_hibit(v_lo0, v_plane, 0); + v_hi0 = hvx_q5k_or_hibit(v_hi0, v_plane, 1); + v_lo1 = hvx_q5k_or_hibit(v_lo1, v_plane, 2); + v_hi1 = hvx_q5k_or_hibit(v_hi1, v_plane, 3); + v_lo2 = hvx_q5k_or_hibit(v_lo2, v_plane, 4); + v_hi2 = hvx_q5k_or_hibit(v_hi2, v_plane, 5); + v_lo3 = hvx_q5k_or_hibit(v_lo3, v_plane, 6); + v_hi3 = hvx_q5k_or_hibit(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, multiply, add offset + HVX_Vector v_grp0_0 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_lo0)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp0_1 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_lo0)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp0_2 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_hi0)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp0_3 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_hi0)), v_scale_duplicated), v_offset_duplicated)); + + HVX_Vector v_grp1_0 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_lo1)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp1_1 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_lo1)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp1_2 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_hi1)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp1_3 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_hi1)), v_scale_duplicated), v_offset_duplicated)); + + HVX_Vector v_grp2_0 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_lo2)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp2_1 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_lo2)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp2_2 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_hi2)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp2_3 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_hi2)), v_scale_duplicated), v_offset_duplicated)); + + HVX_Vector v_grp3_0 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_lo3)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp3_1 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_lo3)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp3_2 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_lo_W(vp_int16_hi3)), v_scale_duplicated), v_offset_duplicated)); + HVX_Vector v_grp3_3 = Q6_Vhf_equals_Vqf16(Q6_Vqf16_vadd_Vqf16Vhf(Q6_Vqf16_vmpy_VhfVhf(Q6_Vhf_equals_Vh(Q6_V_hi_W(vp_int16_hi3)), v_scale_duplicated), v_offset_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_MM_WEIGHT_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 ee5b924419..a03c8550f1 100644 --- a/ggml/src/ggml-hexagon/htp/htp-ops.h +++ b/ggml/src/ggml-hexagon/htp/htp-ops.h @@ -23,6 +23,7 @@ enum htp_data_type { HTP_TYPE_Q4_1 = 3, HTP_TYPE_Q8_0 = 8, HTP_TYPE_Q4_K = 12, + HTP_TYPE_Q5_K = 13, HTP_TYPE_Q6_K = 14, HTP_TYPE_IQ4_NL = 20, HTP_TYPE_I32 = 26, 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 5706259e16..4564b1524c 100644 --- a/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h +++ b/ggml/src/ggml-hexagon/htp/hvx-mm-kernels-tiled.h @@ -396,6 +396,78 @@ static inline HVX_VectorPair accum_q8_0_32x2( return Q6_W_vcombine_VV(v_sum1, v_sum0); } +// Q5_K: OR 0x10 into every lane of v whose flag j is set in the plane (see HTP_MM_WEIGHT_TILE_SIZE_Q5_K) +static inline HVX_Vector hvx_q5k_or_hibit(HVX_Vector v, HVX_Vector v_plane, int j) { + HVX_VectorPred q = Q6_Q_vand_VR(v_plane, 0x01010101u << j); + return Q6_V_vandor_VQR(v, q, 0x10101010); +} + +// 5-bit variant: the high bit comes from the plane, see hvx_q5k_or_hibit +static inline HVX_VectorPair unpack_and_interleave_5bit_x2(HVX_Vector v_src, HVX_Vector v_plane, int i, HVX_Vector mask_h4) { + HVX_Vector v_lo = hvx_q5k_or_hibit(Q6_V_vand_VV(v_src, mask_h4), v_plane, 2 * i); + HVX_Vector v_hi = hvx_q5k_or_hibit(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_5bit_32x1( + const HVX_Vector * restrict vptr, + const HVX_Vector * restrict v_act, + HVX_Vector i8 +) { + 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[5]; + + #pragma unroll + for (int i = 0; i < 4; i++) { + HVX_VectorPair v_W_pair = unpack_and_interleave_5bit_x2(vptr[i], v_plane, i, mask_h4); + HVX_Vector v_W0 = Q6_Vb_vsub_VbVb(Q6_V_lo_W(v_W_pair), i8); + HVX_Vector v_W1 = Q6_Vb_vsub_VbVb(Q6_V_hi_W(v_W_pair), i8); + 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_5bit_32x2( + const HVX_Vector * restrict vptr, + const HVX_Vector * restrict v_act0, + const HVX_Vector * restrict v_act1, + HVX_Vector i8 +) { + 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[5]; + + #pragma unroll + for (int i = 0; i < 4; i++) { + HVX_VectorPair v_W_pair = unpack_and_interleave_5bit_x2(vptr[i], v_plane, i, mask_h4); + HVX_Vector v_W0 = Q6_Vb_vsub_VbVb(Q6_V_lo_W(v_W_pair), i8); + HVX_Vector v_W1 = Q6_Vb_vsub_VbVb(Q6_V_hi_W(v_W_pair), i8); + + 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_MM_WEIGHT_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); @@ -714,6 +786,102 @@ static void tiled_vec_dot_q8_0_32x2(const uint32_t n, float * restrict s0, float } } +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; + + 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 * 1280); + + HVX_Vector v_sum = accum_5bit_32x1(vptr, v_act, Q6_V_vzero()); + HVX_Vector v_sum_sf = Q6_Vsf_equals_Vw(v_sum); + + HVX_Vector v_scale_offset = vptr[4]; + HVX_VectorPair p_deal = Q6_W_vdeal_VVR(v_scale_offset, v_scale_offset, -2); + HVX_Vector v_scale = Q6_V_lo_W(p_deal); + HVX_Vector v_offset = Q6_V_hi_W(p_deal); + + HVX_Vector v_scale_a = v_act[8]; + HVX_Vector v_sum_a = v_act[9]; + + HVX_Vector v_scale_comb = hvx_vec_mul_f16_f16_to_f32_lower32(v_scale, v_scale_a); + HVX_Vector v_offset_comb = hvx_vec_mul_f16_f16_to_f32_lower32(v_offset, v_sum_a); + + HVX_Vector v_scaled_dot = hvx_vec_mul_f32_f32(v_sum_sf, v_scale_comb); + HVX_Vector v_sum_scaled = hvx_vec_add_f32_f32(v_scaled_dot, v_offset_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_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 * 768); + 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_sums = accum_5bit_32x2(vptr, v_act0, v_act1, Q6_V_vzero()); + 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_offset = vptr[4]; + HVX_VectorPair p_deal = Q6_W_vdeal_VVR(v_scale_offset, v_scale_offset, -2); + HVX_Vector v_scale = Q6_V_lo_W(p_deal); + HVX_Vector v_offset = Q6_V_hi_W(p_deal); + + HVX_Vector v_scale_a_c0 = v_act0[8]; + HVX_Vector v_sum_a_c0 = v_act0[9]; + HVX_Vector v_scale_a_c1 = v_act1[8]; + HVX_Vector v_sum_a_c1 = v_act1[9]; + + HVX_Vector v_scale_comb_c0 = hvx_vec_mul_f16_f16_to_f32_lower32(v_scale, v_scale_a_c0); + HVX_Vector v_offset_comb_c0 = hvx_vec_mul_f16_f16_to_f32_lower32(v_offset, v_sum_a_c0); + HVX_Vector v_scale_comb_c1 = hvx_vec_mul_f16_f16_to_f32_lower32(v_scale, v_scale_a_c1); + HVX_Vector v_offset_comb_c1 = hvx_vec_mul_f16_f16_to_f32_lower32(v_offset, v_sum_a_c1); + + HVX_Vector v_scaled_dot_c0 = hvx_vec_mul_f32_f32(v_sum_sf_c0, v_scale_comb_c0); + HVX_Vector v_sum_scaled_c0 = hvx_vec_add_f32_f32(v_scaled_dot_c0, v_offset_comb_c0); + + HVX_Vector v_scaled_dot_c1 = hvx_vec_mul_f32_f32(v_sum_sf_c1, v_scale_comb_c1); + HVX_Vector v_sum_scaled_c1 = hvx_vec_add_f32_f32(v_scaled_dot_c1, v_offset_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_q6_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 2b45112b8c..727dc281f7 100644 --- a/ggml/src/ggml-hexagon/htp/matmul-ops.c +++ b/ggml/src/ggml-hexagon/htp/matmul-ops.c @@ -544,6 +544,7 @@ MATMUL_2D_REPACKED_IMPL(q4_0, 576, tiled_vec_dot_q4_0_32x2, tiled_vec_do MATMUL_2D_REPACKED_IMPL(q4_1, 640, tiled_vec_dot_q4_1_32x2, tiled_vec_dot_q4_1_32x1) 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(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) @@ -714,6 +715,7 @@ static void quantize_f32_q8_1_tiled_block(unsigned int nth, unsigned int ith, vo 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(iq4nl, 576, tiled_vec_dot_iq4nl_32x1) MATVEC_2D_REPACKED_IMPL(mxfp4, 544, tiled_vec_dot_mxfp4_32x1) @@ -723,6 +725,7 @@ MATMUL_NX_2D_REPACKED_IMPL(q4_1, 640, tiled_vec_dot_q4_1_32x2, tiled_vec_do MATMUL_NX_2D_REPACKED_IMPL(q8_0, 1088, tiled_vec_dot_q8_0_32x2, tiled_vec_dot_q8_0_32x1) 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) #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) { \ @@ -855,6 +858,7 @@ MATMUL_4D_REPACKED_IMPL(q4_0, 576, tiled_vec_dot_q4_0_32x2, tiled_vec_do MATMUL_4D_REPACKED_IMPL(q4_1, 640, tiled_vec_dot_q4_1_32x2, tiled_vec_dot_q4_1_32x1) 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(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) @@ -1588,6 +1592,10 @@ static int hvx_mm_init_vec_dot(struct htp_mm_context * mmctx, enum htp_data_type mmctx->type = "q8_0_tiled-f32"; mmctx->vec_dot_32x1 = tiled_vec_dot_q8_0_32x1; return 0; + case HTP_TYPE_Q5_K: + mmctx->type = "q5_k_tiled-f32"; + mmctx->vec_dot_32x1 = tiled_vec_dot_q5_k_32x1; + return 0; case HTP_TYPE_Q6_K: mmctx->type = "q6_k_tiled-f32"; mmctx->vec_dot_32x1 = tiled_vec_dot_q6_k_32x1; @@ -1643,7 +1651,7 @@ 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_Q4_K || src0->type == HTP_TYPE_Q5_K); // Compute src0_nrows_per_thread mmctx->src0_nrows_per_thread = fastdiv(nrows + octx->n_threads - 1, &octx->n_threads_div); @@ -1672,6 +1680,7 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { case HTP_TYPE_Q4_K: matmul_job_func = hvx_mm_4d_repacked_q4_1; break; 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_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; @@ -1687,6 +1696,7 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { case HTP_TYPE_Q4_K: matmul_job_func = hvx_mm_2d_repacked_q4_1; break; 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_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; @@ -1701,6 +1711,7 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { case HTP_TYPE_Q4_1: case HTP_TYPE_Q4_K: matmul_job_func = hvx_mv_2d_repacked_q4_1; break; 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_IQ4_NL: matmul_job_func = hvx_mv_2d_repacked_iq4nl; break; case HTP_TYPE_MXFP4: matmul_job_func = hvx_mv_2d_repacked_mxfp4; break; @@ -1747,7 +1758,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 = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; 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; @@ -1758,9 +1769,9 @@ static int hvx_mm_matmul(struct htp_ops_context * octx) { } } else { n_quant_tasks = MIN(src1_nrows, octx->n_threads); - quant_task_func = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; + quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; } - src1_row_size = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); + 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; } @@ -1838,7 +1849,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 = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + q_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; 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; @@ -2000,6 +2011,7 @@ DEQUANTIZE_WORKER_LOOP_IMPL(iq4_nl) DEQUANTIZE_WORKER_LOOP_IMPL(mxfp4) DEQUANTIZE_WORKER_LOOP_IMPL(q8_0) DEQUANTIZE_WORKER_LOOP_IMPL(q6_k) +DEQUANTIZE_WORKER_LOOP_IMPL(q5_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; @@ -2683,6 +2695,7 @@ static int hmx_mm_2d_f32(struct htp_context *ctx, case HTP_TYPE_Q4_K: dequant_worker_fn = dequantize_tiled_worker_loop_q4_1; break; case HTP_TYPE_MXFP4: dequant_worker_fn = dequantize_tiled_worker_loop_mxfp4; break; 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_F16: dequant_worker_fn = convert_f16_worker_loop; break; case HTP_TYPE_F32: dequant_worker_fn = quantize_f32_worker_loop; break; @@ -2948,6 +2961,7 @@ static int hmx_mm_nx_2d_f32(struct htp_ops_context * octx, const struct htp_mm_k case HTP_TYPE_Q4_K: dequant_worker_fn = dequantize_tiled_worker_loop_q4_1; break; case HTP_TYPE_MXFP4: dequant_worker_fn = dequantize_tiled_worker_loop_mxfp4; break; 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_F16: dequant_worker_fn = convert_f16_worker_loop; break; case HTP_TYPE_F32: dequant_worker_fn = quantize_f32_worker_loop; break; @@ -3543,6 +3557,7 @@ static int hmx_mm_id_2d_f32(struct htp_context *ctx, case HTP_TYPE_Q4_K: dequant_worker_fn = dequantize_tiled_worker_loop_q4_1; break; case HTP_TYPE_MXFP4: dequant_worker_fn = dequantize_tiled_worker_loop_mxfp4; break; 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_F16: dequant_worker_fn = convert_f16_worker_loop; break; case HTP_TYPE_F32: dequant_worker_fn = quantize_f32_worker_loop; break; @@ -3848,7 +3863,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 = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; 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; @@ -3859,9 +3874,9 @@ static int hvx_mm_matmul_id( } } else { n_quant_tasks = MIN(act_nrows, octx->n_threads); - quant_task_func = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; + quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; } - size_t src1_row_size = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); + 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); struct htp_mm_hvx_vtcm_layout L; htp_mm_hvx_vtcm_layout_build(&L, kparams->kernel_type, src0->type, ne10, act_nrows, octx->n_threads, @@ -4001,7 +4016,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 = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; 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; @@ -4012,9 +4027,9 @@ static int hvx_mm_matmul_id_nx( } } else { n_quant_tasks = MIN(act_nrows, octx->n_threads); - quant_task_func = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; + quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; } - size_t src1_row_size = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(act->ne[0]) : htp_mm_q8_0_tiled_row_size(act->ne[0]); + 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]); struct htp_mm_hvx_vtcm_layout L; htp_mm_hvx_vtcm_layout_build(&L, kparams->kernel_type, src0->type, act->ne[0], act_nrows, octx->n_threads, @@ -4378,7 +4393,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_MXFP4 || src0->type == HTP_TYPE_Q4_K || + src0->type == HTP_TYPE_Q5_K); struct htp_mm_context mmctx_struct = {0}; struct htp_mm_context * mmctx = &mmctx_struct; @@ -4405,7 +4421,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 = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; + quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled_block : quantize_f32_q8_0_tiled_block; 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; @@ -4416,10 +4432,10 @@ int op_matmul_nx(struct htp_ops_context * octx) { } } else { n_quant_tasks = MIN(act_nrows, octx->n_threads); - quant_task_func = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; + quant_task_func = htp_mm_weight_has_offset(src0->type) ? quantize_f32_q8_1_tiled : quantize_f32_q8_0_tiled; } - const size_t src1_row_size = (src0->type == HTP_TYPE_Q4_1 || src0->type == HTP_TYPE_Q4_K) + const 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]); @@ -4463,6 +4479,7 @@ int op_matmul_nx(struct htp_ops_context * octx) { case HTP_TYPE_Q4_0: matmul_job_func = hvx_mm_nx_2d_repacked_q4_0; break; 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_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 49a76839ac..cfb3bfbbd3 100644 --- a/ggml/src/ggml-hexagon/htp/matmul-ops.h +++ b/ggml/src/ggml-hexagon/htp/matmul-ops.h @@ -25,6 +25,9 @@ extern "C" { #define HTP_MM_WEIGHT_TILE_SIZE_Q8_0 1088 #define HTP_MM_WEIGHT_TILE_SIZE_IQ4_NL 576 #define HTP_MM_WEIGHT_TILE_SIZE_MXFP4 544 +// Q5_K: the Q4_1 tile (640) followed by a 128-byte plane with the 5th bit of every quant, transposed so that +// plane byte l holds the eight flags of lane l: bit 2i = low nibble of nibble vector i, bit 2i+1 = high nibble +#define HTP_MM_WEIGHT_TILE_SIZE_Q5_K 768 // Q6_K native 6-bit tile (32 rows x 32 k), vrmpy-ready: byte 4*row+b of a vector holds k = 4*group+b // vectors 0..3: low nibbles, vector i holds group 2i (low nibble) and group 2i+1 (high nibble) // vectors 4..5: high 2 bits, vector m holds groups 4m..4m+3 at bit offsets 0,2,4,6 @@ -37,6 +40,7 @@ extern "C" { #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q8_0 1152 #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_IQ4_NL 640 #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_MXFP4 640 +#define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q5_K 768 #define HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q6_K 896 // --- Activation Tiled Block Sizes (including padding) --- @@ -199,6 +203,8 @@ static inline uint32_t htp_mm_get_weight_tile_size(int weight_type) { return HTP_MM_WEIGHT_TILE_SIZE_Q4_1; case HTP_TYPE_Q8_0: return HTP_MM_WEIGHT_TILE_SIZE_Q8_0; + case HTP_TYPE_Q5_K: + return HTP_MM_WEIGHT_TILE_SIZE_Q5_K; case HTP_TYPE_Q6_K: return HTP_MM_WEIGHT_TILE_SIZE_Q6_K; case HTP_TYPE_MXFP4: @@ -218,6 +224,8 @@ static inline uint32_t htp_mm_get_weight_aligned_tile_size(int weight_type) { return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q4_1; case HTP_TYPE_Q8_0: return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q8_0; + case HTP_TYPE_Q5_K: + return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q5_K; case HTP_TYPE_Q6_K: return HTP_MM_WEIGHT_ALIGNED_TILE_SIZE_Q6_K; case HTP_TYPE_MXFP4: @@ -227,6 +235,11 @@ 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) +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; +} + // --- Activation/Row Size Helpers --- static inline size_t htp_mm_q8_0_tiled_row_size(uint32_t ne) { const uint32_t ne_padded = ((ne + 127) / 128) * 128; @@ -248,6 +261,7 @@ static inline size_t htp_mm_get_tiled_row_stride(int weight_type, uint32_t k) { case HTP_TYPE_Q4_1: case HTP_TYPE_Q4_K: case HTP_TYPE_Q8_0: + case HTP_TYPE_Q5_K: case HTP_TYPE_Q6_K: case HTP_TYPE_MXFP4: return (size_t) nb * htp_mm_get_weight_tile_size(weight_type); @@ -487,7 +501,7 @@ 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_Q4_K || wtype == HTP_TYPE_Q5_K); if (is_fused_nx) { const size_t src0_row_size_padded = hex_round_up(src0_row_size, 128); @@ -504,7 +518,7 @@ static inline void htp_mm_hvx_vtcm_layout_build( weight_sz_per_thread = hex_round_up(n_prefetch * src0_row_size_padded, 128); } - size_t tiled_act_row_size = (wtype == HTP_TYPE_Q4_1 || wtype == HTP_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); + size_t tiled_act_row_size = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); size_t act_sz = hex_round_up(tiled_act_row_size * src1_nrows, 128); size_t raw_row_size = hex_round_up(ne10 * sizeof(float), QK_Q8_0_TILED * sizeof(float)); @@ -516,7 +530,7 @@ static inline void htp_mm_hvx_vtcm_layout_build( act_raw_sz = hex_round_up(raw_row_size * src1_nrows, 128); } else if (is_matmul_id) { const size_t src0_row_size_padded = htp_mm_round_up(src0_row_size, 128); - const size_t src1_row_size_tiled = (wtype == HTP_TYPE_Q4_1 || wtype == HTP_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(ne10) + const size_t src1_row_size_tiled = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); size_t src0_sz_per_thread = htp_mm_round_up(n_prefetch * src0_row_size_padded, 256); @@ -560,7 +574,7 @@ static inline void htp_mm_hvx_vtcm_layout_build( } case HTP_MM_KERNEL_HVX_QUANT_BLOCK: case HTP_MM_KERNEL_HVX_QUANT_ROW: { - size_t q_src1_row_size = (wtype == HTP_TYPE_Q4_1 || wtype == HTP_TYPE_Q4_K) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); + size_t q_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); src0_sz = htp_mm_round_up(n_prefetch * src0_row_size_padded, 256); src1_sz = htp_mm_round_up(q_src1_row_size * src1_nrows, 256); @@ -647,7 +661,7 @@ static inline bool htp_mm_hvx_solve_vtcm_params( const size_t avail_act = vtcm_budget - fixed_bytes; size_t row_size = 0; if (kernel_type == HTP_MM_KERNEL_HVX_QUANT_ROW || kernel_type == HTP_MM_KERNEL_HVX_QUANT_BLOCK) { - row_size = (wtype == HTP_TYPE_Q4_1 || wtype == HTP_TYPE_Q4_K) + row_size = htp_mm_weight_has_offset(wtype) ? htp_mm_q8_1_tiled_row_size(ne10) : htp_mm_q8_0_tiled_row_size(ne10); } else if (kernel_type == HTP_MM_KERNEL_HVX_F16_F16_VTCM) {