nikola/llama_cpp
1
0
Fork 0
Code Issues Pull Requests Actions 23 Packages Projects Releases Wiki Activity

SYCL: use native subgroup size for K-quant DMMV (#21700)

Browse Source
This commit is contained in:
Katostrofik
2026-06-15 10:10:53 +03:00
committed by GitHub
parent d8a3f523c8
commit e3bb1add8c
1 changed files with 318 additions and 333 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-sycl/dmmv.cpp
+122 -137
View File
@@ -256,13 +256,6 @@ static void convert_mul_mat_vec_bf16_sycl(const void *vx, const dfloat *y,
} }
#endif #endif
/*
DPCT1110:4: The total declared local variable size in device function
dequantize_mul_mat_vec_q2_k exceeds 128 bytes and may cause high register
pressure. Consult with your hardware vendor to find the total register size
available and adjust the code, or use smaller sub-group size to avoid high
register pressure.
*/
static void dequantize_mul_mat_vec_q2_k(const void *__restrict__ vx, static void dequantize_mul_mat_vec_q2_k(const void *__restrict__ vx,
const float *__restrict__ yy, const float *__restrict__ yy,
float *__restrict__ dst, float *__restrict__ dst,
@@ -284,19 +277,15 @@ static void dequantize_mul_mat_vec_q2_k(const void *__restrict__ vx,
#if QK_K == 256 #if QK_K == 256
const int tid = const int tid =
item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...31 or 0...15 item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...7 or 0...15
const int ix = const int ix =
item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1 item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1
const int step = 16/K_QUANTS_PER_ITERATION; const int step = 16/K_QUANTS_PER_ITERATION;
const int im = tid/step; // 0 or 1. 0 computes 0..., 1 computes 128... const int in = tid % step; // 0...15 or 0...7
const int in = tid - step*im; // 0...15 or 0...7
const int l0 = K_QUANTS_PER_ITERATION*in; // 0...15 or 0...14 in steps of 2 const int l0 = K_QUANTS_PER_ITERATION*in; // 0...15 or 0...14 in steps of 2
const int q_offset = 32*im + l0;
const int s_offset = 8*im;
const int y_offset = 128*im + l0;
uint32_t aux[4]; uint32_t aux[4];
const uint8_t * d = (const uint8_t *)aux; const uint8_t * d = (const uint8_t *)aux;
@@ -304,12 +293,17 @@ static void dequantize_mul_mat_vec_q2_k(const void *__restrict__ vx,
for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) { for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) {
const float * y = yy + i * QK_K + y_offset;
const uint8_t * q = x[i].qs + q_offset;
const float dall = x[i].dm[0]; const float dall = x[i].dm[0];
const float dmin = x[i].dm[1]; const float dmin = x[i].dm[1];
for (int im = 0; im < 2; ++im) {
const int q_offset = 32*im + l0;
const int s_offset = 8*im;
const int y_offset = 128*im + l0;
const float * y = yy + i * QK_K + y_offset;
const uint8_t * q = x[i].qs + q_offset;
const uint32_t * a = (const uint32_t *)(x[i].scales + s_offset); const uint32_t * a = (const uint32_t *)(x[i].scales + s_offset);
aux[0] = a[0] & 0x0f0f0f0f; aux[0] = a[0] & 0x0f0f0f0f;
aux[1] = a[1] & 0x0f0f0f0f; aux[1] = a[1] & 0x0f0f0f0f;
@@ -331,6 +325,7 @@ static void dequantize_mul_mat_vec_q2_k(const void *__restrict__ vx,
} }
tmp += dall * sum1 - dmin * sum2; tmp += dall * sum1 - dmin * sum2;
}
} }
#else #else
@@ -372,7 +367,7 @@ static void dequantize_mul_mat_vec_q2_k(const void *__restrict__ vx,
// sum up partial sums and write back result // sum up partial sums and write back result
#pragma unroll #pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp += tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
} }
@@ -382,13 +377,6 @@ static void dequantize_mul_mat_vec_q2_k(const void *__restrict__ vx,
} }
} }
/*
DPCT1110:5: The total declared local variable size in device function
dequantize_mul_mat_vec_q3_k exceeds 128 bytes and may cause high register
pressure. Consult with your hardware vendor to find the total register size
available and adjust the code, or use smaller sub-group size to avoid high
register pressure.
*/
static void dequantize_mul_mat_vec_q3_k(const void *__restrict__ vx, static void dequantize_mul_mat_vec_q3_k(const void *__restrict__ vx,
const float *__restrict__ yy, const float *__restrict__ yy,
float *__restrict__ dst, float *__restrict__ dst,
@@ -412,31 +400,32 @@ static void dequantize_mul_mat_vec_q3_k(const void *__restrict__ vx,
const uint16_t kmask2 = 0x0f0f; const uint16_t kmask2 = 0x0f0f;
const int tid = const int tid =
item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...31 or 0...16 item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...7 or 0...15
const int ix = const int ix =
item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1 item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1
const int n = K_QUANTS_PER_ITERATION; // iterations in the inner loop const int n = K_QUANTS_PER_ITERATION; // iterations in the inner loop
const int step = 16/K_QUANTS_PER_ITERATION; const int step = 16/K_QUANTS_PER_ITERATION;
const int im = tid/step; // 0 or 1. 0 computes 0..., 1 computes 128... const int in = tid % step; // 0...15 or 0...7
const int in = tid - step*im; // 0....15 or 0...7
const uint8_t m = 1 << (4*im);
const int l0 = n*in; // 0...15 or 0...14 in steps of 2 const int l0 = n*in; // 0...15 or 0...14 in steps of 2
const int q_offset = 32*im + l0;
const int y_offset = 128*im + l0;
uint16_t utmp[4]; uint16_t utmp[4];
const int8_t * s = (const int8_t *)utmp; const int8_t * s = (const int8_t *)utmp;
const uint16_t s_shift = 4*im;
for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) { for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) {
const uint8_t * h = x[i].hmask + l0;
const float d = x[i].d;
for (int im = 0; im < 2; ++im) {
const int q_offset = 32*im + l0;
const int y_offset = 128*im + l0;
const uint16_t s_shift = 4*im;
const uint8_t m = 1 << (4*im);
const float * y = yy + i * QK_K + y_offset; const float * y = yy + i * QK_K + y_offset;
const uint8_t * q = x[i].qs + q_offset; const uint8_t * q = x[i].qs + q_offset;
const uint8_t * h = x[i].hmask + l0;
const uint16_t * a = (const uint16_t *)x[i].scales; const uint16_t * a = (const uint16_t *)x[i].scales;
utmp[0] = ((a[0] >> s_shift) & kmask2) | (((a[4] >> (s_shift + 0)) & kmask1) << 4); utmp[0] = ((a[0] >> s_shift) & kmask2) | (((a[4] >> (s_shift + 0)) & kmask1) << 4);
@@ -444,8 +433,6 @@ static void dequantize_mul_mat_vec_q3_k(const void *__restrict__ vx,
utmp[2] = ((a[2] >> s_shift) & kmask2) | (((a[4] >> (s_shift + 2)) & kmask1) << 4); utmp[2] = ((a[2] >> s_shift) & kmask2) | (((a[4] >> (s_shift + 2)) & kmask1) << 4);
utmp[3] = ((a[3] >> s_shift) & kmask2) | (((a[5] >> (s_shift + 2)) & kmask1) << 4); utmp[3] = ((a[3] >> s_shift) & kmask2) | (((a[5] >> (s_shift + 2)) & kmask1) << 4);
const float d = x[i].d;
float sum = 0; float sum = 0;
for (int l = 0; l < n; ++l) { for (int l = 0; l < n; ++l) {
sum += y[l+ 0] * (s[0] - 32) * (((q[l] >> 0) & 3) - (h[l] & (m << 0) ? 0 : 4)) sum += y[l+ 0] * (s[0] - 32) * (((q[l] >> 0) & 3) - (h[l] & (m << 0) ? 0 : 4))
@@ -458,6 +445,7 @@ static void dequantize_mul_mat_vec_q3_k(const void *__restrict__ vx,
+ y[l+112] * (s[7] - 32) * (((q[l+16] >> 6) & 3) - (h[l+16] & (m << 3) ? 0 : 4)); + y[l+112] * (s[7] - 32) * (((q[l+16] >> 6) & 3) - (h[l+16] & (m << 3) ? 0 : 4));
} }
tmp += d * sum; tmp += d * sum;
}
} }
#else #else
@@ -491,7 +479,7 @@ static void dequantize_mul_mat_vec_q3_k(const void *__restrict__ vx,
// sum up partial sums and write back result // sum up partial sums and write back result
#pragma unroll #pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp += tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
} }
@@ -530,32 +518,34 @@ static void dequantize_mul_mat_vec_q3_k_reorder(const void *__restrict__ vx,
const uint16_t kmask2 = 0x0f0f; const uint16_t kmask2 = 0x0f0f;
const int tid = const int tid =
item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...31 or 0...16 item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...7 or 0...15
const int ix = const int ix =
item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1 item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1
const int n = K_QUANTS_PER_ITERATION; // iterations in the inner loop const int n = K_QUANTS_PER_ITERATION; // iterations in the inner loop
const int step = 16/K_QUANTS_PER_ITERATION; const int step = 16/K_QUANTS_PER_ITERATION;
const int im = tid/step; // 0 or 1. 0 computes 0..., 1 computes 128... const int in = tid % step; // 0...15 or 0...7
const int in = tid - step*im; // 0....15 or 0...7
const uint8_t m = 1 << (4*im);
const int l0 = n*in; // 0...15 or 0...14 in steps of 2 const int l0 = n*in; // 0...15 or 0...14 in steps of 2
const int q_offset = 32*im + l0;
const int y_offset = 128*im + l0;
uint16_t utmp[4]; uint16_t utmp[4];
const int8_t * s = (const int8_t *)utmp; const int8_t * s = (const int8_t *)utmp;
const uint16_t s_shift = 4*im;
for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) { for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) {
const int bi = ib0 + i; const int bi = ib0 + i;
const uint8_t * h = hmask_base + bi * (QK_K / 8) + l0;
const float d = d_base[bi];
for (int im = 0; im < 2; ++im) {
const int q_offset = 32*im + l0;
const int y_offset = 128*im + l0;
const uint16_t s_shift = 4*im;
const uint8_t m = 1 << (4*im);
const float * y = yy + i * QK_K + y_offset; const float * y = yy + i * QK_K + y_offset;
const uint8_t * q = qs_base + bi * (QK_K / 4) + q_offset; const uint8_t * q = qs_base + bi * (QK_K / 4) + q_offset;
const uint8_t * h = hmask_base + bi * (QK_K / 8) + l0;
const uint16_t * a = (const uint16_t *)(scales_base + bi * 12); const uint16_t * a = (const uint16_t *)(scales_base + bi * 12);
utmp[0] = ((a[0] >> s_shift) & kmask2) | (((a[4] >> (s_shift + 0)) & kmask1) << 4); utmp[0] = ((a[0] >> s_shift) & kmask2) | (((a[4] >> (s_shift + 0)) & kmask1) << 4);
@@ -563,8 +553,6 @@ static void dequantize_mul_mat_vec_q3_k_reorder(const void *__restrict__ vx,
utmp[2] = ((a[2] >> s_shift) & kmask2) | (((a[4] >> (s_shift + 2)) & kmask1) << 4); utmp[2] = ((a[2] >> s_shift) & kmask2) | (((a[4] >> (s_shift + 2)) & kmask1) << 4);
utmp[3] = ((a[3] >> s_shift) & kmask2) | (((a[5] >> (s_shift + 2)) & kmask1) << 4); utmp[3] = ((a[3] >> s_shift) & kmask2) | (((a[5] >> (s_shift + 2)) & kmask1) << 4);
const float d = d_base[bi];
float sum = 0; float sum = 0;
for (int l = 0; l < n; ++l) { for (int l = 0; l < n; ++l) {
sum += y[l+ 0] * (s[0] - 32) * (((q[l] >> 0) & 3) - (h[l] & (m << 0) ? 0 : 4)) sum += y[l+ 0] * (s[0] - 32) * (((q[l] >> 0) & 3) - (h[l] & (m << 0) ? 0 : 4))
@@ -578,6 +566,7 @@ static void dequantize_mul_mat_vec_q3_k_reorder(const void *__restrict__ vx,
} }
tmp += d * sum; tmp += d * sum;
} }
}
#else #else
GGML_UNUSED(vx); GGML_UNUSED(vx);
GGML_UNUSED(yy); GGML_UNUSED(yy);
@@ -588,7 +577,7 @@ static void dequantize_mul_mat_vec_q3_k_reorder(const void *__restrict__ vx,
// sum up partial sums and write back result // sum up partial sums and write back result
#pragma unroll #pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp += tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
} }
@@ -598,13 +587,6 @@ static void dequantize_mul_mat_vec_q3_k_reorder(const void *__restrict__ vx,
} }
} }
/*
DPCT1110:6: The total declared local variable size in device function
dequantize_mul_mat_vec_q4_k exceeds 128 bytes and may cause high register
pressure. Consult with your hardware vendor to find the total register size
available and adjust the code, or use smaller sub-group size to avoid high
register pressure.
*/
static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx, static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx,
const float *__restrict__ yy, const float *__restrict__ yy,
float *__restrict__ dst, float *__restrict__ dst,
@@ -625,22 +607,19 @@ static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx,
const uint16_t kmask3 = 0xc0c0; const uint16_t kmask3 = 0xc0c0;
const int tid = const int tid =
item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...31 or 0...16 item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...7 or 0...15
const int ix = const int ix =
item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1 item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1
const int step = 8/K_QUANTS_PER_ITERATION; // 8 or 4 const int step = 8/K_QUANTS_PER_ITERATION; // 8 or 4
const int il = tid/step; // 0...3 const int il_base = tid/step; // 0 or 1 (was 0...3)
const int ir = tid - step*il; // 0...7 or 0...3 const int ir = tid - step*il_base; // 0...7 or 0...3
const int n = 2 * K_QUANTS_PER_ITERATION; // 2 or 4 const int n = 2 * K_QUANTS_PER_ITERATION; // 2 or 4
const int im = il/2; // 0 or 1. 0 computes 0,32 + 128,160, 1 computes 64,96 + 192,224 const int in = il_base%2;
const int in = il%2;
const int l0 = n*(2*ir + in); const int l0 = n*(2*ir + in);
const int q_offset = 32*im + l0;
const int y_offset = 64*im + l0;
uint16_t aux[4]; uint16_t aux[4];
const uint8_t * sc = (const uint8_t *)aux; const uint8_t * sc = (const uint8_t *)aux;
@@ -657,12 +636,16 @@ static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx,
for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) { for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) {
const float * y1 = yy + i*QK_K + y_offset;
const float * y2 = y1 + 128;
const float dall = x[i].dm[0]; const float dall = x[i].dm[0];
const float dmin = x[i].dm[1]; const float dmin = x[i].dm[1];
for (int im = 0; im < 2; ++im) {
const int q_offset = 32*im + l0;
const int y_offset = 64*im + l0;
const float * y1 = yy + i*QK_K + y_offset;
const float * y2 = y1 + 128;
const uint16_t * a = (const uint16_t *)x[i].scales; const uint16_t * a = (const uint16_t *)x[i].scales;
aux[0] = a[im+0] & kmask1; aux[0] = a[im+0] & kmask1;
aux[1] = a[im+2] & kmask1; aux[1] = a[im+2] & kmask1;
@@ -706,6 +689,7 @@ static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx,
} }
tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin; tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin;
#endif #endif
}
} }
#else #else
@@ -741,7 +725,7 @@ static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx,
// sum up partial sums and write back result // sum up partial sums and write back result
#pragma unroll #pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp += tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
} }
@@ -776,22 +760,19 @@ static void dequantize_mul_mat_vec_q4_k_reorder(const void *__restrict__ vx,
const uint16_t kmask3 = 0xc0c0; const uint16_t kmask3 = 0xc0c0;
const int tid = const int tid =
item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...31 or 0...16 item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...7 or 0...15
const int ix = const int ix =
item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1 item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0,1
const int step = 8/K_QUANTS_PER_ITERATION; // 8 or 4 const int step = 8/K_QUANTS_PER_ITERATION; // 8 or 4
const int il = tid/step; // 0...3 const int il_base = tid/step; // 0 or 1 (was 0...3)
const int ir = tid - step*il; // 0...7 or 0...3 const int ir = tid - step*il_base; // 0...7 or 0...3
const int n = 2 * K_QUANTS_PER_ITERATION; // 2 or 4 const int n = 2 * K_QUANTS_PER_ITERATION; // 2 or 4
const int im = il/2; // 0 or 1. 0 computes 0,32 + 128,160, 1 computes 64,96 + 192,224 const int in = il_base%2;
const int in = il%2;
const int l0 = n*(2*ir + in); const int l0 = n*(2*ir + in);
const int q_offset = 32*im + l0;
const int y_offset = 64*im + l0;
uint16_t aux[4]; uint16_t aux[4];
const uint8_t * sc = (const uint8_t *)aux; const uint8_t * sc = (const uint8_t *)aux;
@@ -809,13 +790,17 @@ static void dequantize_mul_mat_vec_q4_k_reorder(const void *__restrict__ vx,
for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) { for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) {
const int bi = ib0 + i; const int bi = ib0 + i;
const float * y1 = yy + i*QK_K + y_offset;
const float * y2 = y1 + 128;
const sycl::half2 dm_val = dm_base[bi]; const sycl::half2 dm_val = dm_base[bi];
const float dall = dm_val[0]; const float dall = dm_val[0];
const float dmin = dm_val[1]; const float dmin = dm_val[1];
for (int im = 0; im < 2; ++im) {
const int q_offset = 32*im + l0;
const int y_offset = 64*im + l0;
const float * y1 = yy + i*QK_K + y_offset;
const float * y2 = y1 + 128;
const uint16_t * a = (const uint16_t *)(scales_base + bi * K_SCALE_SIZE); const uint16_t * a = (const uint16_t *)(scales_base + bi * K_SCALE_SIZE);
aux[0] = a[im+0] & kmask1; aux[0] = a[im+0] & kmask1;
aux[1] = a[im+2] & kmask1; aux[1] = a[im+2] & kmask1;
@@ -859,6 +844,7 @@ static void dequantize_mul_mat_vec_q4_k_reorder(const void *__restrict__ vx,
} }
tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin; tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin;
#endif #endif
}
} }
#else #else
@@ -897,7 +883,7 @@ static void dequantize_mul_mat_vec_q4_k_reorder(const void *__restrict__ vx,
// sum up partial sums and write back result // sum up partial sums and write back result
#pragma unroll #pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp += tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
} }
@@ -907,13 +893,6 @@ static void dequantize_mul_mat_vec_q4_k_reorder(const void *__restrict__ vx,
} }
} }
/*
DPCT1110:7: The total declared local variable size in device function
dequantize_mul_mat_vec_q5_k exceeds 128 bytes and may cause high register
pressure. Consult with your hardware vendor to find the total register size
available and adjust the code, or use smaller sub-group size to avoid high
register pressure.
*/
static void dequantize_mul_mat_vec_q5_k(const void *__restrict__ vx, static void dequantize_mul_mat_vec_q5_k(const void *__restrict__ vx,
const float *__restrict__ yy, const float *__restrict__ yy,
float *__restrict__ dst, float *__restrict__ dst,
@@ -933,22 +912,16 @@ static void dequantize_mul_mat_vec_q5_k(const void *__restrict__ vx,
const uint16_t kmask2 = 0x0f0f; const uint16_t kmask2 = 0x0f0f;
const uint16_t kmask3 = 0xc0c0; const uint16_t kmask3 = 0xc0c0;
const int tid = item_ct1.get_local_id(2) / 2; // 0...15 const int tid = item_ct1.get_local_id(2) / 2; // 0...7
const int ix = item_ct1.get_local_id(2) % 2; const int ix = item_ct1.get_local_id(2) % 2;
const int il = tid/4; // 0...3 const int il_base = tid/4; // 0 or 1 (was 0...3)
const int ir = tid - 4*il;// 0...3 const int ir = tid - 4*il_base;// 0...3
const int n = 2; const int n = 2;
const int im = il/2; // 0 or 1. 0 computes 0,32 + 128,160, 1 computes 64,96 + 192,224 const int in = il_base%2;
const int in = il%2;
const int l0 = n*(2*ir + in); const int l0 = n*(2*ir + in);
const int q_offset = 32*im + l0;
const int y_offset = 64*im + l0;
const uint8_t hm1 = 1 << (2*im);
const uint8_t hm2 = hm1 << 4;
uint16_t aux[4]; uint16_t aux[4];
const uint8_t * sc = (const uint8_t *)aux; const uint8_t * sc = (const uint8_t *)aux;
@@ -958,14 +931,21 @@ static void dequantize_mul_mat_vec_q5_k(const void *__restrict__ vx,
for (int i = ix; i < num_blocks_per_row; i += 2) { for (int i = ix; i < num_blocks_per_row; i += 2) {
const uint8_t * ql1 = x[i].qs + q_offset;
const uint8_t * qh = x[i].qh + l0; const uint8_t * qh = x[i].qh + l0;
const float * y1 = yy + i*QK_K + y_offset;
const float * y2 = y1 + 128;
const float dall = x[i].dm[0]; const float dall = x[i].dm[0];
const float dmin = x[i].dm[1]; const float dmin = x[i].dm[1];
for (int im = 0; im < 2; ++im) {
const int q_offset = 32*im + l0;
const int y_offset = 64*im + l0;
const uint8_t hm1 = 1 << (2*im);
const uint8_t hm2 = hm1 << 4;
const uint8_t * ql1 = x[i].qs + q_offset;
const float * y1 = yy + i*QK_K + y_offset;
const float * y2 = y1 + 128;
const uint16_t * a = (const uint16_t *)x[i].scales; const uint16_t * a = (const uint16_t *)x[i].scales;
aux[0] = a[im+0] & kmask1; aux[0] = a[im+0] & kmask1;
aux[1] = a[im+2] & kmask1; aux[1] = a[im+2] & kmask1;
@@ -1004,6 +984,7 @@ static void dequantize_mul_mat_vec_q5_k(const void *__restrict__ vx,
sum.w() * sc[5]) - sum.w() * sc[5]) -
dmin * smin; dmin * smin;
} }
}
#else #else
const int tid = item_ct1.get_local_id(2)/(2*K_QUANTS_PER_ITERATION); // 0...15 const int tid = item_ct1.get_local_id(2)/(2*K_QUANTS_PER_ITERATION); // 0...15
@@ -1031,7 +1012,7 @@ static void dequantize_mul_mat_vec_q5_k(const void *__restrict__ vx,
// sum up partial sums and write back result // sum up partial sums and write back result
#pragma unroll #pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp += tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
} }
@@ -1058,14 +1039,13 @@ static void dequantize_mul_mat_vec_q6_k(const void * __restrict__ vx, const floa
#if QK_K == 256 #if QK_K == 256
const int tid = const int tid =
item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...31 or 0...16 item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...7 or 0...15
const int ix = const int ix =
item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0, 1 item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0, 1
const int step = 16/K_QUANTS_PER_ITERATION; // 16 or 8 const int step = 16/K_QUANTS_PER_ITERATION; // 16 or 8
const int im = tid/step; // 0 or 1. 0 computes 0..., 1 computes 128... const int in = tid % step; // 0...15 or 0...7
const int in = tid - step*im; // 0...15 or 0...7
#if K_QUANTS_PER_ITERATION == 1 #if K_QUANTS_PER_ITERATION == 1
const int l0 = K_QUANTS_PER_ITERATION*in; // 0...15 const int l0 = K_QUANTS_PER_ITERATION*in; // 0...15
@@ -1074,22 +1054,24 @@ static void dequantize_mul_mat_vec_q6_k(const void * __restrict__ vx, const floa
const int l0 = 4 * in; // 0, 4, 8, ..., 28 const int l0 = 4 * in; // 0, 4, 8, ..., 28
const int is = in / 4; const int is = in / 4;
#endif #endif
const int ql_offset = 64*im + l0;
const int qh_offset = 32*im + l0;
const int s_offset = 8*im + is;
const int y_offset = 128*im + l0;
float tmp = 0; // partial sum for thread in warp float tmp = 0; // partial sum for thread in warp
for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) { for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) {
const float d = x[i].d;
for (int im = 0; im < 2; ++im) {
const int ql_offset = 64*im + l0;
const int qh_offset = 32*im + l0;
const int s_offset = 8*im + is;
const int y_offset = 128*im + l0;
const float * y = yy + i * QK_K + y_offset; const float * y = yy + i * QK_K + y_offset;
const uint8_t * ql = x[i].ql + ql_offset; const uint8_t * ql = x[i].ql + ql_offset;
const uint8_t * qh = x[i].qh + qh_offset; const uint8_t * qh = x[i].qh + qh_offset;
const int8_t * s = x[i].scales + s_offset; const int8_t * s = x[i].scales + s_offset;
const float d = x[i].d;
#if K_QUANTS_PER_ITERATION == 1 #if K_QUANTS_PER_ITERATION == 1
float sum = y[ 0] * s[0] * d * ((int8_t)((ql[ 0] & 0xF) | ((qh[ 0] & 0x03) << 4)) - 32) float sum = y[ 0] * s[0] * d * ((int8_t)((ql[ 0] & 0xF) | ((qh[ 0] & 0x03) << 4)) - 32)
+ y[16] * s[1] * d * ((int8_t)((ql[16] & 0xF) | ((qh[16] & 0x03) << 4)) - 32) + y[16] * s[1] * d * ((int8_t)((ql[16] & 0xF) | ((qh[16] & 0x03) << 4)) - 32)
@@ -1110,6 +1092,7 @@ static void dequantize_mul_mat_vec_q6_k(const void * __restrict__ vx, const floa
} }
tmp += sum; tmp += sum;
#endif #endif
}
} }
@@ -1146,7 +1129,7 @@ static void dequantize_mul_mat_vec_q6_k(const void * __restrict__ vx, const floa
// sum up partial sums and write back result // sum up partial sums and write back result
#pragma unroll #pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp += tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
} }
@@ -1179,14 +1162,13 @@ static void dequantize_mul_mat_vec_q6_k_reorder(const void * __restrict__ vx, co
#if QK_K == 256 #if QK_K == 256
const int tid = const int tid =
item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...31 or 0...16 item_ct1.get_local_id(2) / K_QUANTS_PER_ITERATION; // 0...7 or 0...15
const int ix = const int ix =
item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0, 1 item_ct1.get_local_id(2) % K_QUANTS_PER_ITERATION; // 0 or 0, 1
const int step = 16/K_QUANTS_PER_ITERATION; // 16 or 8 const int step = 16/K_QUANTS_PER_ITERATION; // 16 or 8
const int im = tid/step; // 0 or 1. 0 computes 0..., 1 computes 128... const int in = tid % step; // 0...15 or 0...7
const int in = tid - step*im; // 0...15 or 0...7
#if K_QUANTS_PER_ITERATION == 1 #if K_QUANTS_PER_ITERATION == 1
const int l0 = K_QUANTS_PER_ITERATION*in; // 0...15 const int l0 = K_QUANTS_PER_ITERATION*in; // 0...15
@@ -1195,23 +1177,25 @@ static void dequantize_mul_mat_vec_q6_k_reorder(const void * __restrict__ vx, co
const int l0 = 4 * in; // 0, 4, 8, ..., 28 const int l0 = 4 * in; // 0, 4, 8, ..., 28
const int is = in / 4; const int is = in / 4;
#endif #endif
const int ql_offset = 64*im + l0;
const int qh_offset = 32*im + l0;
const int s_offset = 8*im + is;
const int y_offset = 128*im + l0;
float tmp = 0; // partial sum for thread in warp float tmp = 0; // partial sum for thread in warp
for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) { for (int i = ix; i < num_blocks_per_row; i += K_QUANTS_PER_ITERATION) {
const int bi = ib0 + i; const int bi = ib0 + i;
const float d = d_base[bi];
for (int im = 0; im < 2; ++im) {
const int ql_offset = 64*im + l0;
const int qh_offset = 32*im + l0;
const int s_offset = 8*im + is;
const int y_offset = 128*im + l0;
const float * y = yy + i * QK_K + y_offset; const float * y = yy + i * QK_K + y_offset;
const uint8_t * ql = ql_base + bi * (QK_K / 2) + ql_offset; const uint8_t * ql = ql_base + bi * (QK_K / 2) + ql_offset;
const uint8_t * qh = qh_base + bi * (QK_K / 4) + qh_offset; const uint8_t * qh = qh_base + bi * (QK_K / 4) + qh_offset;
const int8_t * s = scales_base + bi * (QK_K / 16) + s_offset; const int8_t * s = scales_base + bi * (QK_K / 16) + s_offset;
const float d = d_base[bi];
#if K_QUANTS_PER_ITERATION == 1 #if K_QUANTS_PER_ITERATION == 1
float sum = y[ 0] * s[0] * d * ((int8_t)((ql[ 0] & 0xF) | ((qh[ 0] & 0x03) << 4)) - 32) float sum = y[ 0] * s[0] * d * ((int8_t)((ql[ 0] & 0xF) | ((qh[ 0] & 0x03) << 4)) - 32)
+ y[16] * s[1] * d * ((int8_t)((ql[16] & 0xF) | ((qh[16] & 0x03) << 4)) - 32) + y[16] * s[1] * d * ((int8_t)((ql[16] & 0xF) | ((qh[16] & 0x03) << 4)) - 32)
@@ -1232,6 +1216,7 @@ static void dequantize_mul_mat_vec_q6_k_reorder(const void * __restrict__ vx, co
} }
tmp += sum; tmp += sum;
#endif #endif
}
} }
@@ -1269,7 +1254,7 @@ static void dequantize_mul_mat_vec_q6_k_reorder(const void * __restrict__ vx, co
// sum up partial sums and write back result // sum up partial sums and write back result
#pragma unroll #pragma unroll
for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) {
tmp += tmp +=
dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
} }
@@ -1513,10 +1498,10 @@ static void dequantize_mul_mat_vec_q2_K_sycl(const void *vx, const float *y,
const int ny = 2; // very slightly faster than 1 even when K_QUANTS_PER_ITERATION = 2 const int ny = 2; // very slightly faster than 1 even when K_QUANTS_PER_ITERATION = 2
const int block_num_y = (nrows + ny - 1) / ny; const int block_num_y = (nrows + ny - 1) / ny;
const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE); const sycl::range<3> block_dims(1, ny, WARP_SIZE);
stream->parallel_for( stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims), sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q2_k(vx, y, dst, ncols, nrows, item_ct1); dequantize_mul_mat_vec_q2_k(vx, y, dst, ncols, nrows, item_ct1);
}); });
} }
@@ -1529,10 +1514,10 @@ static void dequantize_mul_mat_vec_q3_K_sycl(const void *vx, const float *y,
const int ny = 2 / K_QUANTS_PER_ITERATION; const int ny = 2 / K_QUANTS_PER_ITERATION;
const int block_num_y = (nrows + ny - 1) / ny; const int block_num_y = (nrows + ny - 1) / ny;
const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE); const sycl::range<3> block_dims(1, ny, WARP_SIZE);
stream->parallel_for( stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims), sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q3_k(vx, y, dst, ncols, nrows, item_ct1); dequantize_mul_mat_vec_q3_k(vx, y, dst, ncols, nrows, item_ct1);
}); });
} }
@@ -1545,10 +1530,10 @@ static void dequantize_mul_mat_vec_q3_K_sycl_reorder(const void *vx, const float
const int ny = 2 / K_QUANTS_PER_ITERATION; const int ny = 2 / K_QUANTS_PER_ITERATION;
const int block_num_y = (nrows + ny - 1) / ny; const int block_num_y = (nrows + ny - 1) / ny;
const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE); const sycl::range<3> block_dims(1, ny, WARP_SIZE);
stream->parallel_for( stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims), sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q3_k_reorder(vx, y, dst, ncols, nrows, item_ct1); dequantize_mul_mat_vec_q3_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
}); });
} }
@@ -1561,10 +1546,10 @@ static void dequantize_mul_mat_vec_q4_K_sycl(const void *vx, const float *y,
const int ny = 2 / K_QUANTS_PER_ITERATION; const int ny = 2 / K_QUANTS_PER_ITERATION;
const int block_num_y = (nrows + ny - 1) / ny; const int block_num_y = (nrows + ny - 1) / ny;
const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE); const sycl::range<3> block_dims(1, ny, WARP_SIZE);
stream->parallel_for( stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims), sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q4_k(vx, y, dst, ncols, nrows, item_ct1); dequantize_mul_mat_vec_q4_k(vx, y, dst, ncols, nrows, item_ct1);
}); });
} }
@@ -1574,10 +1559,10 @@ static void dequantize_mul_mat_vec_q5_K_sycl(const void *vx, const float *y,
const int nrows, const int nrows,
dpct::queue_ptr stream) { dpct::queue_ptr stream) {
GGML_ASSERT(ncols % QK_K == 0); GGML_ASSERT(ncols % QK_K == 0);
const sycl::range<3> block_dims(1, 1, QK_WARP_SIZE); const sycl::range<3> block_dims(1, 1, WARP_SIZE);
stream->parallel_for( stream->parallel_for(
sycl::nd_range<3>(sycl::range<3>(1, 1, nrows) * block_dims, block_dims), sycl::nd_range<3>(sycl::range<3>(1, 1, nrows) * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q5_k(vx, y, dst, ncols, item_ct1); dequantize_mul_mat_vec_q5_k(vx, y, dst, ncols, item_ct1);
}); });
} }
@@ -1590,10 +1575,10 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
const int ny = 2 / K_QUANTS_PER_ITERATION; const int ny = 2 / K_QUANTS_PER_ITERATION;
const int block_num_y = (nrows + ny - 1) / ny; const int block_num_y = (nrows + ny - 1) / ny;
const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE); const sycl::range<3> block_dims(1, ny, WARP_SIZE);
stream->parallel_for( stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims), sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q6_k(vx, y, dst, ncols, nrows, item_ct1); dequantize_mul_mat_vec_q6_k(vx, y, dst, ncols, nrows, item_ct1);
}); });
} }
@@ -1606,10 +1591,10 @@ static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float
const int ny = 2 / K_QUANTS_PER_ITERATION; const int ny = 2 / K_QUANTS_PER_ITERATION;
const int block_num_y = (nrows + ny - 1) / ny; const int block_num_y = (nrows + ny - 1) / ny;
const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE); const sycl::range<3> block_dims(1, ny, WARP_SIZE);
stream->parallel_for( stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims), sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q4_k_reorder(vx, y, dst, ncols, nrows, item_ct1); dequantize_mul_mat_vec_q4_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
}); });
} }
@@ -1622,10 +1607,10 @@ static void dequantize_mul_mat_vec_q6_K_sycl_reorder(const void *vx, const float
const int ny = 2 / K_QUANTS_PER_ITERATION; const int ny = 2 / K_QUANTS_PER_ITERATION;
const int block_num_y = (nrows + ny - 1) / ny; const int block_num_y = (nrows + ny - 1) / ny;
const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_nums(1, 1, block_num_y);
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE); const sycl::range<3> block_dims(1, ny, WARP_SIZE);
stream->parallel_for( stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims), sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
dequantize_mul_mat_vec_q6_k_reorder(vx, y, dst, ncols, nrows, item_ct1); dequantize_mul_mat_vec_q6_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
}); });
} }