Skip to content

Commit a53f7f7

Browse files
lslusarczyklslusarczyk
authored
fixed compilation warnings in ggml-sycl (ggml-org#12424)
1 parent 7dfad38 commit a53f7f7

File tree

9 files changed

+101
-106
lines changed

9 files changed

+101
-106
lines changed

‎ggml/src/ggml-sycl/convert.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -138,7 +138,7 @@ static void dequantize_row_q4_0_sycl_reorder(const void *vx, dst_t *y, const int
138138
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, n_warp) *
139139
sycl::range<3>(1, 1, WARP_SIZE),
140140
sycl::range<3>(1, 1, WARP_SIZE)),
141-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]]{
141+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
142142
dequantize_block_q4_0_reorder(vx, y, k, item_ct1);
143143
});
144144

‎ggml/src/ggml-sycl/dmmv.cpp‎

Lines changed: 12 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -210,7 +210,7 @@ static void convert_mul_mat_vec_f16_sycl(const void *vx, const dfloat *y,
210210

211211
stream->parallel_for(
212212
sycl::nd_range<3>(block_nums * block_dims, block_dims),
213-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
213+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
214214
dequantize_mul_mat_vec<1, 1, convert_f16>(vx, y, dst, ncols,
215215
nrows, item_ct1);
216216
});
@@ -879,7 +879,7 @@ static void dequantize_mul_mat_vec_q4_0_sycl_reorder(const void *vx, const dfloa
879879

880880
stream->parallel_for(
881881
sycl::nd_range<3>(block_nums * block_dims, block_dims),
882-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
882+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
883883
dequantize_mul_mat_vec_reorder<QK4_0, QR4_0, dequantize_q4_0_reorder>(
884884
vx, y, dst, ncols, nrows, item_ct1);
885885
});
@@ -902,7 +902,7 @@ static void dequantize_mul_mat_vec_q4_0_sycl(const void *vx, const dfloat *y,
902902

903903
stream->parallel_for(
904904
sycl::nd_range<3>(block_nums * block_dims, block_dims),
905-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
905+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
906906
dequantize_mul_mat_vec<QK4_0, QR4_0, dequantize_q4_0>(
907907
vx, y, dst, ncols, nrows, item_ct1);
908908
});
@@ -923,7 +923,7 @@ static void dequantize_mul_mat_vec_q4_1_sycl(const void *vx, const dfloat *y,
923923

924924
stream->parallel_for(
925925
sycl::nd_range<3>(block_nums * block_dims, block_dims),
926-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
926+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
927927
dequantize_mul_mat_vec<QK4_1, QR4_1, dequantize_q4_1>(
928928
vx, y, dst, ncols, nrows, item_ct1);
929929
});
@@ -944,7 +944,7 @@ static void dequantize_mul_mat_vec_q5_0_sycl(const void *vx, const dfloat *y,
944944

945945
stream->parallel_for(
946946
sycl::nd_range<3>(block_nums * block_dims, block_dims),
947-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
947+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
948948
dequantize_mul_mat_vec<QK5_0, QR5_0, dequantize_q5_0>(
949949
vx, y, dst, ncols, nrows, item_ct1);
950950
});
@@ -965,7 +965,7 @@ static void dequantize_mul_mat_vec_q5_1_sycl(const void *vx, const dfloat *y,
965965

966966
stream->parallel_for(
967967
sycl::nd_range<3>(block_nums * block_dims, block_dims),
968-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
968+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
969969
dequantize_mul_mat_vec<QK5_1, QR5_1, dequantize_q5_1>(
970970
vx, y, dst, ncols, nrows, item_ct1);
971971
});
@@ -986,7 +986,7 @@ static void dequantize_mul_mat_vec_q8_0_sycl(const void *vx, const dfloat *y,
986986

987987
stream->parallel_for(
988988
sycl::nd_range<3>(block_nums * block_dims, block_dims),
989-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(WARP_SIZE)]] {
989+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
990990
dequantize_mul_mat_vec<QK8_0, QR8_0, dequantize_q8_0>(
991991
vx, y, dst, ncols, nrows, item_ct1);
992992
});
@@ -1004,7 +1004,7 @@ static void dequantize_mul_mat_vec_q2_K_sycl(const void *vx, const float *y,
10041004
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE);
10051005
stream->parallel_for(
10061006
sycl::nd_range<3>(block_nums * block_dims, block_dims),
1007-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(QK_WARP_SIZE)]] {
1007+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] {
10081008
dequantize_mul_mat_vec_q2_k(vx, y, dst, ncols, nrows, item_ct1);
10091009
});
10101010
}
@@ -1020,7 +1020,7 @@ static void dequantize_mul_mat_vec_q3_K_sycl(const void *vx, const float *y,
10201020
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE);
10211021
stream->parallel_for(
10221022
sycl::nd_range<3>(block_nums * block_dims, block_dims),
1023-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(QK_WARP_SIZE)]] {
1023+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] {
10241024
dequantize_mul_mat_vec_q3_k(vx, y, dst, ncols, nrows, item_ct1);
10251025
});
10261026
}
@@ -1036,7 +1036,7 @@ static void dequantize_mul_mat_vec_q4_K_sycl(const void *vx, const float *y,
10361036
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE);
10371037
stream->parallel_for(
10381038
sycl::nd_range<3>(block_nums * block_dims, block_dims),
1039-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(QK_WARP_SIZE)]] {
1039+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] {
10401040
dequantize_mul_mat_vec_q4_k(vx, y, dst, ncols, nrows, item_ct1);
10411041
});
10421042
}
@@ -1049,7 +1049,7 @@ static void dequantize_mul_mat_vec_q5_K_sycl(const void *vx, const float *y,
10491049
const sycl::range<3> block_dims(1, 1, QK_WARP_SIZE);
10501050
stream->parallel_for(
10511051
sycl::nd_range<3>(sycl::range<3>(1, 1, nrows) * block_dims, block_dims),
1052-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(QK_WARP_SIZE)]] {
1052+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] {
10531053
dequantize_mul_mat_vec_q5_k(vx, y, dst, ncols, item_ct1);
10541054
});
10551055
}
@@ -1065,7 +1065,7 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
10651065
const sycl::range<3> block_dims(1, ny, QK_WARP_SIZE);
10661066
stream->parallel_for(
10671067
sycl::nd_range<3>(block_nums * block_dims, block_dims),
1068-
[=](sycl::nd_item<3> item_ct1) [[intel::reqd_sub_group_size(QK_WARP_SIZE)]] {
1068+
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] {
10691069
dequantize_mul_mat_vec_q6_k(vx, y, dst, ncols, nrows, item_ct1);
10701070
});
10711071
}
@@ -1143,7 +1143,6 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
11431143
default:
11441144
printf("ggml_sycl_op_dequantize_mul_mat_vec unsupported GGML_TYPE %d\n", src0->type);
11451145
GGML_ABORT("fatal error");
1146-
break;
11471146
}
11481147

11491148
GGML_UNUSED(src1);

0 commit comments

Comments
 (0)