Skip to content

Commit 68a9dd6

Browse files
s-Nicks-Nick
committed
Conver submit and parallel_for in mmvq.cpp
1 parent 126bbbb commit 68a9dd6

File tree

1 file changed

+44
-44
lines changed

1 file changed

+44
-44
lines changed

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

Lines changed: 44 additions & 44 deletions
Original file line numberDiff line numberDiff line change
@@ -544,8 +544,8 @@ static void reorder_mul_mat_vec_q4_0_q8_1_sycl(const void * vx, const void * vy,
544544
const sycl::range<3> global_size(1, GGML_SYCL_MMV_Y, (block_num_y * WARP_SIZE));
545545
const sycl::range<3> workgroup_size(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);
546546

547-
stream->submit([&](sycl::handler & cgh) {
548-
cgh.parallel_for(sycl::nd_range<3>(global_size, workgroup_size),
547+
syclex::submit(*stream,[&](sycl::handler & cgh) {
548+
syclex::nd_launch(cgh,sycl::nd_range<3>(global_size, workgroup_size),
549549
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
550550
mul_mat_vec_q_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q4_0>>(vx, vy, dst, ncols, nrows,
551551
nd_item);
@@ -561,8 +561,8 @@ static void mul_mat_vec_q4_0_q8_1_sycl(const void * vx, const void * vy, float *
561561
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
562562

563563
{
564-
stream->submit([&](sycl::handler & cgh) {
565-
cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims),
564+
syclex::submit(*stream,[&](sycl::handler & cgh) {
565+
syclex::nd_launch(cgh,sycl::nd_range<3>(block_nums * block_dims, block_dims),
566566
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
567567
mul_mat_vec_q<QK4_0, QI4_0, block_q4_0, VDR_Q4_0_Q8_1_MMVQ, vec_dot_q4_0_q8_1>(
568568
vx, vy, dst, ncols, nrows, item_ct1);
@@ -581,9 +581,9 @@ static void mul_mat_vec_q4_1_q8_1_sycl(const void *vx, const void *vy,
581581
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
582582
{
583583

584-
stream->submit([&](sycl::handler &cgh) {
584+
syclex::submit(*stream,[&](sycl::handler &cgh) {
585585

586-
cgh.parallel_for(
586+
syclex::nd_launch(cgh,
587587
sycl::nd_range<3>(block_nums * block_dims, block_dims),
588588
[=](sycl::nd_item<3> item_ct1)
589589
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -605,9 +605,9 @@ static void mul_mat_vec_q5_0_q8_1_sycl(const void *vx, const void *vy,
605605
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
606606
{
607607

608-
stream->submit([&](sycl::handler &cgh) {
608+
syclex::submit(*stream,[&](sycl::handler &cgh) {
609609

610-
cgh.parallel_for(
610+
syclex::nd_launch(cgh,
611611
sycl::nd_range<3>(block_nums * block_dims, block_dims),
612612
[=](sycl::nd_item<3> item_ct1)
613613
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -629,9 +629,9 @@ static void mul_mat_vec_q5_1_q8_1_sycl(const void *vx, const void *vy,
629629
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
630630
{
631631

632-
stream->submit([&](sycl::handler &cgh) {
632+
syclex::submit(*stream,[&](sycl::handler &cgh) {
633633

634-
cgh.parallel_for(
634+
syclex::nd_launch(cgh,
635635
sycl::nd_range<3>(block_nums * block_dims, block_dims),
636636
[=](sycl::nd_item<3> item_ct1)
637637
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -653,9 +653,9 @@ static void mul_mat_vec_q8_0_q8_1_sycl(const void *vx, const void *vy,
653653
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
654654
{
655655

656-
stream->submit([&](sycl::handler &cgh) {
656+
syclex::submit(*stream,[&](sycl::handler &cgh) {
657657

658-
cgh.parallel_for(
658+
syclex::nd_launch(cgh,
659659
sycl::nd_range<3>(block_nums * block_dims, block_dims),
660660
[=](sycl::nd_item<3> item_ct1)
661661
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -677,9 +677,9 @@ static void mul_mat_vec_q2_K_q8_1_sycl(const void *vx, const void *vy,
677677
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
678678
{
679679

680-
stream->submit([&](sycl::handler &cgh) {
680+
syclex::submit(*stream,[&](sycl::handler &cgh) {
681681

682-
cgh.parallel_for(
682+
syclex::nd_launch(cgh,
683683
sycl::nd_range<3>(block_nums * block_dims, block_dims),
684684
[=](sycl::nd_item<3> item_ct1)
685685
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -701,9 +701,9 @@ static void mul_mat_vec_q3_K_q8_1_sycl(const void *vx, const void *vy,
701701
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
702702
{
703703

704-
stream->submit([&](sycl::handler &cgh) {
704+
syclex::submit(*stream,[&](sycl::handler &cgh) {
705705

706-
cgh.parallel_for(
706+
syclex::nd_launch(cgh,
707707
sycl::nd_range<3>(block_nums * block_dims, block_dims),
708708
[=](sycl::nd_item<3> item_ct1)
709709
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -725,9 +725,9 @@ static void mul_mat_vec_q4_K_q8_1_sycl(const void *vx, const void *vy,
725725
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
726726
{
727727

728-
stream->submit([&](sycl::handler &cgh) {
728+
syclex::submit(*stream,[&](sycl::handler &cgh) {
729729

730-
cgh.parallel_for(
730+
syclex::nd_launch(cgh,
731731
sycl::nd_range<3>(block_nums * block_dims, block_dims),
732732
[=](sycl::nd_item<3> item_ct1)
733733
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -750,8 +750,8 @@ static void reorder_mul_mat_vec_q4_k_q8_1_sycl(const void * vx, const void * vy,
750750
const sycl::range<3> global_size(1, GGML_SYCL_MMV_Y, block_num_y * WARP_SIZE);
751751
const sycl::range<3> workgroup_size(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);
752752

753-
stream->submit([&](sycl::handler & cgh) {
754-
cgh.parallel_for(sycl::nd_range<3>(global_size, workgroup_size),
753+
syclex::submit(*stream,[&](sycl::handler & cgh) {
754+
syclex::nd_launch(cgh,sycl::nd_range<3>(global_size, workgroup_size),
755755
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
756756
mul_mat_vec_q_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q4_K>>(vx, vy, dst, ncols,
757757
nrows, nd_item);
@@ -770,9 +770,9 @@ static void mul_mat_vec_q5_K_q8_1_sycl(const void *vx, const void *vy,
770770
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
771771
{
772772

773-
stream->submit([&](sycl::handler &cgh) {
773+
syclex::submit(*stream,[&](sycl::handler &cgh) {
774774

775-
cgh.parallel_for(
775+
syclex::nd_launch(cgh,
776776
sycl::nd_range<3>(block_nums * block_dims, block_dims),
777777
[=](sycl::nd_item<3> item_ct1)
778778
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -794,8 +794,8 @@ static void reorder_mul_mat_vec_q6_k_q8_1_sycl(const void * vx, const void * vy,
794794
const sycl::range<3> global_size(1, GGML_SYCL_MMV_Y, block_num_y * WARP_SIZE);
795795
const sycl::range<3> workgroup_size(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);
796796

797-
stream->submit([&](sycl::handler & cgh) {
798-
cgh.parallel_for(sycl::nd_range<3>(global_size, workgroup_size),
797+
syclex::submit(*stream,[&](sycl::handler & cgh) {
798+
syclex::nd_launch(cgh,sycl::nd_range<3>(global_size, workgroup_size),
799799
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
800800
mul_mat_vec_q_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q6_K>>(vx, vy, dst, ncols, nrows,
801801
nd_item);
@@ -812,9 +812,9 @@ static void mul_mat_vec_q6_K_q8_1_sycl(const void *vx, const void *vy,
812812
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
813813
{
814814

815-
stream->submit([&](sycl::handler &cgh) {
815+
syclex::submit(*stream,[&](sycl::handler &cgh) {
816816

817-
cgh.parallel_for(
817+
syclex::nd_launch(cgh,
818818
sycl::nd_range<3>(block_nums * block_dims, block_dims),
819819
[=](sycl::nd_item<3> item_ct1)
820820
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -836,8 +836,8 @@ static void mul_mat_vec_iq2_xxs_q8_1_sycl(const void *vx, const void *vy,
836836
const sycl::range<3> block_nums(1, 1, block_num_y);
837837
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
838838
{
839-
stream->submit([&](sycl::handler &cgh) {
840-
cgh.parallel_for(
839+
syclex::submit(*stream,[&](sycl::handler &cgh) {
840+
syclex::nd_launch(cgh,
841841
sycl::nd_range<3>(block_nums * block_dims, block_dims),
842842
[=](sycl::nd_item<3> item_ct1)
843843
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -857,8 +857,8 @@ static void mul_mat_vec_iq2_xs_q8_1_sycl(const void *vx, const void *vy,
857857
const sycl::range<3> block_nums(1, 1, block_num_y);
858858
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
859859
{
860-
stream->submit([&](sycl::handler & cgh) {
861-
cgh.parallel_for(
860+
syclex::submit(*stream,[&](sycl::handler & cgh) {
861+
syclex::nd_launch(cgh,
862862
sycl::nd_range<3>(block_nums * block_dims, block_dims),
863863
[=](sycl::nd_item<3> item_ct1)
864864
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -879,8 +879,8 @@ static void mul_mat_vec_iq2_s_q8_1_sycl(const void *vx, const void *vy,
879879
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
880880
{
881881

882-
stream->submit([&](sycl::handler &cgh) {
883-
cgh.parallel_for(
882+
syclex::submit(*stream,[&](sycl::handler &cgh) {
883+
syclex::nd_launch(cgh,
884884
sycl::nd_range<3>(block_nums * block_dims, block_dims),
885885
[=](sycl::nd_item<3> item_ct1)
886886
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -901,8 +901,8 @@ static void mul_mat_vec_iq3_xxs_q8_1_sycl(const void *vx, const void *vy,
901901
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
902902
{
903903

904-
stream->submit([&](sycl::handler &cgh) {
905-
cgh.parallel_for(
904+
syclex::submit(*stream,[&](sycl::handler &cgh) {
905+
syclex::nd_launch(cgh,
906906
sycl::nd_range<3>(block_nums * block_dims, block_dims),
907907
[=](sycl::nd_item<3> item_ct1)
908908
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -923,8 +923,8 @@ static void mul_mat_vec_iq3_s_q8_1_sycl(const void *vx, const void *vy,
923923
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
924924
{
925925

926-
stream->submit([&](sycl::handler &cgh) {
927-
cgh.parallel_for(
926+
syclex::submit(*stream,[&](sycl::handler &cgh) {
927+
syclex::nd_launch(cgh,
928928
sycl::nd_range<3>(block_nums * block_dims, block_dims),
929929
[=](sycl::nd_item<3> item_ct1)
930930
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -945,8 +945,8 @@ static void mul_mat_vec_iq1_s_q8_1_sycl(const void *vx, const void *vy,
945945
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
946946
{
947947

948-
stream->submit([&](sycl::handler &cgh) {
949-
cgh.parallel_for(
948+
syclex::submit(*stream,[&](sycl::handler &cgh) {
949+
syclex::nd_launch(cgh,
950950
sycl::nd_range<3>(block_nums * block_dims, block_dims),
951951
[=](sycl::nd_item<3> item_ct1)
952952
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -966,8 +966,8 @@ static void mul_mat_vec_iq1_m_q8_1_sycl(const void *vx, const void *vy,
966966
const sycl::range<3> block_nums(1, 1, block_num_y);
967967
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
968968
{
969-
stream->submit([&](sycl::handler &cgh) {
970-
cgh.parallel_for(
969+
syclex::submit(*stream,[&](sycl::handler &cgh) {
970+
syclex::nd_launch(cgh,
971971
sycl::nd_range<3>(block_nums * block_dims, block_dims),
972972
[=](sycl::nd_item<3> item_ct1)
973973
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -988,8 +988,8 @@ static void mul_mat_vec_iq4_nl_q8_1_sycl(const void *vx, const void *vy,
988988
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
989989
{
990990

991-
stream->submit([&](sycl::handler &cgh) {
992-
cgh.parallel_for(
991+
syclex::submit(*stream,[&](sycl::handler &cgh) {
992+
syclex::nd_launch(cgh,
993993
sycl::nd_range<3>(block_nums * block_dims, block_dims),
994994
[=](sycl::nd_item<3> item_ct1)
995995
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {
@@ -1010,8 +1010,8 @@ static void mul_mat_vec_iq4_xs_q8_1_sycl(const void *vx, const void *vy,
10101010
const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, WARP_SIZE);
10111011
{
10121012

1013-
stream->submit([&](sycl::handler &cgh) {
1014-
cgh.parallel_for(
1013+
syclex::submit(*stream,[&](sycl::handler &cgh) {
1014+
syclex::nd_launch(cgh,
10151015
sycl::nd_range<3>(block_nums * block_dims, block_dims),
10161016
[=](sycl::nd_item<3> item_ct1)
10171017
[[sycl::reqd_sub_group_size(WARP_SIZE)]] {

0 commit comments

Comments
 (0)