From 885e61372d3b81efc1d5ad8d2d1134b3e4b20c38 Mon Sep 17 00:00:00 2001 From: Alexey Kopytko Date: Tue, 16 Jun 2026 11:01:51 +0900 Subject: [PATCH 1/6] SYCL: fix a bug with async memcpy --- ggml/src/ggml-sycl/ggml-sycl.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index fb8665a02c32..f60192f00258 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4231,6 +4231,8 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, ggml_sycl_pool_alloc dev_row_mapping(ctx.pool(), n_routed_rows); SYCL_CHECK(CHECK_TRY_ERROR( stream->memcpy(dev_row_mapping.get(), routed_row_src.data(), n_routed_rows*sizeof(mmid_row_mapping)))); + // routed_row_src is a host stack vector; wait for the the above async copy + SYCL_CHECK(CHECK_TRY_ERROR(stream->wait())); const unsigned int max_work_group_size = ggml_sycl_info().max_work_group_sizes[ctx.device]; assert(max_work_group_size % (WARP_SIZE * WARP_SIZE) == 0); From dce6b06a3225db38296df251986db2e4503b0e41 Mon Sep 17 00:00:00 2001 From: Alexey Kopytko Date: Tue, 16 Jun 2026 12:00:28 +0900 Subject: [PATCH 2/6] make mmid_row_mapping_host persistent --- ggml/src/ggml-sycl/common.hpp | 7 +++++++ ggml/src/ggml-sycl/ggml-sycl.cpp | 9 +-------- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index 5fb1a1d6bc02..b315ed2e435b 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -323,6 +323,11 @@ void ggml_sycl_free_device(void *ptr, sycl::queue &q); void release_extra_gpu(ggml_tensor_extra_gpu * extra, std::vector streams={}); +struct mmid_row_mapping { + int32_t i1; + int32_t i2; +}; + namespace sycl_ex = sycl::ext::oneapi::experimental; struct ggml_backend_sycl_context { int device; @@ -420,6 +425,8 @@ struct ggml_backend_sycl_context { std::unique_ptr host_pools[GGML_SYCL_MAX_DEVICES]; + std::vector mmid_row_mapping_host; + static std::unique_ptr new_pool_for_device(queue_ptr qptr, int device); static std::unique_ptr new_pool_for_host(queue_ptr qptr, int device); diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index f60192f00258..56821c7ce442 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4007,11 +4007,6 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor } -struct mmid_row_mapping { - int32_t i1; - int32_t i2; -}; - __dpct_inline__ static void k_copy_src1_to_contiguous( const char *__restrict__ src1_original, char *__restrict__ src1_contiguous, const mmid_row_mapping *__restrict__ row_mapping, @@ -4223,7 +4218,7 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, // where each expert's slice starts and the previous ends (row indices, right-exclusive) std::vector expert_row_offsets; // the sources (slot/token pairs) of contiguous rows to guide k_copy_src1_to_contiguous - std::vector routed_row_src; + std::vector & routed_row_src = ctx.mmid_row_mapping_host; mmid_counting_sort_rows(ids, ids_host.data(), n_ids, n_as, n_routed_rows, expert_row_counts, expert_row_offsets, routed_row_src); @@ -4231,8 +4226,6 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, ggml_sycl_pool_alloc dev_row_mapping(ctx.pool(), n_routed_rows); SYCL_CHECK(CHECK_TRY_ERROR( stream->memcpy(dev_row_mapping.get(), routed_row_src.data(), n_routed_rows*sizeof(mmid_row_mapping)))); - // routed_row_src is a host stack vector; wait for the the above async copy - SYCL_CHECK(CHECK_TRY_ERROR(stream->wait())); const unsigned int max_work_group_size = ggml_sycl_info().max_work_group_sizes[ctx.device]; assert(max_work_group_size % (WARP_SIZE * WARP_SIZE) == 0); From b95b31d66edec1cbac9e51a06ff994fc6e6b2a77 Mon Sep 17 00:00:00 2001 From: Alexey Kopytko Date: Tue, 16 Jun 2026 14:04:29 +0900 Subject: [PATCH 3/6] comment on stream->wait --- ggml/src/ggml-sycl/ggml-sycl.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 56821c7ce442..09bd8bfa9013 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4161,6 +4161,8 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, SYCL_CHECK(CHECK_TRY_ERROR( stream->memcpy(ids_host.data(), ids_dev, ggml_nbytes(ids)))); + + // protects few other stream operations from the previous invocation SYCL_CHECK(CHECK_TRY_ERROR(stream->wait())); ggml_tensor src0_row = *src0; From 3d763464fa4f1aeee2a2fd39d89f49ec39c63acb Mon Sep 17 00:00:00 2001 From: Alexey Kopytko Date: Tue, 16 Jun 2026 14:06:02 +0900 Subject: [PATCH 4/6] Apply suggestion from @sanmai --- ggml/src/ggml-sycl/ggml-sycl.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 09bd8bfa9013..f88f5cb93e67 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4162,7 +4162,7 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, SYCL_CHECK(CHECK_TRY_ERROR( stream->memcpy(ids_host.data(), ids_dev, ggml_nbytes(ids)))); - // protects few other stream operations from the previous invocation + // also protects ctx.mmid_row_mapping_host from reuse SYCL_CHECK(CHECK_TRY_ERROR(stream->wait())); ggml_tensor src0_row = *src0; From 1dacb489d3d0e5d37ebc435a7b483356dd57993d Mon Sep 17 00:00:00 2001 From: Alexey Kopytko Date: Tue, 16 Jun 2026 14:20:33 +0900 Subject: [PATCH 5/6] Apply suggestion from @sanmai --- ggml/src/ggml-sycl/ggml-sycl.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index f88f5cb93e67..cced7e65fb49 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4162,7 +4162,7 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, SYCL_CHECK(CHECK_TRY_ERROR( stream->memcpy(ids_host.data(), ids_dev, ggml_nbytes(ids)))); - // also protects ctx.mmid_row_mapping_host from reuse + // also protects ctx.mmid_row_mapping_host from an overwrite SYCL_CHECK(CHECK_TRY_ERROR(stream->wait())); ggml_tensor src0_row = *src0; From 7162e42867f69fbfafb7b1e6aca02458ee654ab8 Mon Sep 17 00:00:00 2001 From: Alexey Kopytko Date: Tue, 16 Jun 2026 14:22:07 +0900 Subject: [PATCH 6/6] Apply suggestion from @sanmai --- ggml/src/ggml-sycl/ggml-sycl.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index cced7e65fb49..0839a05025eb 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4162,7 +4162,7 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, SYCL_CHECK(CHECK_TRY_ERROR( stream->memcpy(ids_host.data(), ids_dev, ggml_nbytes(ids)))); - // also protects ctx.mmid_row_mapping_host from an overwrite + // also ensures ctx.mmid_row_mapping_host is drained before we use it again SYCL_CHECK(CHECK_TRY_ERROR(stream->wait())); ggml_tensor src0_row = *src0;