Commit 5e18eb021 for llama.cpp
commit 5e18eb021b021a133cfa4d9ac84dab644bc24318
Author: Titaniumtown <titaniumtown@proton.me>
Date: Mon Nov 2 03:51:20 2026 -0800
sycl: dmmv code cleanup (#31325)
* sycl: fold q8_0 reorder dmmv into the generic template
* sycl: template the dmmv non-K launchers
* sycl: template the dmmv K-quant launchers
diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp
index d720bc1f8..19af72c0e 100644
--- a/ggml/src/ggml-sycl/dmmv.cpp
+++ b/ggml/src/ggml-sycl/dmmv.cpp
@@ -112,11 +112,12 @@ static void dequantize_mul_mat_vec(const void * __restrict__ vx, const dfloat *
}
}
-template <int qk, int qr, dequantize_kernel_t_reorder dequantize_kernel_reorder>
+template <int qk, int qr, int qbytes, dequantize_kernel_t_reorder dequantize_kernel_reorder>
static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const dfloat * __restrict__ y, float * __restrict__ dst, const int ncols, const int nrows,
const sycl::nd_item<3> &item_ct1) {
// qk = quantized weights per x block
// qr = number of quantized weights per data value in x block
+ // qbytes = bytes of the quants region per block (not qk/qr for bit-packed types)
const int row = item_ct1.get_group(2) * item_ct1.get_local_range(1) +
item_ct1.get_local_id(1);
@@ -127,7 +128,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
const int tid = item_ct1.get_local_id(2);
- const int ncols_left = ncols % (QK4_0*WARP_SIZE);
+ const int ncols_left = ncols % (qk*WARP_SIZE);
const int ncols_align = ncols - ncols_left;
const int iter_stride = 8*2*GGML_SYCL_DMMV_X;
const int vals_per_iter = iter_stride / WARP_SIZE; // num quantized vals per thread and i iter //64/16=4, 512/16/2= 16
@@ -139,7 +140,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
#else
float tmp = 0.0f;
#endif // GGML_SYCL_F16
- const char *d_ptr = (const char*)vx+ncols*nrows/2;
+ const char *d_ptr = (const char*)vx + (size_t) ncols * nrows / qk * qbytes;
int i=0;
for (i = 0; i < ncols_align; i += iter_stride) {
const int col = i + vals_per_iter*tid;
@@ -155,7 +156,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
// dequantize
// for qr = 2 the iqs needs to increase by 1 per j iter because 2 weights per data val
dfloat2 v;
- dequantize_kernel_reorder((const void *)d_ptr, ib, (const void *)vx, ib * QK4_0 / 2 +iqs+j/qr, v);
+ dequantize_kernel_reorder((const void *)d_ptr, ib, (const void *)vx, ib * (qk/qr) +iqs+j/qr, v);
// matrix multiplication
// for qr = 2 the y index needs to increase by 1 per j iter because of y_offset = qk/2
@@ -172,7 +173,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
}
for (; i < ncols; i += iter_stride) {
- if (tid>=ncols_left/QK4_0) continue;
+ if (tid>=ncols_left/qk) continue;
const int col = i + vals_per_iter*tid;
const int ib = (row*ncols + col)/qk; // x block index
const int iqs = (col%qk)/qr; // x quant index
@@ -186,7 +187,7 @@ static void dequantize_mul_mat_vec_reorder(const void * __restrict__ vx, const d
// dequantize
// for qr = 2 the iqs needs to increase by 1 per j iter because 2 weights per data val
dfloat2 v;
- dequantize_kernel_reorder((const void *)d_ptr, ib, (const void *)vx, ib * QK4_0 / 2 +iqs+j/qr, v);
+ dequantize_kernel_reorder((const void *)d_ptr, ib, (const void *)vx, ib * (qk/qr) +iqs+j/qr, v);
// matrix multiplication
// for qr = 2 the y index needs to increase by 1 per j iter because of y_offset = qk/2
@@ -1487,55 +1488,10 @@ static void dequantize_mul_mat_vec_q6_k_reorder(const void * __restrict__ vx, co
}
}
-static void dequantize_mul_mat_vec_q4_0_sycl_reorder(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
- const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
- // the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
- const sycl::range<3> block_nums(1, 1, block_num_y);
- const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
- {
- dpct::has_capability_or_fail(stream->get_device(),
- {sycl::aspect::fp16});
-
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec_reorder<QK4_0, QR4_0, dequantize_q4_0_reorder>(
- vx, y, dst, ncols, nrows, item_ct1);
- });
- }
-}
-
-
-static void dequantize_mul_mat_vec_q4_0_sycl(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
- const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
- // the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
- const sycl::range<3> block_nums(1, 1, block_num_y);
- const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
- {
- dpct::has_capability_or_fail(stream->get_device(),
- {sycl::aspect::fp16});
-
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec<QK4_0, QR4_0, dequantize_q4_0>(
- vx, y, dst, ncols, nrows, item_ct1);
- });
- }
-}
-
-static void dequantize_mul_mat_vec_q1_0_sycl_reorder(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
+template <int qk, int qr, dequantize_kernel_t dequantize_kernel>
+static void launch_dequantize_mul_mat_vec(const void *vx, const dfloat *y, float *dst,
+ const int ncols, const int nrows,
+ dpct::queue_ptr stream) {
GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
// the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
@@ -1548,16 +1504,16 @@ static void dequantize_mul_mat_vec_q1_0_sycl_reorder(const void *vx, const dfloa
stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec_reorder<QK1_0, QR1_0, dequantize_q1_0_reorder>(
+ dequantize_mul_mat_vec<qk, qr, dequantize_kernel>(
vx, y, dst, ncols, nrows, item_ct1);
});
}
}
-static void dequantize_mul_mat_vec_q1_0_sycl(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
+template <int qk, int qr, int qbytes, dequantize_kernel_t_reorder dequantize_kernel>
+static void launch_dequantize_mul_mat_vec_reorder(const void *vx, const dfloat *y, float *dst,
+ const int ncols, const int nrows,
+ dpct::queue_ptr stream) {
GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
// the number of rows may exceed maximum grid size in the y or z dimensions, use the x dimension instead
@@ -1570,302 +1526,24 @@ static void dequantize_mul_mat_vec_q1_0_sycl(const void *vx, const dfloat *y,
stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec<QK1_0, QR1_0, dequantize_q1_0>(
- vx, y, dst, ncols, nrows, item_ct1);
- });
- }
-}
-
-static void dequantize_mul_mat_vec_q4_1_sycl(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
- const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
- const sycl::range<3> block_nums(1, 1, block_num_y);
- const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
- {
- dpct::has_capability_or_fail(stream->get_device(),
- {sycl::aspect::fp16});
-
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec<QK4_1, QR4_1, dequantize_q4_1>(
- vx, y, dst, ncols, nrows, item_ct1);
- });
- }
-}
-
-static void dequantize_mul_mat_vec_q5_0_sycl(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
- const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
- const sycl::range<3> block_nums(1, 1, block_num_y);
- const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
- {
- dpct::has_capability_or_fail(stream->get_device(),
- {sycl::aspect::fp16});
-
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec<QK5_0, QR5_0, dequantize_q5_0>(
- vx, y, dst, ncols, nrows, item_ct1);
- });
- }
-}
-
-static void dequantize_mul_mat_vec_q5_1_sycl(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
- const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
- const sycl::range<3> block_nums(1, 1, block_num_y);
- const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
- {
- dpct::has_capability_or_fail(stream->get_device(),
- {sycl::aspect::fp16});
-
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec<QK5_1, QR5_1, dequantize_q5_1>(
- vx, y, dst, ncols, nrows, item_ct1);
- });
- }
-}
-
-static void dequantize_mul_mat_vec_q8_0_sycl_reorder(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
- const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
- const sycl::range<3> block_nums(1, 1, block_num_y);
- const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
- {
- dpct::has_capability_or_fail(stream->get_device(),
- {sycl::aspect::fp16});
-
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- // Q8_0 reorder layout: [all qs (ncols*nrows bytes)][all d values]
- // Cannot reuse dequantize_mul_mat_vec_reorder template because it has
- // Q4_0-specific constants hardcoded (d_ptr offset and qs stride).
- const int row = item_ct1.get_group(2) * item_ct1.get_local_range(1) +
- item_ct1.get_local_id(1);
- if (row >= nrows) return;
-
- const int tid = item_ct1.get_local_id(2);
- const int iter_stride = 8*2*GGML_SYCL_DMMV_X;
- const int vals_per_iter = iter_stride / WARP_SIZE;
- const int ncols_left = ncols % (QK8_0*WARP_SIZE);
- const int ncols_align = ncols - ncols_left;
-
-#ifdef GGML_SYCL_F16
- sycl::half2 tmp = {0.0f, 0.0f};
-#else
- float tmp = 0.0f;
-#endif
- const char *d_ptr = (const char*)vx + ncols*nrows; // d after all qs
-
- int i = 0;
- for (i = 0; i < ncols_align; i += iter_stride) {
- const int col = i + vals_per_iter*tid;
- const int ib = (row*ncols + col)/QK8_0;
- const int iqs = col % QK8_0;
-
-#pragma unroll
- for (int j = 0; j < vals_per_iter; j += 2) {
- dfloat2 v;
- dequantize_q8_0_reorder((const void *)d_ptr, ib, (const void *)vx,
- ib * QK8_0 + iqs + j, v);
-
-#ifdef GGML_SYCL_F16
- dfloat2 t1{y[col + j + 0], y[col + j + 1]};
- tmp += v * t1;
-#else
- tmp += v.x() * y[col + j + 0];
- tmp += v.y() * y[col + j + 1];
-#endif
- }
- }
-
- // handle remaining columns
- for (; i < ncols; i += iter_stride) {
- if (tid >= ncols_left/QK8_0) continue;
- const int col = i + vals_per_iter*tid;
- const int ib = (row*ncols + col)/QK8_0;
- const int iqs = col % QK8_0;
-
-#pragma unroll
- for (int j = 0; j < vals_per_iter; j += 2) {
- dfloat2 v;
- dequantize_q8_0_reorder((const void *)d_ptr, ib, (const void *)vx,
- ib * QK8_0 + iqs + j, v);
-
-#ifdef GGML_SYCL_F16
- dfloat2 t1{y[col + j + 0], y[col + j + 1]};
- tmp += v * t1;
-#else
- tmp += v.x() * y[col + j + 0];
- tmp += v.y() * y[col + j + 1];
-#endif
- }
- }
-
- // reduce
- const int mask_start = ncols > GGML_SYCL_DMMV_X ? WARP_SIZE >> 1 : WARP_SIZE >> 2;
- for (int mask = mask_start; mask > 0; mask >>= 1) {
- tmp += dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask);
- }
-
- if (tid == 0) {
-#ifdef GGML_SYCL_F16
- dst[row] = tmp.x() + tmp.y();
-#else
- dst[row] = tmp;
-#endif
- }
- });
- }
-}
-
-static void dequantize_mul_mat_vec_q8_0_sycl(const void *vx, const dfloat *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % GGML_SYCL_DMMV_X == 0);
- const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
- const sycl::range<3> block_nums(1, 1, block_num_y);
- const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
- {
- dpct::has_capability_or_fail(stream->get_device(),
- {sycl::aspect::fp16});
-
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec<QK8_0, QR8_0, dequantize_q8_0>(
+ dequantize_mul_mat_vec_reorder<qk, qr, qbytes, dequantize_kernel>(
vx, y, dst, ncols, nrows, item_ct1);
});
}
}
-static void dequantize_mul_mat_vec_q2_K_sycl(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- 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 sycl::range<3> block_nums(1, 1, block_num_y);
- const sycl::range<3> block_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](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);
- });
-}
-
-static void dequantize_mul_mat_vec_q2_K_sycl_reorder(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
- 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_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec_q2_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
- });
-}
-
-static void dequantize_mul_mat_vec_q3_K_sycl(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
- 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_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](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);
- });
-}
-
-static void dequantize_mul_mat_vec_q3_K_sycl_reorder(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
- 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_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](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);
- });
-}
-
-static void dequantize_mul_mat_vec_q4_K_sycl(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
- 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_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](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);
- });
-}
-
-static void dequantize_mul_mat_vec_q5_K_sycl(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
+template <int ny, auto kernel>
+static void launch_dequantize_mul_mat_vec_k(const void *vx, const float *y, float *dst,
+ const int ncols, const int nrows,
+ dpct::queue_ptr stream) {
GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
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_dims(1, ny, WARP_SIZE);
stream->parallel_for(
sycl::nd_range<3>(block_nums * block_dims, block_dims),
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec_q5_k(vx, y, dst, ncols, nrows, item_ct1);
- });
-}
-
-static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
- 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_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](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);
+ kernel(vx, y, dst, ncols, nrows, item_ct1);
});
}
@@ -2248,54 +1926,6 @@ bool ggml_sycl_dequantize_mul_mat_vec_glu_reorder_esimd(enum ggml_type src0_type
#endif // GGML_SYCL_DMMV_HAS_ESIMD
-static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
- 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_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](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);
- });
-}
-
-static void dequantize_mul_mat_vec_q5_K_sycl_reorder(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
- 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_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
- dequantize_mul_mat_vec_q5_k_reorder(vx, y, dst, ncols, nrows, item_ct1);
- });
-}
-
-static void dequantize_mul_mat_vec_q6_K_sycl_reorder(const void *vx, const float *y,
- float *dst, const int ncols,
- const int nrows,
- dpct::queue_ptr stream) {
- GGML_ASSERT(ncols % QK_K == 0);
- const int ny = 2 / K_QUANTS_PER_ITERATION;
- 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_dims(1, ny, WARP_SIZE);
- stream->parallel_for(
- sycl::nd_range<3>(block_nums * block_dims, block_dims),
- [=](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);
- });
-}
-
void ggml_sycl_op_dequantize_mul_mat_vec(
ggml_backend_sycl_context & ctx,
const ggml_tensor *src0, const ggml_tensor *src1, ggml_tensor *dst,
@@ -2344,27 +1974,27 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
case GGML_TYPE_Q1_0:
if ((ggml_tensor_extra_gpu*)dst->src[0]->extra &&
((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
- dequantize_mul_mat_vec_q1_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_reorder<QK1_0, QR1_0, QK1_0/8, dequantize_q1_0_reorder>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
} else {
- dequantize_mul_mat_vec_q1_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec<QK1_0, QR1_0, dequantize_q1_0>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
}
break;
case GGML_TYPE_Q4_0:
if ((ggml_tensor_extra_gpu*)dst->src[0]->extra &&
((ggml_tensor_extra_gpu*)dst->src[0]->extra)->optimized_feature.reorder) {
- dequantize_mul_mat_vec_q4_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_reorder<QK4_0, QR4_0, QK4_0/2, dequantize_q4_0_reorder>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
} else {
- dequantize_mul_mat_vec_q4_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec<QK4_0, QR4_0, dequantize_q4_0>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
}
break;
case GGML_TYPE_Q4_1:
- dequantize_mul_mat_vec_q4_1_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec<QK4_1, QR4_1, dequantize_q4_1>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
break;
case GGML_TYPE_Q5_0:
- dequantize_mul_mat_vec_q5_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec<QK5_0, QR5_0, dequantize_q5_0>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
break;
case GGML_TYPE_Q5_1:
- dequantize_mul_mat_vec_q5_1_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec<QK5_1, QR5_1, dequantize_q5_1>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
break;
case GGML_TYPE_Q8_0:
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
@@ -2375,10 +2005,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
} else
#endif
{
- dequantize_mul_mat_vec_q8_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_reorder<QK8_0, QR8_0, QK8_0, dequantize_q8_0_reorder>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
}
} else {
- dequantize_mul_mat_vec_q8_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec<QK8_0, QR8_0, dequantize_q8_0>(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream);
}
break;
case GGML_TYPE_Q2_K:
@@ -2391,10 +2021,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
else
#endif
{
- dequantize_mul_mat_vec_q2_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q2_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
} else {
- dequantize_mul_mat_vec_q2_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2, dequantize_mul_mat_vec_q2_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
break;
case GGML_TYPE_Q3_K:
@@ -2407,10 +2037,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
else
#endif
{
- dequantize_mul_mat_vec_q3_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q3_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
} else {
- dequantize_mul_mat_vec_q3_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q3_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
break;
case GGML_TYPE_Q4_K:
@@ -2423,10 +2053,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
else
#endif
{
- dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q4_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
} else {
- dequantize_mul_mat_vec_q4_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q4_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
break;
case GGML_TYPE_Q5_K:
@@ -2439,10 +2069,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
else
#endif
{
- dequantize_mul_mat_vec_q5_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q5_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
} else {
- dequantize_mul_mat_vec_q5_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q5_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
break;
case GGML_TYPE_Q6_K:
@@ -2455,10 +2085,10 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
else
#endif
{
- dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q6_k_reorder>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
} else {
- dequantize_mul_mat_vec_q6_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ launch_dequantize_mul_mat_vec_k<2/K_QUANTS_PER_ITERATION, dequantize_mul_mat_vec_q6_k>(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
break;
case GGML_TYPE_F16: