@@ -894,7 +894,7 @@ static void dequantize_mul_mat_vec_q4_0_sycl_reorder(const void *vx, const dfloa
894
894
895
895
stream->parallel_for (
896
896
sycl::nd_range<3 >(block_nums * block_dims, block_dims),
897
- [=](sycl::nd_item<3 > item_ct1) [[intel ::reqd_sub_group_size (WARP_SIZE)]] {
897
+ [=](sycl::nd_item<3 > item_ct1) [[sycl ::reqd_sub_group_size (WARP_SIZE)]] {
898
898
dequantize_mul_mat_vec_reorder<QK4_0, QR4_0, dequantize_q4_0_reorder>(
899
899
vx, y, dst, ncols, nrows, item_ct1);
900
900
});
@@ -927,7 +927,7 @@ static void dequantize_mul_mat_vec_q4_0_sycl(const void *vx, const dfloat *y,
927
927
const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
928
928
stream->parallel_for (
929
929
sycl::nd_range<3 >(block_nums * block_dims, block_dims),
930
- [=](sycl::nd_item<3 > item_ct1) [[intel ::reqd_sub_group_size (WARP_SIZE)]] {
930
+ [=](sycl::nd_item<3 > item_ct1) [[sycl ::reqd_sub_group_size (WARP_SIZE)]] {
931
931
dequantize_mul_mat_vec<QK4_0, QR4_0, dequantize_q4_0>(
932
932
vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
933
933
});
@@ -955,7 +955,7 @@ static void dequantize_mul_mat_vec_q4_1_sycl(const void *vx, const dfloat *y,
955
955
const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
956
956
stream->parallel_for (
957
957
sycl::nd_range<3 >(block_nums * block_dims, block_dims),
958
- [=](sycl::nd_item<3 > item_ct1) [[intel ::reqd_sub_group_size (WARP_SIZE)]] {
958
+ [=](sycl::nd_item<3 > item_ct1) [[sycl ::reqd_sub_group_size (WARP_SIZE)]] {
959
959
dequantize_mul_mat_vec<QK4_1, QR4_1, dequantize_q4_1>(
960
960
vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
961
961
});
@@ -983,7 +983,7 @@ static void dequantize_mul_mat_vec_q5_0_sycl(const void *vx, const dfloat *y,
983
983
const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
984
984
stream->parallel_for (
985
985
sycl::nd_range<3 >(block_nums * block_dims, block_dims),
986
- [=](sycl::nd_item<3 > item_ct1) [[intel ::reqd_sub_group_size (WARP_SIZE)]] {
986
+ [=](sycl::nd_item<3 > item_ct1) [[sycl ::reqd_sub_group_size (WARP_SIZE)]] {
987
987
dequantize_mul_mat_vec<QK5_0, QR5_0, dequantize_q5_0>(
988
988
vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
989
989
});
@@ -1011,7 +1011,7 @@ static void dequantize_mul_mat_vec_q5_1_sycl(const void *vx, const dfloat *y,
1011
1011
const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
1012
1012
stream->parallel_for (
1013
1013
sycl::nd_range<3 >(block_nums * block_dims, block_dims),
1014
- [=](sycl::nd_item<3 > item_ct1) [[intel ::reqd_sub_group_size (WARP_SIZE)]] {
1014
+ [=](sycl::nd_item<3 > item_ct1) [[sycl ::reqd_sub_group_size (WARP_SIZE)]] {
1015
1015
dequantize_mul_mat_vec<QK5_1, QR5_1, dequantize_q5_1>(
1016
1016
vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
1017
1017
});
@@ -1038,7 +1038,7 @@ static void dequantize_mul_mat_vec_q8_0_sycl(const void *vx, const dfloat *y,
1038
1038
const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
1039
1039
stream->parallel_for (
1040
1040
sycl::nd_range<3 >(block_nums * block_dims, block_dims),
1041
- [=](sycl::nd_item<3 > item_ct1) [[intel ::reqd_sub_group_size (WARP_SIZE)]] {
1041
+ [=](sycl::nd_item<3 > item_ct1) [[sycl ::reqd_sub_group_size (WARP_SIZE)]] {
1042
1042
dequantize_mul_mat_vec<QK8_0, QR8_0, dequantize_q8_0>(
1043
1043
vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
1044
1044
});
@@ -1195,7 +1195,6 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
1195
1195
default :
1196
1196
printf (" ggml_sycl_op_dequantize_mul_mat_vec unsupported GGML_TYPE %d\n " , src0->type );
1197
1197
GGML_ABORT (" fatal error" );
1198
- break ;
1199
1198
}
1200
1199
1201
1200
GGML_UNUSED (src1);
0 commit comments