From df1e23a68503f86a8f08b8eecd7797414a2906b7 Mon Sep 17 00:00:00 2001 From: Hrishith Thadicherla Date: Thu, 1 Oct 2026 03:22:51 -0700 Subject: [PATCH 1/5] llama : avoid repeated synchronization of ready outputs Assisted-by: OpenAI Codex --- common/sampling.cpp | 4 ++-- src/llama-context.cpp | 38 +++++++++++++++++++++++--------------- src/llama-context.h | 4 ++++ 3 files changed, 29 insertions(+), 17 deletions(-) diff --git a/common/sampling.cpp b/common/sampling.cpp index e9e1cb372eb9..9b75a33df40e 100644 --- a/common/sampling.cpp +++ b/common/sampling.cpp @@ -592,7 +592,7 @@ struct llama_sampler * common_sampler_get(const struct common_sampler * gsmpl) { } llama_token common_sampler_sample(struct common_sampler * gsmpl, struct llama_context * ctx, int idx, bool grammar_first) { - llama_synchronize(ctx); + const llama_token backend_id = llama_get_sampled_token_ith(ctx, idx); // start measuring sampling time after the llama_context synchronization in order to not measure any ongoing async operations const auto tm = gsmpl->tm(); @@ -609,7 +609,7 @@ llama_token common_sampler_sample(struct common_sampler * gsmpl, struct llama_co // Check if a backend sampler has already sampled a token in which case we // return that token id directly. { - id = llama_get_sampled_token_ith(ctx, idx); + id = backend_id; if (id != LLAMA_TOKEN_NULL) { LOG_DBG("%s: Backend sampler selected token: '%d'. Will not run any CPU samplers\n", __func__, id); diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 4f04db0ee67f..8cb14dd9e789 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -772,6 +772,7 @@ void llama_context::synchronize() { } ggml_backend_sched_synchronize(sched.get()); + output_pending = false; // FIXME: if multiple single tokens are evaluated without a synchronization, // the stats will be added to the prompt evaluation stats @@ -800,6 +801,12 @@ void llama_context::synchronize() { t_compute_start_us = 0; } +void llama_context::output_synchronize() { + if (output_pending) { + synchronize(); + } +} + const llama_model & llama_context::get_model() const { return model; } @@ -2623,6 +2630,7 @@ ggml_status llama_context::graph_compute( } auto status = ggml_backend_sched_graph_compute_async(sched.get(), gf); + output_pending = true; if (status != GGML_STATUS_SUCCESS) { LLAMA_LOG_ERROR("%s: ggml_backend_sched_graph_compute_async failed with error %d\n", __func__, status); } @@ -3993,13 +4001,13 @@ void llama_synchronize(llama_context * ctx) { } float * llama_get_logits(llama_context * ctx) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_logits(); } float * llama_get_logits_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); float * res = nullptr; @@ -4013,19 +4021,19 @@ float * llama_get_logits_ith(llama_context * ctx, int32_t i) { } float * llama_get_embeddings(llama_context * ctx) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_embeddings(); } float * llama_get_embeddings_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_embeddings_ith(i); } float * llama_get_embeddings_seq(llama_context * ctx, llama_seq_id seq_id) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_embeddings_seq(seq_id); } @@ -4051,19 +4059,19 @@ llama_memory_t llama_get_memory(const struct llama_context * ctx) { } float * llama_get_embeddings_nextn(llama_context * ctx) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_embeddings_nextn(); } float * llama_get_embeddings_nextn_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_embeddings_nextn_ith(i); } float * llama_get_embeddings_layer_inp(llama_context * ctx, uint32_t lid) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_embeddings_layer_inp(lid); } @@ -4073,43 +4081,43 @@ bool llama_set_sampler(llama_context * ctx, llama_seq_id seq_id, llama_sampler * } llama_token llama_get_sampled_token_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_sampled_token_ith(i); } float * llama_get_sampled_probs_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_sampled_probs_ith(i); } float * llama_get_sampled_logits_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return ctx->get_sampled_logits_ith(i); } llama_token * llama_get_sampled_candidates_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return const_cast(ctx->get_sampled_candidates_ith(i)); } uint32_t llama_get_sampled_candidates_count_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return static_cast(ctx->get_sampled_candidates_count(i)); } uint32_t llama_get_sampled_logits_count_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return static_cast(ctx->get_sampled_logits_count(i)); } uint32_t llama_get_sampled_probs_count_ith(llama_context * ctx, int32_t i) { - ctx->synchronize(); + ctx->output_synchronize(); return static_cast(ctx->get_sampled_probs_count(i)); } diff --git a/src/llama-context.h b/src/llama-context.h index 95b7226288b4..bea2a6b5d50b 100644 --- a/src/llama-context.h +++ b/src/llama-context.h @@ -57,6 +57,7 @@ struct llama_context { void sched_reserve(); void synchronize(); + void output_synchronize(); const llama_model & get_model() const; const llama_cparams & get_cparams() const; @@ -356,6 +357,9 @@ struct llama_context { bool sched_need_reserve = true; + // Output getters share completion of the last graph and its host copies. + bool output_pending = false; + ggml_backend_t backend_cpu = nullptr; std::vector backends; From 30256a27e532dba9eb0c75e5041ac3dc926a6d63 Mon Sep 17 00:00:00 2001 From: Hrishith Thadicherla Date: Thu, 1 Oct 2026 06:47:36 -0700 Subject: [PATCH 2/5] llama : centralize pending-work synchronization Keep timing and token accounting outside the pending-work guard. Clear the flag after pipeline synchronization and restore the original getter and sampling call sites. Assisted-by: OpenAI Codex --- common/sampling.cpp | 4 ++-- src/llama-context.cpp | 43 ++++++++++++++++++++----------------------- src/llama-context.h | 3 +-- 3 files changed, 23 insertions(+), 27 deletions(-) diff --git a/common/sampling.cpp b/common/sampling.cpp index 9b75a33df40e..e9e1cb372eb9 100644 --- a/common/sampling.cpp +++ b/common/sampling.cpp @@ -592,7 +592,7 @@ struct llama_sampler * common_sampler_get(const struct common_sampler * gsmpl) { } llama_token common_sampler_sample(struct common_sampler * gsmpl, struct llama_context * ctx, int idx, bool grammar_first) { - const llama_token backend_id = llama_get_sampled_token_ith(ctx, idx); + llama_synchronize(ctx); // start measuring sampling time after the llama_context synchronization in order to not measure any ongoing async operations const auto tm = gsmpl->tm(); @@ -609,7 +609,7 @@ llama_token common_sampler_sample(struct common_sampler * gsmpl, struct llama_co // Check if a backend sampler has already sampled a token in which case we // return that token id directly. { - id = backend_id; + id = llama_get_sampled_token_ith(ctx, idx); if (id != LLAMA_TOKEN_NULL) { LOG_DBG("%s: Backend sampler selected token: '%d'. Will not run any CPU samplers\n", __func__, id); diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 8cb14dd9e789..176989aaa38b 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -771,8 +771,10 @@ void llama_context::synchronize() { return; } - ggml_backend_sched_synchronize(sched.get()); - output_pending = false; + if (output_pending) { + ggml_backend_sched_synchronize(sched.get()); + output_pending = false; + } // FIXME: if multiple single tokens are evaluated without a synchronization, // the stats will be added to the prompt evaluation stats @@ -801,12 +803,6 @@ void llama_context::synchronize() { t_compute_start_us = 0; } -void llama_context::output_synchronize() { - if (output_pending) { - synchronize(); - } -} - const llama_model & llama_context::get_model() const { return model; } @@ -1424,6 +1420,7 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll // that the previous compute is still reading. if (cparams.pipeline_parallel) { ggml_backend_sched_synchronize(sched.get()); + output_pending = false; } n_reused++; @@ -4001,13 +3998,13 @@ void llama_synchronize(llama_context * ctx) { } float * llama_get_logits(llama_context * ctx) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_logits(); } float * llama_get_logits_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); float * res = nullptr; @@ -4021,19 +4018,19 @@ float * llama_get_logits_ith(llama_context * ctx, int32_t i) { } float * llama_get_embeddings(llama_context * ctx) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_embeddings(); } float * llama_get_embeddings_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_embeddings_ith(i); } float * llama_get_embeddings_seq(llama_context * ctx, llama_seq_id seq_id) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_embeddings_seq(seq_id); } @@ -4059,19 +4056,19 @@ llama_memory_t llama_get_memory(const struct llama_context * ctx) { } float * llama_get_embeddings_nextn(llama_context * ctx) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_embeddings_nextn(); } float * llama_get_embeddings_nextn_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_embeddings_nextn_ith(i); } float * llama_get_embeddings_layer_inp(llama_context * ctx, uint32_t lid) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_embeddings_layer_inp(lid); } @@ -4081,43 +4078,43 @@ bool llama_set_sampler(llama_context * ctx, llama_seq_id seq_id, llama_sampler * } llama_token llama_get_sampled_token_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_sampled_token_ith(i); } float * llama_get_sampled_probs_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_sampled_probs_ith(i); } float * llama_get_sampled_logits_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return ctx->get_sampled_logits_ith(i); } llama_token * llama_get_sampled_candidates_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return const_cast(ctx->get_sampled_candidates_ith(i)); } uint32_t llama_get_sampled_candidates_count_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return static_cast(ctx->get_sampled_candidates_count(i)); } uint32_t llama_get_sampled_logits_count_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return static_cast(ctx->get_sampled_logits_count(i)); } uint32_t llama_get_sampled_probs_count_ith(llama_context * ctx, int32_t i) { - ctx->output_synchronize(); + ctx->synchronize(); return static_cast(ctx->get_sampled_probs_count(i)); } diff --git a/src/llama-context.h b/src/llama-context.h index bea2a6b5d50b..9b0a840703ea 100644 --- a/src/llama-context.h +++ b/src/llama-context.h @@ -57,7 +57,6 @@ struct llama_context { void sched_reserve(); void synchronize(); - void output_synchronize(); const llama_model & get_model() const; const llama_cparams & get_cparams() const; @@ -357,7 +356,7 @@ struct llama_context { bool sched_need_reserve = true; - // Output getters share completion of the last graph and its host copies. + // An async graph or its host output copies may still be pending. bool output_pending = false; ggml_backend_t backend_cpu = nullptr; From 3bec97917ba7b8ff5319a7e8ecc145b5d97457c7 Mon Sep 17 00:00:00 2001 From: Hrishith Thadicherla Date: Thu, 1 Oct 2026 07:29:24 -0700 Subject: [PATCH 3/5] ggml: track synchronization at backend submission boundaries Assisted-by: OpenAI Codex --- ggml/src/ggml-backend-impl.h | 9 ++- ggml/src/ggml-backend.cpp | 42 ++++++++-- ggml/src/ggml-cuda/ggml-cuda.cu | 3 + ggml/src/ggml-sycl/ggml-sycl.cpp | 3 + .../backend/backend-dispatched-backend.cpp | 4 +- src/llama-context.cpp | 7 +- src/llama-context.h | 3 - tests/test-alloc.cpp | 79 +++++++++++++++++++ 8 files changed, 131 insertions(+), 19 deletions(-) diff --git a/ggml/src/ggml-backend-impl.h b/ggml/src/ggml-backend-impl.h index ef05905cf9ab..bf9b2df64bd6 100644 --- a/ggml/src/ggml-backend-impl.h +++ b/ggml/src/ggml-backend-impl.h @@ -8,7 +8,7 @@ extern "C" { #endif - #define GGML_BACKEND_API_VERSION 2 + #define GGML_BACKEND_API_VERSION 3 // // Backend buffer type @@ -160,6 +160,13 @@ extern "C" { struct ggml_backend_i iface; ggml_backend_dev_t device; void * context; + + // Async submissions clear this; only a full backend synchronization sets it. +#ifdef __cplusplus + bool is_synchronized = false; +#else + bool is_synchronized; +#endif }; struct ggml_backend_event { diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index 273dc92b2571..fd17fa4e9237 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -273,6 +273,7 @@ void ggml_backend_tensor_set_async(ggml_backend_t backend, struct ggml_tensor * ggml_backend_tensor_set(tensor, data, offset, size); } else { backend->iface.set_tensor_async(backend, tensor, data, offset, size); + backend->is_synchronized = false; } } @@ -287,6 +288,7 @@ void ggml_backend_tensor_get_async(ggml_backend_t backend, const struct ggml_ten ggml_backend_tensor_get(tensor, data, offset, size); } else { backend->iface.get_tensor_async(backend, tensor, data, offset, size); + backend->is_synchronized = false; } } @@ -309,6 +311,7 @@ void ggml_backend_tensor_set_2d_async(ggml_backend_t backend, struct ggml_tensor GGML_ASSERT(tensor->data != NULL && "tensor not allocated"); GGML_ASSERT(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && "tensor write out of bounds"); backend->iface.set_tensor_2d_async(backend, tensor, data, offset, size, n_copies, stride_tensor, stride_data); + backend->is_synchronized = false; } void ggml_backend_tensor_get_2d_async(ggml_backend_t backend, const struct ggml_tensor * tensor, void * data, size_t offset, size_t size, @@ -330,6 +333,7 @@ void ggml_backend_tensor_get_2d_async(ggml_backend_t backend, const struct ggml_ GGML_ASSERT(tensor->data != NULL && "tensor not allocated"); GGML_ASSERT(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && "tensor read out of bounds"); backend->iface.get_tensor_2d_async(backend, tensor, data, offset, size, n_copies, stride_tensor, stride_data); + backend->is_synchronized = false; } void ggml_backend_tensor_set(struct ggml_tensor * tensor, const void * data, size_t offset, size_t size) { @@ -428,7 +432,13 @@ void ggml_backend_synchronize(ggml_backend_t backend) { return; } + // Meta backends must also check work submitted directly to their child backends. + if (backend->is_synchronized && !ggml_backend_is_meta(backend)) { + return; + } + backend->iface.synchronize(backend); + backend->is_synchronized = true; } ggml_backend_graph_plan_t ggml_backend_graph_plan_create(ggml_backend_t backend, struct ggml_cgraph * cgraph) { @@ -449,7 +459,10 @@ enum ggml_status ggml_backend_graph_plan_compute(ggml_backend_t backend, ggml_ba GGML_ASSERT(backend); GGML_ASSERT(backend->iface.graph_plan_compute != NULL); - return backend->iface.graph_plan_compute(backend, plan); + const auto status = backend->iface.graph_plan_compute(backend, plan); + // A failed submission may still have queued work. + backend->is_synchronized = false; + return status; } enum ggml_status ggml_backend_graph_compute(ggml_backend_t backend, struct ggml_cgraph * cgraph) { @@ -460,7 +473,10 @@ enum ggml_status ggml_backend_graph_compute(ggml_backend_t backend, struct ggml_ enum ggml_status ggml_backend_graph_compute_async(ggml_backend_t backend, struct ggml_cgraph * cgraph) { GGML_ASSERT(backend); - return backend->iface.graph_compute(backend, cgraph); + const auto status = backend->iface.graph_compute(backend, cgraph); + // A failed submission may still have queued work. + backend->is_synchronized = false; + return status; } bool ggml_backend_supports_op(ggml_backend_t backend, const struct ggml_tensor * op) { @@ -508,6 +524,18 @@ void ggml_backend_tensor_copy(const struct ggml_tensor * src, struct ggml_tensor } } +static bool ggml_backend_tensor_try_copy_async(ggml_backend_t backend_src, ggml_backend_t backend_dst, const struct ggml_tensor * src, struct ggml_tensor * dst) { + if (backend_dst->iface.cpy_tensor_async == NULL) { + return false; + } + + const bool copied = backend_dst->iface.cpy_tensor_async(backend_src, backend_dst, src, dst); + // A cross-backend copy may queue work on both streams. + backend_src->is_synchronized = false; + backend_dst->is_synchronized = false; + return copied; +} + void ggml_backend_tensor_copy_async(ggml_backend_t backend_src, ggml_backend_t backend_dst, const struct ggml_tensor * src, struct ggml_tensor * dst) { GGML_ASSERT(ggml_are_same_layout(src, dst) && "cannot copy tensors with different layouts"); @@ -516,10 +544,8 @@ void ggml_backend_tensor_copy_async(ggml_backend_t backend_src, ggml_backend_t b } GGML_ASSERT(backend_dst); - if (backend_dst->iface.cpy_tensor_async != NULL) { - if (backend_dst->iface.cpy_tensor_async(backend_src, backend_dst, src, dst)) { - return; - } + if (ggml_backend_tensor_try_copy_async(backend_src, backend_dst, src, dst)) { + return; } // an async copy would normally happen after all the queued operations on both backends are completed @@ -551,6 +577,7 @@ void ggml_backend_event_record(ggml_backend_event_t event, ggml_backend_t backen GGML_ASSERT(backend->iface.event_record != NULL); backend->iface.event_record(backend, event); + backend->is_synchronized = false; } void ggml_backend_event_synchronize(ggml_backend_event_t event) { @@ -565,6 +592,7 @@ void ggml_backend_event_wait(ggml_backend_t backend, ggml_backend_event_t event) GGML_ASSERT(backend->iface.event_wait != NULL); backend->iface.event_wait(backend, event); + backend->is_synchronized = false; } static void ggml_backend_graph_optimize(ggml_backend_t backend, struct ggml_cgraph * cgraph, struct ggml_backend_graph_optimize_params * params) { @@ -1798,7 +1826,7 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } else { // try async copy, but if not possible, we can still use a sync copy without synchronizing the dst backend, since we handle the synchronization here with multiple copies and events // TODO: add public function to facilitate this, since applications do not have direct access to the backend interface - if (!split_backend->iface.cpy_tensor_async || !split_backend->iface.cpy_tensor_async(input_backend, split_backend, input, input_cpy)) { + if (!ggml_backend_tensor_try_copy_async(input_backend, split_backend, input, input_cpy)) { ggml_backend_synchronize(input_backend); if (sched->events[split_backend_id][sched->cur_copy] != NULL) { ggml_backend_event_synchronize(sched->events[split_backend_id][sched->cur_copy]); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index ebbb5c10615c..59eac909756b 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1249,6 +1249,9 @@ static bool ggml_backend_cuda_comm_allreduce_tensor(void * comm_ctx_v, struct gg return false; } auto * comm_ctx = static_cast(comm_ctx_v); + for (ggml_backend_t backend : comm_ctx->backends) { + backend->is_synchronized = false; + } return comm_ctx->try_allreduce(comm_ctx, tensors); } diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 2bc2aaa326f4..4540fe381688 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -7088,6 +7088,9 @@ bool ggml_backend_sycl_comm_allreduce_tensor(void * comm_ctx_v, struct ggml_tens } auto * comm_ctx = static_cast(comm_ctx_v); + for (ggml_backend_t backend : comm_ctx->backends) { + backend->is_synchronized = false; + } const size_t n_backends = comm_ctx->backends.size(); // Fast path: N=2, F32/F16, contiguous, matching shapes. diff --git a/ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp b/ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp index 03a037f1cbd9..6cd54675278a 100644 --- a/ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp +++ b/ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp @@ -90,10 +90,10 @@ uint32_t backend_backend_graph_compute(apir_encoder * enc, apir_decoder * dec, v return 1; } - status = bck->iface.graph_compute(bck, cgraph); + status = ggml_backend_graph_compute_async(bck, cgraph); if (async_backend && bck->iface.synchronize) { - bck->iface.synchronize(bck); + ggml_backend_synchronize(bck); } apir_encode_ggml_status(enc, &status); diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 176989aaa38b..4f04db0ee67f 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -771,10 +771,7 @@ void llama_context::synchronize() { return; } - if (output_pending) { - ggml_backend_sched_synchronize(sched.get()); - output_pending = false; - } + ggml_backend_sched_synchronize(sched.get()); // FIXME: if multiple single tokens are evaluated without a synchronization, // the stats will be added to the prompt evaluation stats @@ -1420,7 +1417,6 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll // that the previous compute is still reading. if (cparams.pipeline_parallel) { ggml_backend_sched_synchronize(sched.get()); - output_pending = false; } n_reused++; @@ -2627,7 +2623,6 @@ ggml_status llama_context::graph_compute( } auto status = ggml_backend_sched_graph_compute_async(sched.get(), gf); - output_pending = true; if (status != GGML_STATUS_SUCCESS) { LLAMA_LOG_ERROR("%s: ggml_backend_sched_graph_compute_async failed with error %d\n", __func__, status); } diff --git a/src/llama-context.h b/src/llama-context.h index 9b0a840703ea..95b7226288b4 100644 --- a/src/llama-context.h +++ b/src/llama-context.h @@ -356,9 +356,6 @@ struct llama_context { bool sched_need_reserve = true; - // An async graph or its host output copies may still be pending. - bool output_pending = false; - ggml_backend_t backend_cpu = nullptr; std::vector backends; diff --git a/tests/test-alloc.cpp b/tests/test-alloc.cpp index 8f1a98aa03c3..12bdcc68308c 100644 --- a/tests/test-alloc.cpp +++ b/tests/test-alloc.cpp @@ -650,6 +650,84 @@ static void test_graph_optimize_alloc_dep() { GGML_ASSERT(!graph_reuses_allocation(true)); } +// Check that every async entry point invalidates a completed synchronization. +static void test_backend_synchronize() { + int n_sync = 0; + ggml_backend backend = {}; + backend.context = &n_sync; + backend.iface.synchronize = [](ggml_backend_t b) { ++*(int *) b->context; }; + backend.iface.set_tensor_async = [](ggml_backend_t, ggml_tensor *, const void *, size_t, size_t) {}; + backend.iface.get_tensor_async = [](ggml_backend_t, const ggml_tensor *, void *, size_t, size_t) {}; + backend.iface.set_tensor_2d_async = [](ggml_backend_t, ggml_tensor *, const void *, size_t, size_t, size_t, size_t, size_t) {}; + backend.iface.get_tensor_2d_async = [](ggml_backend_t, const ggml_tensor *, void *, size_t, size_t, size_t, size_t, size_t) {}; + backend.iface.graph_compute = [](ggml_backend_t, ggml_cgraph *) { return GGML_STATUS_SUCCESS; }; + backend.iface.graph_plan_compute = [](ggml_backend_t, ggml_backend_graph_plan_t) { return GGML_STATUS_SUCCESS; }; + backend.iface.event_record = [](ggml_backend_t, ggml_backend_event_t) {}; + backend.iface.event_wait = [](ggml_backend_t, ggml_backend_event_t) {}; + backend.iface.cpy_tensor_async = [](ggml_backend_t, ggml_backend_t, const ggml_tensor *, ggml_tensor *) { return true; }; + + auto check_sync = [&]() { + const int before = n_sync; + ggml_backend_synchronize(&backend); + GGML_ASSERT(n_sync == before + 1); + ggml_backend_synchronize(&backend); + GGML_ASSERT(n_sync == before + 1); + }; + check_sync(); + + float data[2] = {}; + ggml_tensor tensor = {}; + tensor.type = GGML_TYPE_F32; + tensor.ne[0] = 2; + tensor.nb[0] = sizeof(float); + for (int i = 1; i < GGML_MAX_DIMS; ++i) { + tensor.ne[i] = 1; + tensor.nb[i] = sizeof(data); + } + tensor.data = data; + + ggml_backend_graph_compute_async(&backend, nullptr); + check_sync(); + // A copy after a completed graph must still be synchronized. + ggml_backend_tensor_get_async(&backend, &tensor, data, 0, sizeof(data)); + check_sync(); + ggml_backend_tensor_set_async(&backend, &tensor, data, 0, sizeof(data)); + check_sync(); + ggml_backend_tensor_get_2d_async(&backend, &tensor, data, 0, sizeof(float), 2, sizeof(float), sizeof(float)); + check_sync(); + ggml_backend_tensor_set_2d_async(&backend, &tensor, data, 0, sizeof(float), 2, sizeof(float), sizeof(float)); + check_sync(); + // Also cover the fallback through the one-dimensional callbacks. + backend.iface.get_tensor_2d_async = nullptr; + backend.iface.set_tensor_2d_async = nullptr; + ggml_backend_tensor_get_2d_async(&backend, &tensor, data, 0, sizeof(float), 2, sizeof(float), sizeof(float)); + check_sync(); + ggml_backend_tensor_set_2d_async(&backend, &tensor, data, 0, sizeof(float), 2, sizeof(float), sizeof(float)); + check_sync(); + ggml_backend_graph_plan_compute(&backend, nullptr); + check_sync(); + // Failure does not imply that no work was submitted. + backend.iface.graph_compute = [](ggml_backend_t, ggml_cgraph *) { return GGML_STATUS_ABORTED; }; + GGML_ASSERT(ggml_backend_graph_compute_async(&backend, nullptr) == GGML_STATUS_ABORTED); + check_sync(); + ggml_backend_event_record(nullptr, &backend); + check_sync(); + ggml_backend_event_wait(&backend, nullptr); + check_sync(); + + int n_sync_dst = 0; + ggml_backend dst_backend = backend; + dst_backend.context = &n_sync_dst; + ggml_tensor dst = tensor; + ggml_backend_tensor_copy_async(&backend, &dst_backend, &tensor, &dst); + check_sync(); + ggml_backend_synchronize(&dst_backend); + ggml_backend_synchronize(&dst_backend); + GGML_ASSERT(n_sync_dst == 1); + ggml_backend_tensor_copy_async(&backend, &backend, &tensor, &dst); + check_sync(); +} + static void run(const char * name, void (*f)()) { printf("%s ", name); fflush(stdout); @@ -658,6 +736,7 @@ static void run(const char * name, void (*f)()) { } int main() { + run("test_backend_synchronize", test_backend_synchronize); run("test_max_size_too_many_tensors", test_max_size_too_many_tensors); run("test_max_size_tensor_too_large", test_max_size_tensor_too_large); run("test_tensor_larger_than_max_size", test_tensor_larger_than_max_size); From 2ce37f0e99b7c07c57f477be0cb2d442d26cf1ac Mon Sep 17 00:00:00 2001 From: Hrishith Thadicherla Date: Thu, 1 Oct 2026 08:45:50 -0700 Subject: [PATCH 4/5] ggml: opt in to synchronization tracking only for CUDA Preserve unconditional synchronization for backends that have not opted in. Restore the SYCL and virtgpu paths and check the default fallback in test-alloc. Assisted-by: OpenAI Codex --- ggml/src/ggml-backend-impl.h | 3 +++ ggml/src/ggml-backend.cpp | 3 +-- ggml/src/ggml-cuda/ggml-cuda.cu | 4 ++++ ggml/src/ggml-sycl/ggml-sycl.cpp | 3 --- .../ggml-virtgpu/backend/backend-dispatched-backend.cpp | 4 ++-- tests/test-alloc.cpp | 7 +++++++ 6 files changed, 17 insertions(+), 7 deletions(-) diff --git a/ggml/src/ggml-backend-impl.h b/ggml/src/ggml-backend-impl.h index bf9b2df64bd6..1eeeeee66313 100644 --- a/ggml/src/ggml-backend-impl.h +++ b/ggml/src/ggml-backend-impl.h @@ -164,8 +164,11 @@ extern "C" { // Async submissions clear this; only a full backend synchronization sets it. #ifdef __cplusplus bool is_synchronized = false; + // Enable only when all async submission paths invalidate is_synchronized. + bool supports_sync_tracking = false; #else bool is_synchronized; + bool supports_sync_tracking; #endif }; diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index fd17fa4e9237..44a9cea67d61 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -432,8 +432,7 @@ void ggml_backend_synchronize(ggml_backend_t backend) { return; } - // Meta backends must also check work submitted directly to their child backends. - if (backend->is_synchronized && !ggml_backend_is_meta(backend)) { + if (backend->supports_sync_tracking && backend->is_synchronized) { return; } diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 59eac909756b..dabe07011560 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -5890,6 +5890,10 @@ ggml_backend_t ggml_backend_cuda_init(int device) { /* .context = */ ctx, }; +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + cuda_backend->supports_sync_tracking = true; +#endif + return cuda_backend; } diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 4540fe381688..2bc2aaa326f4 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -7088,9 +7088,6 @@ bool ggml_backend_sycl_comm_allreduce_tensor(void * comm_ctx_v, struct ggml_tens } auto * comm_ctx = static_cast(comm_ctx_v); - for (ggml_backend_t backend : comm_ctx->backends) { - backend->is_synchronized = false; - } const size_t n_backends = comm_ctx->backends.size(); // Fast path: N=2, F32/F16, contiguous, matching shapes. diff --git a/ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp b/ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp index 6cd54675278a..03a037f1cbd9 100644 --- a/ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp +++ b/ggml/src/ggml-virtgpu/backend/backend-dispatched-backend.cpp @@ -90,10 +90,10 @@ uint32_t backend_backend_graph_compute(apir_encoder * enc, apir_decoder * dec, v return 1; } - status = ggml_backend_graph_compute_async(bck, cgraph); + status = bck->iface.graph_compute(bck, cgraph); if (async_backend && bck->iface.synchronize) { - ggml_backend_synchronize(bck); + bck->iface.synchronize(bck); } apir_encode_ggml_status(enc, &status); diff --git a/tests/test-alloc.cpp b/tests/test-alloc.cpp index 12bdcc68308c..6b9ff9d10f8f 100644 --- a/tests/test-alloc.cpp +++ b/tests/test-alloc.cpp @@ -666,6 +666,13 @@ static void test_backend_synchronize() { backend.iface.event_wait = [](ggml_backend_t, ggml_backend_event_t) {}; backend.iface.cpy_tensor_async = [](ggml_backend_t, ggml_backend_t, const ggml_tensor *, ggml_tensor *) { return true; }; + // Backends that do not opt in retain unconditional synchronization. + ggml_backend_synchronize(&backend); + ggml_backend_synchronize(&backend); + GGML_ASSERT(n_sync == 2); + backend.supports_sync_tracking = true; + backend.is_synchronized = false; + auto check_sync = [&]() { const int before = n_sync; ggml_backend_synchronize(&backend); From fc92d666c734635b708ca078683ab93b6535aa58 Mon Sep 17 00:00:00 2001 From: Hrishith Thadicherla Date: Tue, 6 Oct 2026 08:43:53 -0700 Subject: [PATCH 5/5] cuda: track synchronization in the stream owner Assisted-by: OpenAI Codex --- ggml/src/ggml-backend-impl.h | 12 +---- ggml/src/ggml-backend.cpp | 41 +++------------- ggml/src/ggml-cuda/common.cuh | 19 ++++++++ ggml/src/ggml-cuda/ggml-cuda.cu | 9 +--- tests/test-alloc.cpp | 86 --------------------------------- 5 files changed, 28 insertions(+), 139 deletions(-) diff --git a/ggml/src/ggml-backend-impl.h b/ggml/src/ggml-backend-impl.h index 1eeeeee66313..ef05905cf9ab 100644 --- a/ggml/src/ggml-backend-impl.h +++ b/ggml/src/ggml-backend-impl.h @@ -8,7 +8,7 @@ extern "C" { #endif - #define GGML_BACKEND_API_VERSION 3 + #define GGML_BACKEND_API_VERSION 2 // // Backend buffer type @@ -160,16 +160,6 @@ extern "C" { struct ggml_backend_i iface; ggml_backend_dev_t device; void * context; - - // Async submissions clear this; only a full backend synchronization sets it. -#ifdef __cplusplus - bool is_synchronized = false; - // Enable only when all async submission paths invalidate is_synchronized. - bool supports_sync_tracking = false; -#else - bool is_synchronized; - bool supports_sync_tracking; -#endif }; struct ggml_backend_event { diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index 44a9cea67d61..273dc92b2571 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -273,7 +273,6 @@ void ggml_backend_tensor_set_async(ggml_backend_t backend, struct ggml_tensor * ggml_backend_tensor_set(tensor, data, offset, size); } else { backend->iface.set_tensor_async(backend, tensor, data, offset, size); - backend->is_synchronized = false; } } @@ -288,7 +287,6 @@ void ggml_backend_tensor_get_async(ggml_backend_t backend, const struct ggml_ten ggml_backend_tensor_get(tensor, data, offset, size); } else { backend->iface.get_tensor_async(backend, tensor, data, offset, size); - backend->is_synchronized = false; } } @@ -311,7 +309,6 @@ void ggml_backend_tensor_set_2d_async(ggml_backend_t backend, struct ggml_tensor GGML_ASSERT(tensor->data != NULL && "tensor not allocated"); GGML_ASSERT(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && "tensor write out of bounds"); backend->iface.set_tensor_2d_async(backend, tensor, data, offset, size, n_copies, stride_tensor, stride_data); - backend->is_synchronized = false; } void ggml_backend_tensor_get_2d_async(ggml_backend_t backend, const struct ggml_tensor * tensor, void * data, size_t offset, size_t size, @@ -333,7 +330,6 @@ void ggml_backend_tensor_get_2d_async(ggml_backend_t backend, const struct ggml_ GGML_ASSERT(tensor->data != NULL && "tensor not allocated"); GGML_ASSERT(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && "tensor read out of bounds"); backend->iface.get_tensor_2d_async(backend, tensor, data, offset, size, n_copies, stride_tensor, stride_data); - backend->is_synchronized = false; } void ggml_backend_tensor_set(struct ggml_tensor * tensor, const void * data, size_t offset, size_t size) { @@ -432,12 +428,7 @@ void ggml_backend_synchronize(ggml_backend_t backend) { return; } - if (backend->supports_sync_tracking && backend->is_synchronized) { - return; - } - backend->iface.synchronize(backend); - backend->is_synchronized = true; } ggml_backend_graph_plan_t ggml_backend_graph_plan_create(ggml_backend_t backend, struct ggml_cgraph * cgraph) { @@ -458,10 +449,7 @@ enum ggml_status ggml_backend_graph_plan_compute(ggml_backend_t backend, ggml_ba GGML_ASSERT(backend); GGML_ASSERT(backend->iface.graph_plan_compute != NULL); - const auto status = backend->iface.graph_plan_compute(backend, plan); - // A failed submission may still have queued work. - backend->is_synchronized = false; - return status; + return backend->iface.graph_plan_compute(backend, plan); } enum ggml_status ggml_backend_graph_compute(ggml_backend_t backend, struct ggml_cgraph * cgraph) { @@ -472,10 +460,7 @@ enum ggml_status ggml_backend_graph_compute(ggml_backend_t backend, struct ggml_ enum ggml_status ggml_backend_graph_compute_async(ggml_backend_t backend, struct ggml_cgraph * cgraph) { GGML_ASSERT(backend); - const auto status = backend->iface.graph_compute(backend, cgraph); - // A failed submission may still have queued work. - backend->is_synchronized = false; - return status; + return backend->iface.graph_compute(backend, cgraph); } bool ggml_backend_supports_op(ggml_backend_t backend, const struct ggml_tensor * op) { @@ -523,18 +508,6 @@ void ggml_backend_tensor_copy(const struct ggml_tensor * src, struct ggml_tensor } } -static bool ggml_backend_tensor_try_copy_async(ggml_backend_t backend_src, ggml_backend_t backend_dst, const struct ggml_tensor * src, struct ggml_tensor * dst) { - if (backend_dst->iface.cpy_tensor_async == NULL) { - return false; - } - - const bool copied = backend_dst->iface.cpy_tensor_async(backend_src, backend_dst, src, dst); - // A cross-backend copy may queue work on both streams. - backend_src->is_synchronized = false; - backend_dst->is_synchronized = false; - return copied; -} - void ggml_backend_tensor_copy_async(ggml_backend_t backend_src, ggml_backend_t backend_dst, const struct ggml_tensor * src, struct ggml_tensor * dst) { GGML_ASSERT(ggml_are_same_layout(src, dst) && "cannot copy tensors with different layouts"); @@ -543,8 +516,10 @@ void ggml_backend_tensor_copy_async(ggml_backend_t backend_src, ggml_backend_t b } GGML_ASSERT(backend_dst); - if (ggml_backend_tensor_try_copy_async(backend_src, backend_dst, src, dst)) { - return; + if (backend_dst->iface.cpy_tensor_async != NULL) { + if (backend_dst->iface.cpy_tensor_async(backend_src, backend_dst, src, dst)) { + return; + } } // an async copy would normally happen after all the queued operations on both backends are completed @@ -576,7 +551,6 @@ void ggml_backend_event_record(ggml_backend_event_t event, ggml_backend_t backen GGML_ASSERT(backend->iface.event_record != NULL); backend->iface.event_record(backend, event); - backend->is_synchronized = false; } void ggml_backend_event_synchronize(ggml_backend_event_t event) { @@ -591,7 +565,6 @@ void ggml_backend_event_wait(ggml_backend_t backend, ggml_backend_event_t event) GGML_ASSERT(backend->iface.event_wait != NULL); backend->iface.event_wait(backend, event); - backend->is_synchronized = false; } static void ggml_backend_graph_optimize(ggml_backend_t backend, struct ggml_cgraph * cgraph, struct ggml_backend_graph_optimize_params * params) { @@ -1825,7 +1798,7 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } else { // try async copy, but if not possible, we can still use a sync copy without synchronizing the dst backend, since we handle the synchronization here with multiple copies and events // TODO: add public function to facilitate this, since applications do not have direct access to the backend interface - if (!ggml_backend_tensor_try_copy_async(input_backend, split_backend, input, input_cpy)) { + if (!split_backend->iface.cpy_tensor_async || !split_backend->iface.cpy_tensor_async(input_backend, split_backend, input, input_cpy)) { ggml_backend_synchronize(input_backend); if (sched->events[split_backend_id][sched->cur_copy] != NULL) { ggml_backend_event_synchronize(sched->events[split_backend_id][sched->cur_copy]); diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 50713c012cce..4a78d0039095 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1448,11 +1448,15 @@ struct ggml_backend_cuda_context { std::string name; cudaEvent_t copy_event = nullptr; +private: cudaStream_t streams[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = { { nullptr } }; + bool stream_used[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {}; + cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {nullptr}; void * cublas_workspaces[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = {nullptr}; size_t cublas_workspace_sizes[GGML_CUDA_MAX_DEVICES] = {0}; +public: int curr_stream_no = 0; #ifdef USE_CUDA_GRAPH @@ -1516,7 +1520,9 @@ struct ggml_backend_cuda_context { ~ggml_backend_cuda_context(); + // Borrow handles for the current backend call, not across synchronize(). cudaStream_t stream(int device, int stream) { + stream_used[device][stream] = true; if (streams[device][stream] == nullptr) { ggml_cuda_set_device(device); CUDA_CHECK(cudaStreamCreateWithFlags(&streams[device][stream], cudaStreamNonBlocking)); @@ -1526,9 +1532,22 @@ struct ggml_backend_cuda_context { cudaStream_t stream() { return stream(device, curr_stream_no); } + void synchronize() { +#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) + // Multi-device execution retains its existing synchronization behavior. + if (ggml_cuda_info().device_count == 1 && !stream_used[device][curr_stream_no]) { + return; + } +#endif + CUDA_CHECK(cudaStreamSynchronize(stream())); + stream_used[device][curr_stream_no] = false; + } + ggml_cuda_stream_context & stream_context() { return concurrent_stream_context; } cublasHandle_t cublas_handle() { + // A cached BLAS handle can submit work without requesting the stream again. + stream(); if (cublas_handles[device][curr_stream_no] == nullptr) { ggml_cuda_set_device(device); CUBLAS_CHECK(cublasCreate(&cublas_handles[device][curr_stream_no])); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index dabe07011560..d4aec9bed9c5 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -1249,9 +1249,6 @@ static bool ggml_backend_cuda_comm_allreduce_tensor(void * comm_ctx_v, struct gg return false; } auto * comm_ctx = static_cast(comm_ctx_v); - for (ggml_backend_t backend : comm_ctx->backends) { - backend->is_synchronized = false; - } return comm_ctx->try_allreduce(comm_ctx, tensors); } @@ -2544,7 +2541,7 @@ static bool ggml_backend_cuda_cpy_tensor_async(ggml_backend_t backend_src, ggml_ static void ggml_backend_cuda_synchronize(ggml_backend_t backend) { ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context; - CUDA_CHECK(cudaStreamSynchronize(cuda_ctx->stream())); + cuda_ctx->synchronize(); GGML_UNUSED(backend); } @@ -5890,10 +5887,6 @@ ggml_backend_t ggml_backend_cuda_init(int device) { /* .context = */ ctx, }; -#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA) - cuda_backend->supports_sync_tracking = true; -#endif - return cuda_backend; } diff --git a/tests/test-alloc.cpp b/tests/test-alloc.cpp index 6b9ff9d10f8f..8f1a98aa03c3 100644 --- a/tests/test-alloc.cpp +++ b/tests/test-alloc.cpp @@ -650,91 +650,6 @@ static void test_graph_optimize_alloc_dep() { GGML_ASSERT(!graph_reuses_allocation(true)); } -// Check that every async entry point invalidates a completed synchronization. -static void test_backend_synchronize() { - int n_sync = 0; - ggml_backend backend = {}; - backend.context = &n_sync; - backend.iface.synchronize = [](ggml_backend_t b) { ++*(int *) b->context; }; - backend.iface.set_tensor_async = [](ggml_backend_t, ggml_tensor *, const void *, size_t, size_t) {}; - backend.iface.get_tensor_async = [](ggml_backend_t, const ggml_tensor *, void *, size_t, size_t) {}; - backend.iface.set_tensor_2d_async = [](ggml_backend_t, ggml_tensor *, const void *, size_t, size_t, size_t, size_t, size_t) {}; - backend.iface.get_tensor_2d_async = [](ggml_backend_t, const ggml_tensor *, void *, size_t, size_t, size_t, size_t, size_t) {}; - backend.iface.graph_compute = [](ggml_backend_t, ggml_cgraph *) { return GGML_STATUS_SUCCESS; }; - backend.iface.graph_plan_compute = [](ggml_backend_t, ggml_backend_graph_plan_t) { return GGML_STATUS_SUCCESS; }; - backend.iface.event_record = [](ggml_backend_t, ggml_backend_event_t) {}; - backend.iface.event_wait = [](ggml_backend_t, ggml_backend_event_t) {}; - backend.iface.cpy_tensor_async = [](ggml_backend_t, ggml_backend_t, const ggml_tensor *, ggml_tensor *) { return true; }; - - // Backends that do not opt in retain unconditional synchronization. - ggml_backend_synchronize(&backend); - ggml_backend_synchronize(&backend); - GGML_ASSERT(n_sync == 2); - backend.supports_sync_tracking = true; - backend.is_synchronized = false; - - auto check_sync = [&]() { - const int before = n_sync; - ggml_backend_synchronize(&backend); - GGML_ASSERT(n_sync == before + 1); - ggml_backend_synchronize(&backend); - GGML_ASSERT(n_sync == before + 1); - }; - check_sync(); - - float data[2] = {}; - ggml_tensor tensor = {}; - tensor.type = GGML_TYPE_F32; - tensor.ne[0] = 2; - tensor.nb[0] = sizeof(float); - for (int i = 1; i < GGML_MAX_DIMS; ++i) { - tensor.ne[i] = 1; - tensor.nb[i] = sizeof(data); - } - tensor.data = data; - - ggml_backend_graph_compute_async(&backend, nullptr); - check_sync(); - // A copy after a completed graph must still be synchronized. - ggml_backend_tensor_get_async(&backend, &tensor, data, 0, sizeof(data)); - check_sync(); - ggml_backend_tensor_set_async(&backend, &tensor, data, 0, sizeof(data)); - check_sync(); - ggml_backend_tensor_get_2d_async(&backend, &tensor, data, 0, sizeof(float), 2, sizeof(float), sizeof(float)); - check_sync(); - ggml_backend_tensor_set_2d_async(&backend, &tensor, data, 0, sizeof(float), 2, sizeof(float), sizeof(float)); - check_sync(); - // Also cover the fallback through the one-dimensional callbacks. - backend.iface.get_tensor_2d_async = nullptr; - backend.iface.set_tensor_2d_async = nullptr; - ggml_backend_tensor_get_2d_async(&backend, &tensor, data, 0, sizeof(float), 2, sizeof(float), sizeof(float)); - check_sync(); - ggml_backend_tensor_set_2d_async(&backend, &tensor, data, 0, sizeof(float), 2, sizeof(float), sizeof(float)); - check_sync(); - ggml_backend_graph_plan_compute(&backend, nullptr); - check_sync(); - // Failure does not imply that no work was submitted. - backend.iface.graph_compute = [](ggml_backend_t, ggml_cgraph *) { return GGML_STATUS_ABORTED; }; - GGML_ASSERT(ggml_backend_graph_compute_async(&backend, nullptr) == GGML_STATUS_ABORTED); - check_sync(); - ggml_backend_event_record(nullptr, &backend); - check_sync(); - ggml_backend_event_wait(&backend, nullptr); - check_sync(); - - int n_sync_dst = 0; - ggml_backend dst_backend = backend; - dst_backend.context = &n_sync_dst; - ggml_tensor dst = tensor; - ggml_backend_tensor_copy_async(&backend, &dst_backend, &tensor, &dst); - check_sync(); - ggml_backend_synchronize(&dst_backend); - ggml_backend_synchronize(&dst_backend); - GGML_ASSERT(n_sync_dst == 1); - ggml_backend_tensor_copy_async(&backend, &backend, &tensor, &dst); - check_sync(); -} - static void run(const char * name, void (*f)()) { printf("%s ", name); fflush(stdout); @@ -743,7 +658,6 @@ static void run(const char * name, void (*f)()) { } int main() { - run("test_backend_synchronize", test_backend_synchronize); run("test_max_size_too_many_tensors", test_max_size_too_many_tensors); run("test_max_size_tensor_too_large", test_max_size_tensor_too_large); run("test_tensor_larger_than_max_size", test_tensor_larger_than_max_size);