Skip to content

Commit ffcd91f

Browse files
author
Aidan
committed
Update ggml_sycl_op_mul_mat_vec_q
1 parent 7e1ae37 commit ffcd91f

File tree

1 file changed

+64
-188
lines changed

1 file changed

+64
-188
lines changed

‎ggml-sycl.cpp

Lines changed: 64 additions & 188 deletions
Original file line numberDiff line numberDiff line change
@@ -9188,174 +9188,22 @@ static void convert_mul_mat_vec_f16_sycl(const void *vx, const dfloat *y,
91889188
}
91899189
}
91909190

9191-
static void mul_mat_vec_q4_0_q8_1_sycl(const void *vx, const void *vy,
9192-
float *dst, const int ncols,
9193-
const int nrows,
9194-
dpct::queue_ptr stream) {
9195-
GGML_ASSERT(ncols % QK4_0 == 0);
9196-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9197-
const sycl::range<3> block_nums(1, 1, block_num_y);
9198-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9199-
stream->parallel_for(
9200-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9201-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9202-
mul_mat_vec_q<QK4_0, QI4_0, block_q4_0, VDR_Q4_0_Q8_1_MMVQ,
9203-
vec_dot_q4_0_q8_1>(vx, vy, dst, ncols, nrows,
9204-
item_ct1);
9205-
});
9206-
}
9207-
9208-
static void mul_mat_vec_q4_1_q8_1_sycl(const void *vx, const void *vy,
9209-
float *dst, const int ncols,
9210-
const int nrows,
9211-
dpct::queue_ptr stream) {
9212-
GGML_ASSERT(ncols % QK4_1 == 0);
9213-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9214-
const sycl::range<3> block_nums(1, 1, block_num_y);
9215-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9216-
stream->parallel_for(
9217-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9218-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9219-
mul_mat_vec_q<QK4_0, QI4_1, block_q4_1, VDR_Q4_1_Q8_1_MMVQ,
9220-
vec_dot_q4_1_q8_1>(vx, vy, dst, ncols, nrows,
9221-
item_ct1);
9222-
});
9223-
}
9224-
9225-
static void mul_mat_vec_q5_0_q8_1_sycl(const void *vx, const void *vy,
9226-
float *dst, const int ncols,
9227-
const int nrows,
9228-
dpct::queue_ptr stream) {
9229-
GGML_ASSERT(ncols % QK5_0 == 0);
9230-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9231-
const sycl::range<3> block_nums(1, 1, block_num_y);
9232-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9233-
stream->parallel_for(
9234-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9235-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9236-
mul_mat_vec_q<QK5_0, QI5_0, block_q5_0, VDR_Q5_0_Q8_1_MMVQ,
9237-
vec_dot_q5_0_q8_1>(vx, vy, dst, ncols, nrows,
9238-
item_ct1);
9239-
});
9240-
}
9241-
9242-
static void mul_mat_vec_q5_1_q8_1_sycl(const void *vx, const void *vy,
9243-
float *dst, const int ncols,
9244-
const int nrows,
9245-
dpct::queue_ptr stream) {
9246-
GGML_ASSERT(ncols % QK5_1 == 0);
9247-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9248-
const sycl::range<3> block_nums(1, 1, block_num_y);
9249-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9250-
stream->parallel_for(
9251-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9252-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9253-
mul_mat_vec_q<QK5_1, QI5_1, block_q5_1, VDR_Q5_1_Q8_1_MMVQ,
9254-
vec_dot_q5_1_q8_1>(vx, vy, dst, ncols, nrows,
9255-
item_ct1);
9256-
});
9257-
}
9258-
9259-
static void mul_mat_vec_q8_0_q8_1_sycl(const void *vx, const void *vy,
9260-
float *dst, const int ncols,
9261-
const int nrows,
9262-
dpct::queue_ptr stream) {
9263-
GGML_ASSERT(ncols % QK8_0 == 0);
9264-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9265-
const sycl::range<3> block_nums(1, 1, block_num_y);
9266-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9267-
stream->parallel_for(
9268-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9269-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9270-
mul_mat_vec_q<QK8_0, QI8_0, block_q8_0, VDR_Q8_0_Q8_1_MMVQ,
9271-
vec_dot_q8_0_q8_1>(vx, vy, dst, ncols, nrows,
9272-
item_ct1);
9273-
});
9274-
}
9275-
9276-
static void mul_mat_vec_q2_K_q8_1_sycl(const void *vx, const void *vy,
9277-
float *dst, const int ncols,
9278-
const int nrows,
9279-
dpct::queue_ptr stream) {
9280-
GGML_ASSERT(ncols % QK_K == 0);
9281-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9282-
const sycl::range<3> block_nums(1, 1, block_num_y);
9283-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9284-
stream->parallel_for(
9285-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9286-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9287-
mul_mat_vec_q<QK_K, QI2_K, block_q2_K, VDR_Q2_K_Q8_1_MMVQ,
9288-
vec_dot_q2_K_q8_1>(vx, vy, dst, ncols, nrows,
9289-
item_ct1);
9290-
});
9291-
}
9292-
9293-
static void mul_mat_vec_q3_K_q8_1_sycl(const void *vx, const void *vy,
9294-
float *dst, const int ncols,
9295-
const int nrows,
9296-
dpct::queue_ptr stream) {
9297-
GGML_ASSERT(ncols % QK_K == 0);
9298-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9299-
const sycl::range<3> block_nums(1, 1, block_num_y);
9300-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9301-
stream->parallel_for(
9302-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9303-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9304-
mul_mat_vec_q<QK_K, QI3_K, block_q3_K, VDR_Q3_K_Q8_1_MMVQ,
9305-
vec_dot_q3_K_q8_1>(vx, vy, dst, ncols, nrows,
9306-
item_ct1);
9307-
});
9308-
}
9309-
9310-
static void mul_mat_vec_q4_K_q8_1_sycl(const void *vx, const void *vy,
9311-
float *dst, const int ncols,
9312-
const int nrows,
9313-
dpct::queue_ptr stream) {
9314-
GGML_ASSERT(ncols % QK_K == 0);
9315-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9316-
const sycl::range<3> block_nums(1, 1, block_num_y);
9317-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9318-
stream->parallel_for(
9319-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9320-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9321-
mul_mat_vec_q<QK_K, QI4_K, block_q4_K, VDR_Q4_K_Q8_1_MMVQ,
9322-
vec_dot_q4_K_q8_1>(vx, vy, dst, ncols, nrows,
9323-
item_ct1);
9324-
});
9325-
}
9326-
9327-
static void mul_mat_vec_q5_K_q8_1_sycl(const void *vx, const void *vy,
9328-
float *dst, const int ncols,
9329-
const int nrows,
9330-
dpct::queue_ptr stream) {
9331-
GGML_ASSERT(ncols % QK_K == 0);
9332-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9333-
const sycl::range<3> block_nums(1, 1, block_num_y);
9334-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9335-
stream->parallel_for(
9336-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9337-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9338-
mul_mat_vec_q<QK_K, QI5_K, block_q5_K, VDR_Q5_K_Q8_1_MMVQ,
9339-
vec_dot_q5_K_q8_1>(vx, vy, dst, ncols, nrows,
9340-
item_ct1);
9341-
});
9342-
}
9343-
9344-
static void mul_mat_vec_q6_K_q8_1_sycl(const void *vx, const void *vy,
9345-
float *dst, const int ncols,
9346-
const int nrows,
9347-
dpct::queue_ptr stream) {
9348-
GGML_ASSERT(ncols % QK_K == 0);
9349-
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9350-
const sycl::range<3> block_nums(1, 1, block_num_y);
9351-
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9352-
stream->parallel_for(
9353-
sycl::nd_range<3>(block_nums * block_dims, block_dims),
9354-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9355-
mul_mat_vec_q<QK_K, QI6_K, block_q6_K, VDR_Q6_K_Q8_1_MMVQ,
9356-
vec_dot_q6_K_q8_1>(vx, vy, dst, ncols, nrows,
9357-
item_ct1);
9358-
});
9191+
template <int qk, int qi, typename block_q_t, int vdr,
9192+
vec_dot_q_sycl_t vec_dot_q_sycl>
9193+
static void mul_mat_vec_q_sycl_submitter(const void *vx, const void *vy,
9194+
float *dst, const int ncols,
9195+
const int nrows,
9196+
dpct::queue_ptr stream) {
9197+
GGML_ASSERT(ncols % QK4_0 == 0);
9198+
const int block_num_y = (nrows + GGML_SYCL_MMV_Y - 1) / GGML_SYCL_MMV_Y;
9199+
const sycl::range<3> block_nums(1, 1, block_num_y);
9200+
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
9201+
stream->parallel_for(
9202+
sycl::nd_range<3>(block_nums * block_dims, block_dims), [=
9203+
](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(32)]] {
9204+
mul_mat_vec_q<qk, qi, block_q_t, vdr, vec_dot_q_sycl>(
9205+
vx, vy, dst, ncols, nrows, item_ct1);
9206+
});
93599207
}
93609208

