@@ -210,7 +210,7 @@ static void convert_mul_mat_vec_f16_sycl(const void *vx, const dfloat *y,
210
210
211
211
stream->parallel_for (
212
212
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)]] {
214
214
dequantize_mul_mat_vec<1 , 1 , convert_f16>(vx, y, dst, ncols,
215
215
nrows, item_ct1);
216
216
});
@@ -879,7 +879,7 @@ static void dequantize_mul_mat_vec_q4_0_sycl_reorder(const void *vx, const dfloa
879
879
880
880
stream->parallel_for (
881
881
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)]] {
883
883
dequantize_mul_mat_vec_reorder<QK4_0, QR4_0, dequantize_q4_0_reorder>(
884
884
vx, y, dst, ncols, nrows, item_ct1);
885
885
});
@@ -902,7 +902,7 @@ static void dequantize_mul_mat_vec_q4_0_sycl(const void *vx, const dfloat *y,
902
902
903
903
stream->parallel_for (
904
904
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)]] {
906
906
dequantize_mul_mat_vec<QK4_0, QR4_0, dequantize_q4_0>(
907
907
vx, y, dst, ncols, nrows, item_ct1);
908
908
});
@@ -923,7 +923,7 @@ static void dequantize_mul_mat_vec_q4_1_sycl(const void *vx, const dfloat *y,
923
923
924
924
stream->parallel_for (
925
925
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)]] {
927
927
dequantize_mul_mat_vec<QK4_1, QR4_1, dequantize_q4_1>(
928
928
vx, y, dst, ncols, nrows, item_ct1);
929
929
});
@@ -944,7 +944,7 @@ static void dequantize_mul_mat_vec_q5_0_sycl(const void *vx, const dfloat *y,
944
944
945
945
stream->parallel_for (
946
946
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)]] {
948
948
dequantize_mul_mat_vec<QK5_0, QR5_0, dequantize_q5_0>(
949
949
vx, y, dst, ncols, nrows, item_ct1);
950
950
});
@@ -965,7 +965,7 @@ static void dequantize_mul_mat_vec_q5_1_sycl(const void *vx, const dfloat *y,
965
965
966
966
stream->parallel_for (
967
967
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)]] {
969
969
dequantize_mul_mat_vec<QK5_1, QR5_1, dequantize_q5_1>(
970
970
vx, y, dst, ncols, nrows, item_ct1);
971
971
});
@@ -986,7 +986,7 @@ static void dequantize_mul_mat_vec_q8_0_sycl(const void *vx, const dfloat *y,
986
986
987
987
stream->parallel_for (
988
988
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)]] {
990
990
dequantize_mul_mat_vec<QK8_0, QR8_0, dequantize_q8_0>(
991
991
vx, y, dst, ncols, nrows, item_ct1);
992
992
});
@@ -1004,7 +1004,7 @@ static void dequantize_mul_mat_vec_q2_K_sycl(const void *vx, const float *y,
1004
1004
const sycl::range<3 > block_dims (1 , ny, QK_WARP_SIZE);
1005
1005
stream->parallel_for (
1006
1006
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)]] {
1008
1008
dequantize_mul_mat_vec_q2_k (vx, y, dst, ncols, nrows, item_ct1);
1009
1009
});
1010
1010
}
@@ -1020,7 +1020,7 @@ static void dequantize_mul_mat_vec_q3_K_sycl(const void *vx, const float *y,
1020
1020
const sycl::range<3 > block_dims (1 , ny, QK_WARP_SIZE);
1021
1021
stream->parallel_for (
1022
1022
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)]] {
1024
1024
dequantize_mul_mat_vec_q3_k (vx, y, dst, ncols, nrows, item_ct1);
1025
1025
});
1026
1026
}
@@ -1036,7 +1036,7 @@ static void dequantize_mul_mat_vec_q4_K_sycl(const void *vx, const float *y,
1036
1036
const sycl::range<3 > block_dims (1 , ny, QK_WARP_SIZE);
1037
1037
stream->parallel_for (
1038
1038
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)]] {
1040
1040
dequantize_mul_mat_vec_q4_k (vx, y, dst, ncols, nrows, item_ct1);
1041
1041
});
1042
1042
}
@@ -1049,7 +1049,7 @@ static void dequantize_mul_mat_vec_q5_K_sycl(const void *vx, const float *y,
1049
1049
const sycl::range<3 > block_dims (1 , 1 , QK_WARP_SIZE);
1050
1050
stream->parallel_for (
1051
1051
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)]] {
1053
1053
dequantize_mul_mat_vec_q5_k (vx, y, dst, ncols, item_ct1);
1054
1054
});
1055
1055
}
@@ -1065,7 +1065,7 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
1065
1065
const sycl::range<3 > block_dims (1 , ny, QK_WARP_SIZE);
1066
1066
stream->parallel_for (
1067
1067
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)]] {
1069
1069
dequantize_mul_mat_vec_q6_k (vx, y, dst, ncols, nrows, item_ct1);
1070
1070
});
1071
1071
}
@@ -1143,7 +1143,6 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
1143
1143
default :
1144
1144
printf (" ggml_sycl_op_dequantize_mul_mat_vec unsupported GGML_TYPE %d\n " , src0->type );
1145
1145
GGML_ABORT (" fatal error" );
1146
- break ;
1147
1146
}
1148
1147
1149
1148
GGML_UNUSED (src1);
0 commit comments