Skip to content

Commit 25febe5

Browse files
s-NickMinh141120
authored andcommitted
sycl: Add reorder to Q6_K mmvq implementation (ggml-org#13885)
* Add Reorder to Q6_K mmvq implementation * Address PR comments: clean up comments * Remove unused parameter after refactoring q4_k * Adding inline to function and removing unnecessary reference to int --------- Signed-off-by: nscipione <nicolo.scipione@codeplay.com>
1 parent 0afa563 commit 25febe5

File tree

2 files changed

+9
-9
lines changed

2 files changed

+9
-9
lines changed

ggml/src/ggml-sycl/convert.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -236,9 +236,9 @@ static void dequantize_row_q6_K_sycl_reorder(const void * vx, dst_t * y, const i
236236

237237
dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 });
238238

239-
sycl_parallel_for(stream,
240-
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
241-
[=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y, item_ct1, nb); });
239+
stream->parallel_for(
240+
sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)),
241+
[=](sycl::nd_item<3> item_ct1) { dequantize_block_q6_K_reorder(vx, y, item_ct1, nb); });
242242
}
243243

244244
template <typename dst_t>

ggml/src/ggml-sycl/mmvq.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -754,12 +754,12 @@ static void reorder_mul_mat_vec_q6_k_q8_1_sycl(const void * vx, const void * vy,
754754
const sycl::range<3> global_size(1, GGML_SYCL_MMV_Y, block_num_y * WARP_SIZE);
755755
const sycl::range<3> workgroup_size(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE);
756756

757-
sycl_launch(stream, [&](sycl::handler & cgh) {
758-
sycl_parallel_for(cgh, sycl::nd_range<3>(global_size, workgroup_size),
759-
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
760-
mul_mat_vec_q_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q6_K>>(vx, vy, dst, ncols, nrows,
761-
nd_item);
762-
});
757+
stream->submit([&](sycl::handler & cgh) {
758+
cgh.parallel_for(sycl::nd_range<3>(global_size, workgroup_size),
759+
[=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
760+
mul_mat_vec_q_reorder<reorder_vec_dot_q_sycl<GGML_TYPE_Q6_K>>(vx, vy, dst, ncols, nrows,
761+
nd_item);
762+
});
763763
});
764764
}
765765
static void mul_mat_vec_q6_K_q8_1_sycl(const void *vx, const void *vy,

0 commit comments

Comments
 (0)