@@ -894,7 +894,7 @@ static void dequantize_mul_mat_vec_q4_0_sycl_reorder(const void *vx, const dfloa
894894
895895 stream->parallel_for (
896896 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)]] {
898898 dequantize_mul_mat_vec_reorder<QK4_0, QR4_0, dequantize_q4_0_reorder>(
899899 vx, y, dst, ncols, nrows, item_ct1);
900900 });
@@ -927,7 +927,7 @@ static void dequantize_mul_mat_vec_q4_0_sycl(const void *vx, const dfloat *y,
927927 const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
928928 stream->parallel_for (
929929 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)]] {
931931 dequantize_mul_mat_vec<QK4_0, QR4_0, dequantize_q4_0>(
932932 vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
933933 });
@@ -955,7 +955,7 @@ static void dequantize_mul_mat_vec_q4_1_sycl(const void *vx, const dfloat *y,
955955 const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
956956 stream->parallel_for (
957957 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)]] {
959959 dequantize_mul_mat_vec<QK4_1, QR4_1, dequantize_q4_1>(
960960 vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
961961 });
@@ -983,7 +983,7 @@ static void dequantize_mul_mat_vec_q5_0_sycl(const void *vx, const dfloat *y,
983983 const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
984984 stream->parallel_for (
985985 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)]] {
987987 dequantize_mul_mat_vec<QK5_0, QR5_0, dequantize_q5_0>(
988988 vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
989989 });
@@ -1011,7 +1011,7 @@ static void dequantize_mul_mat_vec_q5_1_sycl(const void *vx, const dfloat *y,
10111011 const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
10121012 stream->parallel_for (
10131013 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)]] {
10151015 dequantize_mul_mat_vec<QK5_1, QR5_1, dequantize_q5_1>(
10161016 vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
10171017 });
@@ -1038,7 +1038,7 @@ static void dequantize_mul_mat_vec_q8_0_sycl(const void *vx, const dfloat *y,
10381038 const sycl::range<3 > block_dims (1 , GGML_SYCL_MMV_Y, WARP_SIZE);
10391039 stream->parallel_for (
10401040 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)]] {
10421042 dequantize_mul_mat_vec<QK8_0, QR8_0, dequantize_q8_0>(
10431043 vx, y, dst, ncols, nrows, WARP_SIZE, item_ct1);
10441044 });
@@ -1195,7 +1195,6 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
11951195 default :
11961196 printf (" ggml_sycl_op_dequantize_mul_mat_vec unsupported GGML_TYPE %d\n " , src0->type );
11971197 GGML_ABORT (" fatal error" );
1198- break ;
11991198 }
12001199
12011200 GGML_UNUSED (src1);
0 commit comments