@@ -15575,11 +15575,11 @@ __dpct_inline__ static void k_copy_dst_from_contiguous(
15575
15575
static void ggml_sycl_mul_mat_id(const ggml_tensor *src0,
15576
15576
const ggml_tensor *src1,
15577
15577
ggml_tensor *dst) try {
15578
- GGML_ASSERT(src0->backend != GGML_BACKEND_TYPE_GPU_SPLIT &&
15579
- "mul_mat_id does not support split buffers");
15578
+ GGML_ASSERT(!ggml_backend_buffer_is_sycl_split( src0->buffer) && "mul_mat_id does not support split buffers");
15579
+
15580
15580
const ggml_tensor *ids = dst->src[2];
15581
15581
GGML_TENSOR_BINARY_OP_LOCALS
15582
- GGML_ASSERT(!ggml_backend_buffer_is_sycl_split(src0->buffer) && "mul_mat_id does not support split buffers");
15582
+
15583
15583
const dpct::queue_ptr stream = g_syclStreams[g_main_device][0];
15584
15584
15585
15585
const int64_t n_as = ne02;
@@ -15698,13 +15698,13 @@ static void ggml_sycl_mul_mat_id(const ggml_tensor *src0,
15698
15698
sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne10, 768u));
15699
15699
sycl::range<3> grid_dims(1, n_ids, ids->ne[1]);
15700
15700
stream->submit([&](sycl::handler &cgh) {
15701
- sycl::local_accessor<int, 0> src1_row_acc_ct1 (cgh);
15701
+ sycl::local_accessor<int, 0> src1_row_acc (cgh);
15702
15702
15703
- char *__restrict src1_contiguous_get_ct1 =
15703
+ char *__restrict src1_contiguous_get =
15704
15704
src1_contiguous.get();
15705
- int *__restrict dev_cur_src1_row_get_ct2 =
15705
+ int *__restrict dev_cur_src1_row_get =
15706
15706
dev_cur_src1_row.get();
15707
- mmid_row_mapping *__restrict dev_row_mapping_get_ct3 =
15707
+ mmid_row_mapping *__restrict dev_row_mapping_get =
15708
15708
dev_row_mapping.get();
15709
15709
size_t ids_nb_ct6 = ids->nb[1];
15710
15710
size_t ids_nb_ct7 = ids->nb[0];
@@ -15713,11 +15713,11 @@ static void ggml_sycl_mul_mat_id(const ggml_tensor *src0,
15713
15713
sycl::nd_range<3>(grid_dims * block_dims, block_dims),
15714
15714
[=](sycl::nd_item<3> item_ct1) {
15715
15715
k_copy_src1_to_contiguous(
15716
- src1_original, src1_contiguous_get_ct1 ,
15717
- dev_cur_src1_row_get_ct2 ,
15718
- dev_row_mapping_get_ct3 , ids_dev, i02,
15716
+ src1_original, src1_contiguous_get ,
15717
+ dev_cur_src1_row_get ,
15718
+ dev_row_mapping_get , ids_dev, i02,
15719
15719
ids_nb_ct6, ids_nb_ct7, ne11, ne10, nb11, nb12,
15720
- item_ct1, src1_row_acc_ct1 );
15720
+ item_ct1, src1_row_acc );
15721
15721
});
15722
15722
});
15723
15723
}
@@ -15743,25 +15743,22 @@ static void ggml_sycl_mul_mat_id(const ggml_tensor *src0,
15743
15743
sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne0, 768u));
15744
15744
sycl::range<3> grid_dims(1, 1, num_src1_rows);
15745
15745
stream->submit([&](sycl::handler &cgh) {
15746
- const char *__restrict dst_contiguous_get_ct1 =
15746
+ const char *__restrict dst_contiguous_get =
15747
15747
dst_contiguous.get();
15748
- const mmid_row_mapping *__restrict dev_row_mapping_get_ct2 =
15748
+ const mmid_row_mapping *__restrict dev_row_mapping_get =
15749
15749
dev_row_mapping.get();
15750
15750
15751
15751
cgh.parallel_for(
15752
15752
sycl::nd_range<3>(grid_dims * block_dims, block_dims),
15753
15753
[=](sycl::nd_item<3> item_ct1) {
15754
15754
k_copy_dst_from_contiguous(dst_original,
15755
- dst_contiguous_get_ct1 ,
15756
- dev_row_mapping_get_ct2 ,
15755
+ dst_contiguous_get ,
15756
+ dev_row_mapping_get ,
15757
15757
ne0, nb1, nb2, item_ct1);
15758
15758
});
15759
15759
});
15760
15760
}
15761
15761
}
15762
- if (dst->backend == GGML_BACKEND_TYPE_CPU) {
15763
- SYCL_CHECK(CHECK_TRY_ERROR(stream->wait()));
15764
- }
15765
15762
}
15766
15763
}
15767
15764
catch (sycl::exception const &exc) {
0 commit comments