mirror of
https://github.com/ggml-org/whisper.cpp.git
synced 2026-10-12 09:15:38 +02:00
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 <maxk@qti.qualcomm.com>
This commit is contained in:
Jhen-Jie Hong
co-authored by
Max Krasnyansky
Max Krasnyansky
parent
396f68d33e
commit
0295ef6846
@@ -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");
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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++;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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);
|
||||
|
||||
Reference in New Issue
Block a user