93619209
int get_device_index_by_id(int id){
@@ -12098,37 +11946,62 @@ inline void ggml_sycl_op_mul_mat_vec_q(
1209811946
const int64_t ne00 = src0->ne[0];
1209911947
const int64_t row_diff = row_high - row_low;
1210011948

11949+
// TODO: support these quantization types
11950+
GGML_ASSERT(!(src0->type == GGML_TYPE_IQ2_XXS ||
11951+
src0->type == GGML_TYPE_IQ2_XS ||
11952+
src0->type == GGML_TYPE_IQ3_XXS));
11953+
1210111954
switch (src0->type) {
1210211955
case GGML_TYPE_Q4_0:
12103-
mul_mat_vec_q4_0_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12104-
break;
11956+
mul_mat_vec_q_sycl_submitter<QK4_0, QI4_0, block_q4_0,
11957+
VDR_Q4_0_Q8_1_MMVQ, vec_dot_q4_0_q8_1>(
11958+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11959+
break;
1210511960
case GGML_TYPE_Q4_1:
12106-
mul_mat_vec_q4_1_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12107-
break;
11961+
mul_mat_vec_q_sycl_submitter<QK4_1, QI4_1, block_q4_1,
11962+
VDR_Q4_1_Q8_1_MMVQ, vec_dot_q4_1_q8_1>(
11963+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11964+
break;
1210811965
case GGML_TYPE_Q5_0:
12109-
mul_mat_vec_q5_0_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12110-
break;
11966+
mul_mat_vec_q_sycl_submitter<QK5_0, QI5_0, block_q5_0,
11967+
VDR_Q5_0_Q8_1_MMVQ, vec_dot_q5_0_q8_1>(
11968+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11969+
break;
1211111970
case GGML_TYPE_Q5_1:
12112-
mul_mat_vec_q5_1_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12113-
break;
11971+
mul_mat_vec_q_sycl_submitter<QK5_1, QI5_1, block_q5_1,
11972+
VDR_Q5_1_Q8_1_MMVQ, vec_dot_q5_1_q8_1>(
11973+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11974+
break;
1211411975
case GGML_TYPE_Q8_0:
12115-
mul_mat_vec_q8_0_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12116-
break;
11976+
mul_mat_vec_q_sycl_submitter<QK8_0, QI8_0, block_q8_0,
11977+
VDR_Q8_0_Q8_1_MMVQ, vec_dot_q8_0_q8_1>(
11978+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11979+
break;
1211711980
case GGML_TYPE_Q2_K:
12118-
mul_mat_vec_q2_K_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12119-
break;
11981+
mul_mat_vec_q_sycl_submitter<QK_K, QI2_K, block_q2_K,
11982+
VDR_Q2_K_Q8_1_MMVQ, vec_dot_q2_K_q8_1>(
11983+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11984+
break;
1212011985
case GGML_TYPE_Q3_K:
12121-
mul_mat_vec_q3_K_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12122-
break;
11986+
mul_mat_vec_q_sycl_submitter<QK_K, QI3_K, block_q3_K,
11987+
VDR_Q3_K_Q8_1_MMVQ, vec_dot_q3_K_q8_1>(
11988+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11989+
break;
1212311990
case GGML_TYPE_Q4_K:
12124-
mul_mat_vec_q4_K_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12125-
break;
11991+
mul_mat_vec_q_sycl_submitter<QK_K, QI4_K, block_q4_K,
11992+
VDR_Q4_K_Q8_1_MMVQ, vec_dot_q4_K_q8_1>(
11993+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11994+
break;
1212611995
case GGML_TYPE_Q5_K:
12127-
mul_mat_vec_q5_K_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12128-
break;
11996+
mul_mat_vec_q_sycl_submitter<QK_K, QI5_K, block_q5_K,
11997+
VDR_Q5_K_Q8_1_MMVQ, vec_dot_q5_K_q8_1>(
11998+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
11999+
break;
1212912000
case GGML_TYPE_Q6_K:
12130-
mul_mat_vec_q6_K_q8_1_sycl(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12131-
break;
12001+
mul_mat_vec_q_sycl_submitter<QK_K, QI6_K, block_q6_K,
12002+
VDR_Q6_K_Q8_1_MMVQ, vec_dot_q6_K_q8_1>(
12003+
src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, stream);
12004+
break;
1213212005
default:
1213312006
GGML_ASSERT(false);
1213412007
break;
@@ -15138,6 +15011,9 @@ static bool ggml_backend_sycl_supports_op(ggml_backend_t backend, const ggml_ten
1513815011
return false;
1513915012
}
1514015013

15014+
if (a->type == GGML_TYPE_IQ3_XXS) {
15015+
return false;
15016+
}
1514115017
if (a->type == GGML_TYPE_IQ2_XXS) {
1514215018
return false;
1514315019
}

0 commit comments

Comments
 (0)