From de85989fccc2460153a2536b8068e473eeb0dff5 Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Wed, 25 Mar 2026 22:23:27 +0800 Subject: [PATCH 1/4] ggml-backend: prefetch weights async --- common/arg.cpp | 7 + common/common.cpp | 1 + common/common.h | 1 + ggml/include/ggml-backend.h | 3 + ggml/src/ggml-backend-impl.h | 6 + ggml/src/ggml-backend.cpp | 227 ++++++++++++++++++----- ggml/src/ggml-blas/ggml-blas.cpp | 2 + ggml/src/ggml-cann/ggml-cann.cpp | 2 + ggml/src/ggml-cpu/ggml-cpu.cpp | 2 + ggml/src/ggml-cuda/common.cuh | 3 + ggml/src/ggml-cuda/ggml-cuda.cu | 48 +++++ ggml/src/ggml-opencl/ggml-opencl.cpp | 2 + ggml/src/ggml-openvino/ggml-openvino.cpp | 2 + ggml/src/ggml-rpc/ggml-rpc.cpp | 2 + ggml/src/ggml-sycl/ggml-sycl.cpp | 2 + ggml/src/ggml-webgpu/ggml-webgpu.cpp | 2 + ggml/src/ggml-zdnn/ggml-zdnn.cpp | 34 ++-- ggml/src/ggml-zendnn/ggml-zendnn.cpp | 2 + include/llama.h | 2 + src/llama-context.cpp | 8 +- src/llama-cparams.h | 1 + tools/llama-bench/llama-bench.cpp | 34 +++- 22 files changed, 328 insertions(+), 65 deletions(-) diff --git a/common/arg.cpp b/common/arg.cpp index d844c4360942..ea9fdec62dfd 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -2849,6 +2849,13 @@ common_params_context common_params_parser_init(common_params & params, llama_ex params.no_op_offload = !value; } )); + add_opt(common_arg( + {"-pw", "--prefetch-weights"}, "0|1", + string_format("prefetch weight transfers to overlap CPU->GPU copies with compute (default: %d)", (int) params.prefetch_weights), + [](common_params & params, int value) { + params.prefetch_weights = value != 0; + } + )); add_opt(common_arg( {"--lora"}, "FNAME", "path to LoRA adapter (use comma-separated values to load multiple adapters)", diff --git a/common/common.cpp b/common/common.cpp index 19a6b9dd8017..abc5207587f2 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -1663,6 +1663,7 @@ struct llama_context_params common_context_params_to_llama(const common_params & cparams.op_offload = !params.no_op_offload; cparams.swa_full = params.swa_full; cparams.kv_unified = params.kv_unified; + cparams.prefetch_weights = params.prefetch_weights; cparams.type_k = params.cache_type_k; cparams.type_v = params.cache_type_v; diff --git a/common/common.h b/common/common.h index 53dd37d86aa6..a473a5e4c68b 100644 --- a/common/common.h +++ b/common/common.h @@ -583,6 +583,7 @@ struct common_params { bool warmup = true; // warmup run bool check_tensors = false; // validate tensor data bool no_op_offload = false; // globally disable offload host tensor operations to device + bool prefetch_weights = false; // prefetch weight transfers to overlap CPU->GPU copies with compute bool no_extra_bufts = false; // disable extra buffer types (used for weight repacking) bool no_host = false; // bypass host buffer allowing extra buffers to be used diff --git a/ggml/include/ggml-backend.h b/ggml/include/ggml-backend.h index 2924fdbe9884..f83fdba69af5 100644 --- a/ggml/include/ggml-backend.h +++ b/ggml/include/ggml-backend.h @@ -351,6 +351,9 @@ extern "C" { // Set a callback to be called for each resulting node during graph compute GGML_API void ggml_backend_sched_set_eval_callback(ggml_backend_sched_t sched, ggml_backend_sched_eval_callback callback, void * user_data); + // Enable async weight prefetching to overlap CPU->GPU transfers with compute + GGML_API void ggml_backend_sched_set_prefetch_weights(ggml_backend_sched_t sched, bool enabled); + // // Meta backend // diff --git a/ggml/src/ggml-backend-impl.h b/ggml/src/ggml-backend-impl.h index 9c56ec30c5f1..5c7031d8189a 100644 --- a/ggml/src/ggml-backend-impl.h +++ b/ggml/src/ggml-backend-impl.h @@ -137,6 +137,12 @@ extern "C" { // (optional) sort/optimize the nodes in the graph void (*graph_optimize) (ggml_backend_t backend, struct ggml_cgraph * cgraph); + + // (optional) prefetch tensor data on a separate copy stream for overlap with compute + void (*prefetch_tensor_async)(ggml_backend_t backend, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size); + + // (optional) wait for a pending prefetch to complete + void (*prefetch_event_wait)(ggml_backend_t backend); }; struct ggml_backend { diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index 7f4e252dca39..d0b36bc82f37 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -817,6 +817,7 @@ struct ggml_backend_sched { size_t context_buffer_size; bool op_offload; + bool prefetch_weights; int debug; @@ -1309,6 +1310,47 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra } } + // when prefetch is enabled and a new split is needed, progressively fuse + // consecutive nodes that also require weight offloading to the same backend. + // intermediate nodes without weight offload (e.g. GLU between up and down in MoE) + // are included in the fused range. a max lookahead distance prevents scanning + // too far past the last offloading node. + int fuse_end = i; // inclusive end index in graph + if (sched->prefetch_weights && need_new_split) { + const int max_lookahead = 8; + int last_offload_k = i; + for (int k = i + 1; k < graph->n_nodes; k++) { + struct ggml_tensor * next = graph->nodes[k]; + if (ggml_is_view_op(next->op)) { + continue; + } + if (tensor_backend_id(next) != node_backend_id) { + break; + } + if (k - last_offload_k > max_lookahead) { + break; + } + bool has_weight_offload = false; + for (int j = 0; j < GGML_MAX_SRC; j++) { + struct ggml_tensor * src = next->src[j]; + if (src == NULL) { + continue; + } + if (src->buffer != NULL && src->buffer->usage == GGML_BACKEND_BUFFER_USAGE_WEIGHTS) { + int src_backend_id = tensor_backend_id(src); + if (src_backend_id != node_backend_id && !ggml_backend_sched_buffer_supported(sched, src, node_backend_id)) { + has_weight_offload = true; + break; + } + } + } + if (has_weight_offload) { + fuse_end = k; + last_offload_k = k; + } + } + } + if (node_backend_id != cur_backend_id || need_new_split) { split->i_end = i; i_split++; @@ -1326,59 +1368,67 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra } // find inputs that are not on the same backend - for (int j = 0; j < GGML_MAX_SRC; j++) { - struct ggml_tensor * src = node->src[j]; - if (src == NULL) { + for (int fi = i; fi <= fuse_end; fi++) { + struct ggml_tensor * fnode = graph->nodes[fi]; + if (ggml_is_view_op(fnode->op)) { continue; } + for (int j = 0; j < GGML_MAX_SRC; j++) { + struct ggml_tensor * src = fnode->src[j]; + if (src == NULL) { + continue; + } - size_t src_id = hash_id(src); - const int src_backend_id = sched->hv_tensor_backend_ids[src_id]; - GGML_ASSERT(src_backend_id != -1); // all inputs should be assigned by now - - if (src->flags & GGML_TENSOR_FLAG_INPUT && sched->n_copies > 1) { - if (tensor_id_copy(src_id, src_backend_id, 0) == NULL) { - ggml_backend_t backend = sched->backends[src_backend_id]; - for (int c = 0; c < sched->n_copies; c++) { - struct ggml_tensor * tensor_copy; - if (c == sched->cur_copy) { - tensor_copy = src; // use the original tensor as the current copy - } else { - tensor_copy = ggml_dup_tensor_layout(sched->ctx, src); - ggml_format_name(tensor_copy, "%s#%s#%d", ggml_backend_name(backend), src->name, c); + size_t src_id = hash_id(src); + const int src_backend_id = sched->hv_tensor_backend_ids[src_id]; + GGML_ASSERT(src_backend_id != -1); // all inputs should be assigned by now + + if (src->flags & GGML_TENSOR_FLAG_INPUT && sched->n_copies > 1) { + if (tensor_id_copy(src_id, src_backend_id, 0) == NULL) { + ggml_backend_t backend = sched->backends[src_backend_id]; + for (int c = 0; c < sched->n_copies; c++) { + struct ggml_tensor * tensor_copy; + if (c == sched->cur_copy) { + tensor_copy = src; // use the original tensor as the current copy + } else { + tensor_copy = ggml_dup_tensor_layout(sched->ctx, src); + ggml_format_name(tensor_copy, "%s#%s#%d", ggml_backend_name(backend), src->name, c); + } + ggml_set_input(tensor_copy); + ggml_set_output(tensor_copy); // prevent ggml-alloc from overwriting the tensor + tensor_id_copy(src_id, src_backend_id, c) = tensor_copy; + SET_CAUSE(tensor_copy, "4.cpy"); } - ggml_set_input(tensor_copy); - ggml_set_output(tensor_copy); // prevent ggml-alloc from overwriting the tensor - tensor_id_copy(src_id, src_backend_id, c) = tensor_copy; - SET_CAUSE(tensor_copy, "4.cpy"); + int n_graph_inputs = sched->n_graph_inputs++; + GGML_ASSERT(n_graph_inputs < GGML_SCHED_MAX_SPLIT_INPUTS); + sched->graph_inputs[n_graph_inputs] = src; } - int n_graph_inputs = sched->n_graph_inputs++; - GGML_ASSERT(n_graph_inputs < GGML_SCHED_MAX_SPLIT_INPUTS); - sched->graph_inputs[n_graph_inputs] = src; } - } - if (src_backend_id != cur_backend_id && !ggml_backend_sched_buffer_supported(sched, src, cur_backend_id)) { - // create a copy of the input in the split's backend - if (tensor_id_copy(src_id, cur_backend_id, 0) == NULL) { - ggml_backend_t backend = sched->backends[cur_backend_id]; - for (int c = 0; c < sched->n_copies; c++) { - struct ggml_tensor * tensor_copy = ggml_dup_tensor_layout(sched->ctx, src); - ggml_format_name(tensor_copy, "%s#%s#%d", ggml_backend_name(backend), src->name, c); - if (sched->n_copies > 1) { - ggml_set_input(tensor_copy); - ggml_set_output(tensor_copy); // prevent ggml-alloc from overwriting the tensor + if (src_backend_id != cur_backend_id && !ggml_backend_sched_buffer_supported(sched, src, cur_backend_id)) { + // create a copy of the input in the split's backend + if (tensor_id_copy(src_id, cur_backend_id, 0) == NULL) { + ggml_backend_t backend = sched->backends[cur_backend_id]; + for (int c = 0; c < sched->n_copies; c++) { + struct ggml_tensor * tensor_copy = ggml_dup_tensor_layout(sched->ctx, src); + ggml_format_name(tensor_copy, "%s#%s#%d", ggml_backend_name(backend), src->name, c); + if (sched->n_copies > 1) { + ggml_set_input(tensor_copy); + ggml_set_output(tensor_copy); // prevent ggml-alloc from overwriting the tensor + } + tensor_id_copy(src_id, cur_backend_id, c) = tensor_copy; + SET_CAUSE(tensor_copy, "4.cpy"); } - tensor_id_copy(src_id, cur_backend_id, c) = tensor_copy; - SET_CAUSE(tensor_copy, "4.cpy"); + int n_inputs = split->n_inputs++; + GGML_ASSERT(n_inputs < GGML_SCHED_MAX_SPLIT_INPUTS); + split->inputs[n_inputs] = src; } - int n_inputs = split->n_inputs++; - GGML_ASSERT(n_inputs < GGML_SCHED_MAX_SPLIT_INPUTS); - split->inputs[n_inputs] = src; + fnode->src[j] = tensor_id_copy(src_id, cur_backend_id, sched->cur_copy); } - node->src[j] = tensor_id_copy(src_id, cur_backend_id, sched->cur_copy); } } + + i = fuse_end; } split->i_end = graph->n_nodes; sched->n_splits = i_split + 1; @@ -1399,7 +1449,9 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra sched->prev_leaf_backend_ids = tmp; } - int graph_size = std::max(graph->n_nodes, graph->n_leafs) + sched->n_splits*GGML_SCHED_MAX_SPLIT_INPUTS*2*sched->n_copies; + // extra nodes per split input: 2 (dep + copy), plus 2 more for keepalive nodes when prefetching + const int nodes_per_input = sched->prefetch_weights ? 4 : 2; + int graph_size = std::max(graph->n_nodes, graph->n_leafs) + sched->n_splits*GGML_SCHED_MAX_SPLIT_INPUTS*nodes_per_input*sched->n_copies; // remember the actual graph_size for performing reallocation checks later [GGML_SCHED_DEBUG_REALLOC] sched->debug_prev_graph_size = sched->debug_graph_size; @@ -1444,11 +1496,51 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra graph_copy->nodes[graph_copy->n_nodes++] = input_cpy; } + // prefetch double-buffer: reserve next split's weight copy memory BEFORE compute + // so the allocator doesn't reuse it for intermediates in this split + if (sched->prefetch_weights && i + 1 < sched->n_splits) { + struct ggml_backend_sched_split * next = &sched->splits[i + 1]; + if (next->backend_id == split->backend_id) { + for (int j = 0; j < next->n_inputs; j++) { + struct ggml_tensor * next_input = next->inputs[j]; + if (next_input->buffer != NULL && + ggml_backend_buffer_get_usage(next_input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && + ggml_backend_buffer_is_host(next_input->buffer)) { + const size_t id = hash_id(next_input); + struct ggml_tensor * next_cpy = tensor_id_copy(id, next->backend_id, sched->cur_copy); + assert(graph_copy->size > graph_copy->n_nodes); + struct ggml_tensor * keepalive = ggml_view_tensor(sched->ctx, next_cpy); + keepalive->src[0] = next_cpy; + sched->node_backend_ids[graph_copy->n_nodes] = next->backend_id; + graph_copy->nodes[graph_copy->n_nodes++] = keepalive; + } + } + } + } + for (int j = split->i_start; j < split->i_end; j++) { assert(graph_copy->size > graph_copy->n_nodes); sched->node_backend_ids[graph_copy->n_nodes] = tensor_backend_id(graph->nodes[j]); graph_copy->nodes[graph_copy->n_nodes++] = graph->nodes[j]; } + + // extend current split's weight copies lifetime to here (after next's are allocated above) + if (sched->prefetch_weights) { + for (int j = 0; j < split->n_inputs; j++) { + struct ggml_tensor * input = split->inputs[j]; + if (input->buffer != NULL && + ggml_backend_buffer_get_usage(input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && + ggml_backend_buffer_is_host(input->buffer)) { + const size_t id = hash_id(input); + struct ggml_tensor * curr_cpy = tensor_id_copy(id, split->backend_id, sched->cur_copy); + assert(graph_copy->size > graph_copy->n_nodes); + struct ggml_tensor * keepalive = ggml_view_tensor(sched->ctx, curr_cpy); + keepalive->src[0] = curr_cpy; + sched->node_backend_ids[graph_copy->n_nodes] = split->backend_id; + graph_copy->nodes[graph_copy->n_nodes++] = keepalive; + } + } + } } if (sched->n_copies > 1) { @@ -1555,17 +1647,57 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s std::vector ids; std::vector used_ids; + bool next_weights_prefetched = false; + for (int split_id = 0; split_id < sched->n_splits; split_id++) { struct ggml_backend_sched_split * split = &splits[split_id]; int split_backend_id = split->backend_id; ggml_backend_t split_backend = sched->backends[split_backend_id]; + bool weights_prefetched = next_weights_prefetched; + next_weights_prefetched = false; + + if (sched->prefetch_weights) { + // wait for the previous prefetch to finish before issuing new ones + if (split_backend->iface.prefetch_event_wait) { + split_backend->iface.prefetch_event_wait(split_backend); + } + + // prefetch next split's weights on the copy stream (overlaps with current split's compute) + if (split_backend->iface.prefetch_tensor_async && split_id + 1 < sched->n_splits) { + struct ggml_backend_sched_split * next = &splits[split_id + 1]; + // only prefetch when next split is on the same device (multi-GPU safety) + // and only for CPU->GPU transfers (host buffers) + if (next->backend_id == split_backend_id) { + for (int input_id = 0; input_id < next->n_inputs; input_id++) { + struct ggml_tensor * input = next->inputs[input_id]; + if (input->buffer != NULL && + ggml_backend_buffer_get_usage(input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && + ggml_backend_buffer_is_host(input->buffer)) { + struct ggml_tensor * input_cpy = tensor_copy(input, next->backend_id, sched->cur_copy); + split_backend->iface.prefetch_tensor_async( + split_backend, input_cpy, input->data, 0, ggml_nbytes(input)); + next_weights_prefetched = true; + } + } + } + } + } + // copy the input tensors to the split backend for (int input_id = 0; input_id < split->n_inputs; input_id++) { ggml_backend_t input_backend = ggml_backend_sched_get_tensor_backend(sched, split->inputs[input_id]); struct ggml_tensor * input = split->inputs[input_id]; struct ggml_tensor * input_cpy = tensor_copy(input, split_backend_id, sched->cur_copy); + // skip weight inputs that were already prefetched by the previous split + if (weights_prefetched && + input->buffer != NULL && + ggml_backend_buffer_get_usage(input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && + ggml_backend_buffer_is_host(input->buffer)) { + continue; + } + if (input->flags & GGML_TENSOR_FLAG_INPUT) { // inputs from the user must be copied immediately to prevent the user overwriting the data before the copy is done if (sched->events[split_backend_id][sched->cur_copy] != NULL) { @@ -1588,7 +1720,6 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s ggml_backend_buffer_get_usage(input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && ggml_backend_buffer_is_host(input->buffer) && ( (node->src[0] == input_cpy && node->op == GGML_OP_MUL_MAT_ID) - //|| (node->src[1] == input_cpy && node->op == GGML_OP_ADD_ID) /* GGML_OP_ADD_ID weights are small and not worth splitting */ )) { const int64_t n_expert = node->op == GGML_OP_MUL_MAT_ID ? input->ne[2] : input->ne[1]; @@ -1728,6 +1859,7 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s ggml_backend_event_record(sched->events[split_backend_id][sched->cur_copy], split_backend); } } + } return GGML_STATUS_SUCCESS; @@ -1766,7 +1898,7 @@ ggml_backend_sched_t ggml_backend_sched_new( sched->hv_tensor_copies = (ggml_tensor **) malloc(sched->hash_set.size * sched->n_backends * sched->n_copies * sizeof(struct ggml_tensor *)); const size_t ggml_sched_max_splits = graph_size; // at most there is one split for each node in the graph - const size_t nodes_size = graph_size + ggml_sched_max_splits*GGML_SCHED_MAX_SPLIT_INPUTS*2; + const size_t nodes_size = graph_size + ggml_sched_max_splits*GGML_SCHED_MAX_SPLIT_INPUTS*4; sched->node_backend_ids = (int *) calloc(nodes_size, sizeof(sched->node_backend_ids[0])); sched->leaf_backend_ids = (int *) calloc(nodes_size, sizeof(sched->leaf_backend_ids[0])); sched->prev_node_backend_ids = (int *) calloc(nodes_size, sizeof(sched->prev_node_backend_ids[0])); @@ -1775,7 +1907,7 @@ ggml_backend_sched_t ggml_backend_sched_new( sched->debug_graph_size = 0; sched->debug_prev_graph_size = 0; - sched->context_buffer_size = ggml_sched_max_splits*GGML_SCHED_MAX_SPLIT_INPUTS*2*sizeof(struct ggml_tensor) + ggml_graph_overhead_custom(graph_size, false); + sched->context_buffer_size = ggml_sched_max_splits*GGML_SCHED_MAX_SPLIT_INPUTS*4*sizeof(struct ggml_tensor) + ggml_graph_overhead_custom(graph_size, false); sched->context_buffer = (char *) malloc(sched->context_buffer_size); const int initial_splits_capacity = 16; @@ -1929,6 +2061,11 @@ void ggml_backend_sched_set_eval_callback(ggml_backend_sched_t sched, ggml_backe sched->callback_eval_user_data = user_data; } +void ggml_backend_sched_set_prefetch_weights(ggml_backend_sched_t sched, bool enabled) { + GGML_ASSERT(sched); + sched->prefetch_weights = enabled; +} + int ggml_backend_sched_get_n_splits(ggml_backend_sched_t sched) { GGML_ASSERT(sched); return sched->n_splits; diff --git a/ggml/src/ggml-blas/ggml-blas.cpp b/ggml/src/ggml-blas/ggml-blas.cpp index 9745fa29f5db..997db1f50c39 100644 --- a/ggml/src/ggml-blas/ggml-blas.cpp +++ b/ggml/src/ggml-blas/ggml-blas.cpp @@ -276,6 +276,8 @@ static struct ggml_backend_i blas_backend_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_blas_guid(void) { diff --git a/ggml/src/ggml-cann/ggml-cann.cpp b/ggml/src/ggml-cann/ggml-cann.cpp index 5f51ea3bb3c8..681628c879c7 100644 --- a/ggml/src/ggml-cann/ggml-cann.cpp +++ b/ggml/src/ggml-cann/ggml-cann.cpp @@ -2758,6 +2758,8 @@ static const ggml_backend_i ggml_backend_cann_interface = { /* .event_record = */ ggml_backend_cann_event_record, /* .event_wait = */ ggml_backend_cann_event_wait, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; /** diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index 16cc5116c545..8047898a13b0 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -207,6 +207,8 @@ static const struct ggml_backend_i ggml_backend_cpu_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_cpu_guid(void) { diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index d0a61051cfeb..e84d4cd235ef 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1420,6 +1420,9 @@ struct ggml_backend_cuda_context { int device; std::string name; cudaEvent_t copy_event = nullptr; + cudaEvent_t prefetch_sync_event = nullptr; // signals copy stream that compute is done + cudaEvent_t prefetch_event = nullptr; // signals compute stream that prefetch is done + bool prefetch_pending = false; cudaStream_t streams[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = { { nullptr } }; cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES] = {nullptr}; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 09984d0ee94a..c80699ab7305 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -711,6 +711,12 @@ ggml_backend_cuda_context::~ggml_backend_cuda_context() { if (copy_event != nullptr) { CUDA_CHECK(cudaEventDestroy(copy_event)); } + if (prefetch_sync_event != nullptr) { + CUDA_CHECK(cudaEventDestroy(prefetch_sync_event)); + } + if (prefetch_event != nullptr) { + CUDA_CHECK(cudaEventDestroy(prefetch_event)); + } for (int i = 0; i < GGML_CUDA_MAX_DEVICES; ++i) { for (int j = 0; j < GGML_CUDA_MAX_STREAMS; ++j) { if (streams[i][j] != nullptr) { @@ -4278,9 +4284,49 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, ggml_cuda_graph_evaluate_and_capture(cuda_ctx, cgraph, use_cuda_graph, cuda_graph_update_required, graph_key); + // record event on compute stream so the copy stream can wait for compute to finish + // before starting prefetches that may reuse memory from this split's intermediates + if (cuda_ctx->prefetch_sync_event == nullptr) { + CUDA_CHECK(cudaEventCreateWithFlags(&cuda_ctx->prefetch_sync_event, cudaEventDisableTiming)); + } + CUDA_CHECK(cudaEventRecord(cuda_ctx->prefetch_sync_event, cuda_ctx->stream())); + return GGML_STATUS_SUCCESS; } +static void ggml_backend_cuda_prefetch_tensor_async(ggml_backend_t backend, ggml_tensor * tensor, const void * data, size_t offset, size_t size) { + ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context; + + ggml_cuda_set_device(cuda_ctx->device); + + // use stream 1 as dedicated copy stream for overlap with compute + cudaStream_t copy_stream = cuda_ctx->stream(cuda_ctx->device, 1); + + // wait for previous compute to finish before writing to memory that may alias + // with the previous split's intermediate tensors + if (cuda_ctx->prefetch_sync_event != nullptr) { + CUDA_CHECK(cudaStreamWaitEvent(copy_stream, cuda_ctx->prefetch_sync_event, 0)); + } + + CUDA_CHECK(cudaMemcpyAsync((char *)tensor->data + offset, data, size, cudaMemcpyHostToDevice, copy_stream)); + + // record prefetch event so the scheduler can wait for it + if (cuda_ctx->prefetch_event == nullptr) { + CUDA_CHECK(cudaEventCreateWithFlags(&cuda_ctx->prefetch_event, cudaEventDisableTiming)); + } + CUDA_CHECK(cudaEventRecord(cuda_ctx->prefetch_event, copy_stream)); + cuda_ctx->prefetch_pending = true; +} + +static void ggml_backend_cuda_prefetch_event_wait(ggml_backend_t backend) { + ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context; + if (cuda_ctx->prefetch_pending) { + // compute stream waits for copy stream (prefetched data is ready) + CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx->stream(), cuda_ctx->prefetch_event, 0)); + cuda_ctx->prefetch_pending = false; + } +} + static void ggml_backend_cuda_event_record(ggml_backend_t backend, ggml_backend_event_t event) { ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context; @@ -4567,6 +4613,8 @@ static const ggml_backend_i ggml_backend_cuda_interface = { /* .event_record = */ ggml_backend_cuda_event_record, /* .event_wait = */ ggml_backend_cuda_event_wait, /* .graph_optimize = */ ggml_backend_cuda_graph_optimize, + /* .prefetch_tensor_async = */ ggml_backend_cuda_prefetch_tensor_async, + /* .prefetch_event_wait = */ ggml_backend_cuda_prefetch_event_wait, }; static ggml_guid_t ggml_backend_cuda_guid() { diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 915c8e7b903f..ef8decf92bd1 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -7490,6 +7490,8 @@ static ggml_backend_i ggml_backend_opencl_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; ggml_backend_t ggml_backend_opencl_init(void) { diff --git a/ggml/src/ggml-openvino/ggml-openvino.cpp b/ggml/src/ggml-openvino/ggml-openvino.cpp index 0e7501fefe38..046543bac3f3 100644 --- a/ggml/src/ggml-openvino/ggml-openvino.cpp +++ b/ggml/src/ggml-openvino/ggml-openvino.cpp @@ -653,6 +653,8 @@ static const ggml_backend_i ggml_backend_openvino_interface = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; int ggml_backend_openvino_get_device_count() { diff --git a/ggml/src/ggml-rpc/ggml-rpc.cpp b/ggml/src/ggml-rpc/ggml-rpc.cpp index 17c53a5f049e..e44a1f4f5e79 100644 --- a/ggml/src/ggml-rpc/ggml-rpc.cpp +++ b/ggml/src/ggml-rpc/ggml-rpc.cpp @@ -758,6 +758,8 @@ static ggml_backend_i ggml_backend_rpc_interface = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; ggml_backend_buffer_type_t ggml_backend_rpc_buffer_type(const char * endpoint, uint32_t device) { diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 131bff714f9c..16de8364b73e 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5570,6 +5570,8 @@ static ggml_backend_i ggml_backend_sycl_interface = { /* .event_record = */ ggml_backend_sycl_event_record, /* .event_wait = */ ggml_backend_sycl_event_wait, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_sycl_guid() { diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index c001cda7d116..eec6460a4e8a 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -3613,6 +3613,8 @@ static ggml_backend_i ggml_backend_webgpu_i = { /* .event_record = */ ggml_backend_webgpu_event_record, /* .event_wait = */ ggml_backend_webgpu_event_wait, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; /* End GGML Backend Interface */ diff --git a/ggml/src/ggml-zdnn/ggml-zdnn.cpp b/ggml/src/ggml-zdnn/ggml-zdnn.cpp index 639b818d128e..a550feba7ee8 100644 --- a/ggml/src/ggml-zdnn/ggml-zdnn.cpp +++ b/ggml/src/ggml-zdnn/ggml-zdnn.cpp @@ -419,22 +419,24 @@ static enum ggml_status ggml_backend_zdnn_graph_compute(ggml_backend_t backend, } static ggml_backend_i ggml_backend_zdnn_i = { - /* .get_name = */ ggml_backend_zdnn_name, - /* .free = */ ggml_backend_zdnn_free, - /* .set_tensor_async = */ NULL, - /* .get_tensor_async = */ NULL, - /* .set_tensor_2d_async = */ NULL, - /* .get_tensor_2d_async = */ NULL, - /* .cpy_tensor_async = */ NULL, - /* .synchronize = */ NULL, - /* .graph_plan_create = */ NULL, - /* .graph_plan_free = */ NULL, - /* .graph_plan_update = */ NULL, - /* .graph_plan_compute = */ NULL, - /* .graph_compute = */ ggml_backend_zdnn_graph_compute, - /* .event_record = */ NULL, - /* .event_wait = */ NULL, - /* .graph_optimize = */ NULL, + /* .get_name = */ ggml_backend_zdnn_name, + /* .free = */ ggml_backend_zdnn_free, + /* .set_tensor_async = */ NULL, + /* .get_tensor_async = */ NULL, + /* .set_tensor_2d_async = */ NULL, + /* .get_tensor_2d_async = */ NULL, + /* .cpy_tensor_async = */ NULL, + /* .synchronize = */ NULL, + /* .graph_plan_create = */ NULL, + /* .graph_plan_free = */ NULL, + /* .graph_plan_update = */ NULL, + /* .graph_plan_compute = */ NULL, + /* .graph_compute = */ ggml_backend_zdnn_graph_compute, + /* .event_record = */ NULL, + /* .event_wait = */ NULL, + /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_zdnn_guid(void) { diff --git a/ggml/src/ggml-zendnn/ggml-zendnn.cpp b/ggml/src/ggml-zendnn/ggml-zendnn.cpp index e6a9b51b7925..dc7d4d4f10c1 100644 --- a/ggml/src/ggml-zendnn/ggml-zendnn.cpp +++ b/ggml/src/ggml-zendnn/ggml-zendnn.cpp @@ -583,6 +583,8 @@ static struct ggml_backend_i ggml_backend_zendnn_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, + /* .prefetch_tensor_async = */ NULL, + /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_zendnn_guid(void) { diff --git a/include/llama.h b/include/llama.h index 38d0078048ad..f3fd37598244 100644 --- a/include/llama.h +++ b/include/llama.h @@ -399,6 +399,8 @@ extern "C" { // try to disable when n_seq_max > 1 for improved performance when the sequences do not share a large prefix // ref: https://github.com/ggml-org/llama.cpp/pull/14363 + bool prefetch_weights; // prefetch weight transfers to overlap CPU->GPU copies with compute + // [EXPERIMENTAL] // backend sampler chain configuration (make sure the caller keeps the sampler chains alive) // note: the samplers must be sampler chains (i.e. use llama_sampler_chain_init) diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 386c8ee49eb1..3f4a323f3804 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -267,8 +267,9 @@ llama_context::llama_context( cparams.n_outputs_max = params.n_outputs_max == 0 || llama_model_has_encoder(&model) ? cparams.n_batch : params.n_outputs_max; - cparams.op_offload = params.op_offload; - cparams.kv_unified = params.kv_unified; + cparams.op_offload = params.op_offload; + cparams.kv_unified = params.kv_unified; + cparams.prefetch_weights = params.prefetch_weights; // initialized later cparams.pipeline_parallel = false; @@ -599,6 +600,7 @@ void llama_context::sched_reserve() { gf_res_reserve.reset(new llm_graph_result(max_nodes)); sched.reset(ggml_backend_sched_new(backend_ptrs.data(), backend_buft.data(), backend_ptrs.size(), max_nodes, cparams.pipeline_parallel, cparams.op_offload)); + ggml_backend_sched_set_prefetch_weights(sched.get(), cparams.prefetch_weights); llama_memory_context_ptr mctx; if (memory) { @@ -688,6 +690,7 @@ void llama_context::sched_reserve() { LLAMA_LOG_WARN("%s: compute buffer allocation failed, retrying without pipeline parallelism\n", __func__); cparams.pipeline_parallel = false; sched.reset(ggml_backend_sched_new(backend_ptrs.data(), backend_buft.data(), backend_ptrs.size(), max_nodes, false, cparams.op_offload)); + ggml_backend_sched_set_prefetch_weights(sched.get(), cparams.prefetch_weights); gf = graph_reserve(n_tokens, n_seqs, n_outputs_pp, mctx.get()); } if (!gf) { @@ -3569,6 +3572,7 @@ llama_context_params llama_context_default_params() { /*.op_offload =*/ true, /*.swa_full =*/ true, /*.kv_unified =*/ false, + /*.prefetch_weights =*/ false, /*.sampler =*/ nullptr, /*.n_sampler =*/ 0, /*.ctx_other =*/ nullptr, diff --git a/src/llama-cparams.h b/src/llama-cparams.h index 8b6ab1ecf407..7654a196963e 100644 --- a/src/llama-cparams.h +++ b/src/llama-cparams.h @@ -52,6 +52,7 @@ struct llama_cparams { bool op_offload; bool kv_unified; bool pipeline_parallel; + bool prefetch_weights; std::vector embeddings_layer_inp; // [n_layer()] extract input embeddings for layer diff --git a/tools/llama-bench/llama-bench.cpp b/tools/llama-bench/llama-bench.cpp index e2e9160709c3..90497271ea12 100644 --- a/tools/llama-bench/llama-bench.cpp +++ b/tools/llama-bench/llama-bench.cpp @@ -349,6 +349,7 @@ struct cmd_params { std::vector> tensor_buft_overrides; std::vector embeddings; std::vector no_op_offload; + std::vector prefetch_weights; std::vector no_host; std::vector fit_params_target; std::vector fit_params_min_ctx; @@ -393,6 +394,7 @@ static const cmd_params cmd_params_defaults = { /* tensor_buft_overrides*/ { std::vector{ { nullptr, nullptr } } }, /* embeddings */ { false }, /* no_op_offload */ { false }, + /* prefetch_weights */ { false }, /* no_host */ { false }, /* fit_params_target */ { 0 }, /* fit_params_min_ctx */ { 0 }, @@ -906,6 +908,13 @@ static cmd_params parse_cmd_params(int argc, char ** argv) { } auto p = string_split(argv[i], split_delim); params.no_op_offload.insert(params.no_op_offload.end(), p.begin(), p.end()); + } else if (arg == "-pw" || arg == "--prefetch-weights") { + if (++i >= argc) { + invalid_param = true; + break; + } + auto p = string_split(argv[i], split_delim); + params.prefetch_weights.insert(params.prefetch_weights.end(), p.begin(), p.end()); } else if (arg == "--no-host") { if (++i >= argc) { invalid_param = true; @@ -1174,6 +1183,9 @@ static cmd_params parse_cmd_params(int argc, char ** argv) { if (params.no_op_offload.empty()) { params.no_op_offload = cmd_params_defaults.no_op_offload; } + if (params.prefetch_weights.empty()) { + params.prefetch_weights = cmd_params_defaults.prefetch_weights; + } if (params.no_host.empty()) { params.no_host = cmd_params_defaults.no_host; } @@ -1224,6 +1236,7 @@ struct cmd_params_instance { std::vector tensor_buft_overrides; bool embeddings; bool no_op_offload; + bool prefetch_weights; bool no_host; size_t fit_target; uint32_t fit_min_ctx; @@ -1300,6 +1313,7 @@ struct cmd_params_instance { cparams.flash_attn_type = flash_attn; cparams.embeddings = embeddings; cparams.op_offload = !no_op_offload; + cparams.prefetch_weights = prefetch_weights; cparams.swa_full = false; return cparams; @@ -1325,6 +1339,7 @@ static std::vector get_cmd_params_instances(const cmd_param for (const auto & noh : params.no_host) for (const auto & embd : params.embeddings) for (const auto & nopo : params.no_op_offload) + for (const auto & pw : params.prefetch_weights) for (const auto & nb : params.n_batch) for (const auto & nub : params.n_ubatch) for (const auto & tk : params.type_k) @@ -1365,6 +1380,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .tensor_buft_overrides = */ ot, /* .embeddings = */ embd, /* .no_op_offload = */ nopo, + /* .prefetch_weights = */ pw, /* .no_host = */ noh, /* .fit_target = */ fpt, /* .fit_min_ctx = */ fpc, @@ -1401,6 +1417,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .tensor_buft_overrides = */ ot, /* .embeddings = */ embd, /* .no_op_offload = */ nopo, + /* .prefetch_weights = */ pw, /* .no_host = */ noh, /* .fit_target = */ fpt, /* .fit_min_ctx = */ fpc, @@ -1437,6 +1454,7 @@ static std::vector get_cmd_params_instances(const cmd_param /* .tensor_buft_overrides = */ ot, /* .embeddings = */ embd, /* .no_op_offload = */ nopo, + /* .prefetch_weights = */ pw, /* .no_host = */ noh, /* .fit_target = */ fpt, /* .fit_min_ctx = */ fpc, @@ -1478,6 +1496,7 @@ struct test { std::vector tensor_buft_overrides; bool embeddings; bool no_op_offload; + bool prefetch_weights; bool no_host; size_t fit_target; uint32_t fit_min_ctx; @@ -1517,6 +1536,7 @@ struct test { tensor_buft_overrides = inst.tensor_buft_overrides; embeddings = inst.embeddings; no_op_offload = inst.no_op_offload; + prefetch_weights = inst.prefetch_weights; no_host = inst.no_host; fit_target = inst.fit_target; fit_min_ctx = inst.fit_min_ctx; @@ -1577,7 +1597,7 @@ struct test { "type_k", "type_v", "n_gpu_layers", "n_cpu_moe", "split_mode", "main_gpu", "no_kv_offload", "flash_attn", "devices", "tensor_split", "tensor_buft_overrides", "load_mode", "embeddings", - "no_op_offload", "no_host", "fit_target", "fit_min_ctx", + "no_op_offload", "prefetch_weights", "no_host", "fit_target", "fit_min_ctx", "n_prompt", "n_gen", "n_depth", "test_time", "avg_ns", "stddev_ns", "avg_ts", "stddev_ts" }; @@ -1590,7 +1610,7 @@ struct test { if (field == "build_number" || field == "n_batch" || field == "n_ubatch" || field == "n_threads" || field == "poll" || field == "model_size" || field == "model_n_params" || field == "n_gpu_layers" || field == "main_gpu" || field == "n_prompt" || field == "n_gen" || field == "n_depth" || field == "avg_ns" || - field == "stddev_ns" || field == "no_op_offload" || field == "n_cpu_moe" || + field == "stddev_ns" || field == "no_op_offload" || field == "prefetch_weights" || field == "n_cpu_moe" || field == "fit_target" || field == "fit_min_ctx" || field == "flash_attn") { return INT; } @@ -1673,6 +1693,7 @@ struct test { llama_load_mode_name(load_mode), std::to_string(embeddings), std::to_string(no_op_offload), + std::to_string(prefetch_weights), std::to_string(no_host), std::to_string(fit_target), std::to_string(fit_min_ctx), @@ -1864,6 +1885,9 @@ struct markdown_printer : public printer { if (field == "no_op_offload") { return 4; } + if (field == "prefetch_weights") { + return 2; + } if (field == "no_host") { return 4; } @@ -1901,6 +1925,9 @@ struct markdown_printer : public printer { if (field == "no_op_offload") { return "nopo"; } + if (field == "prefetch_weights") { + return "pw"; + } if (field == "no_host") { return "noh"; } @@ -1991,6 +2018,9 @@ struct markdown_printer : public printer { if (params.no_op_offload.size() > 1 || params.no_op_offload != cmd_params_defaults.no_op_offload) { fields.emplace_back("no_op_offload"); } + if (params.prefetch_weights.size() > 1 || params.prefetch_weights != cmd_params_defaults.prefetch_weights) { + fields.emplace_back("prefetch_weights"); + } if (params.no_host.size() > 1 || params.no_host != cmd_params_defaults.no_host) { fields.emplace_back("no_host"); } From 1c6df79f388607bc0972ee772f08bfc2df0f3857 Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Tue, 31 Mar 2026 23:45:42 +0800 Subject: [PATCH 2/4] simplify --- ggml/include/ggml-backend.h | 2 + ggml/src/ggml-backend-impl.h | 6 -- ggml/src/ggml-backend.cpp | 91 +++++++++++++++++++----- ggml/src/ggml-blas/ggml-blas.cpp | 2 - ggml/src/ggml-cann/ggml-cann.cpp | 2 - ggml/src/ggml-cpu/ggml-cpu.cpp | 2 - ggml/src/ggml-cuda/common.cuh | 3 - ggml/src/ggml-cuda/ggml-cuda.cu | 49 +------------ ggml/src/ggml-opencl/ggml-opencl.cpp | 2 - ggml/src/ggml-openvino/ggml-openvino.cpp | 2 - ggml/src/ggml-rpc/ggml-rpc.cpp | 2 - ggml/src/ggml-sycl/ggml-sycl.cpp | 2 - ggml/src/ggml-webgpu/ggml-webgpu.cpp | 2 - ggml/src/ggml-zdnn/ggml-zdnn.cpp | 2 - ggml/src/ggml-zendnn/ggml-zendnn.cpp | 2 - 15 files changed, 75 insertions(+), 96 deletions(-) diff --git a/ggml/include/ggml-backend.h b/ggml/include/ggml-backend.h index f83fdba69af5..99090fef7887 100644 --- a/ggml/include/ggml-backend.h +++ b/ggml/include/ggml-backend.h @@ -154,6 +154,8 @@ extern "C" { bool buffer_from_host_ptr; // event synchronization bool events; + // dedicated copy stream for compute/transfer overlap + bool copy_stream; }; // all the device properties diff --git a/ggml/src/ggml-backend-impl.h b/ggml/src/ggml-backend-impl.h index 5c7031d8189a..9c56ec30c5f1 100644 --- a/ggml/src/ggml-backend-impl.h +++ b/ggml/src/ggml-backend-impl.h @@ -137,12 +137,6 @@ extern "C" { // (optional) sort/optimize the nodes in the graph void (*graph_optimize) (ggml_backend_t backend, struct ggml_cgraph * cgraph); - - // (optional) prefetch tensor data on a separate copy stream for overlap with compute - void (*prefetch_tensor_async)(ggml_backend_t backend, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size); - - // (optional) wait for a pending prefetch to complete - void (*prefetch_event_wait)(ggml_backend_t backend); }; struct ggml_backend { diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index d0b36bc82f37..e13ed0686d42 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -819,6 +819,10 @@ struct ggml_backend_sched { bool op_offload; bool prefetch_weights; + // prefetch support: copy backends and events for compute/transfer overlap + ggml_backend_t copy_backends[GGML_SCHED_MAX_BACKENDS]; + ggml_backend_event_t copy_events[GGML_SCHED_MAX_BACKENDS]; + int debug; // used for debugging graph reallocations [GGML_SCHED_DEBUG_REALLOC] @@ -1627,6 +1631,9 @@ static bool ggml_backend_sched_alloc_splits(ggml_backend_sched_t sched) { // synchronize without ggml_backend_sched_synchronize to avoid changing cur_copy for (int i = 0; i < sched->n_backends; i++) { ggml_backend_synchronize(sched->backends[i]); + if (sched->copy_backends[i] != NULL) { + ggml_backend_synchronize(sched->copy_backends[i]); + } } ggml_gallocr_reserve_n(sched->galloc, &sched->graph, sched->node_backend_ids, sched->leaf_backend_ids); @@ -1658,27 +1665,35 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s next_weights_prefetched = false; if (sched->prefetch_weights) { - // wait for the previous prefetch to finish before issuing new ones - if (split_backend->iface.prefetch_event_wait) { - split_backend->iface.prefetch_event_wait(split_backend); - } + ggml_backend_t copy_backend = sched->copy_backends[split_backend_id]; + if (copy_backend != NULL) { + // compute stream waits for previous prefetch to complete + ggml_backend_event_wait(split_backend, sched->copy_events[split_backend_id]); + + // prefetch next split's weights on the copy stream (overlaps with current split's compute) + if (split_id + 1 < sched->n_splits) { + struct ggml_backend_sched_split * next = &splits[split_id + 1]; + // only prefetch when next split is on the same device (multi-GPU safety) + // and only for CPU->GPU transfers (host buffers) + if (next->backend_id == split_backend_id) { + // copy stream waits for previous compute to complete + if (sched->events[split_backend_id][sched->cur_copy] != NULL) { + ggml_backend_event_wait(copy_backend, sched->events[split_backend_id][sched->cur_copy]); + } - // prefetch next split's weights on the copy stream (overlaps with current split's compute) - if (split_backend->iface.prefetch_tensor_async && split_id + 1 < sched->n_splits) { - struct ggml_backend_sched_split * next = &splits[split_id + 1]; - // only prefetch when next split is on the same device (multi-GPU safety) - // and only for CPU->GPU transfers (host buffers) - if (next->backend_id == split_backend_id) { - for (int input_id = 0; input_id < next->n_inputs; input_id++) { - struct ggml_tensor * input = next->inputs[input_id]; - if (input->buffer != NULL && - ggml_backend_buffer_get_usage(input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && - ggml_backend_buffer_is_host(input->buffer)) { - struct ggml_tensor * input_cpy = tensor_copy(input, next->backend_id, sched->cur_copy); - split_backend->iface.prefetch_tensor_async( - split_backend, input_cpy, input->data, 0, ggml_nbytes(input)); - next_weights_prefetched = true; + for (int input_id = 0; input_id < next->n_inputs; input_id++) { + struct ggml_tensor * input = next->inputs[input_id]; + if (input->buffer != NULL && + ggml_backend_buffer_get_usage(input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS && + ggml_backend_buffer_is_host(input->buffer)) { + struct ggml_tensor * input_cpy = tensor_copy(input, next->backend_id, sched->cur_copy); + ggml_backend_tensor_set_async(copy_backend, input_cpy, input->data, 0, ggml_nbytes(input)); + next_weights_prefetched = true; + } } + + // signal that prefetch is done + ggml_backend_event_record(sched->copy_events[split_backend_id], copy_backend); } } } @@ -1942,6 +1957,10 @@ void ggml_backend_sched_free(ggml_backend_sched_t sched) { for (int c = 0; c < sched->n_copies; c++) { ggml_backend_event_free(sched->events[b][c]); } + ggml_backend_event_free(sched->copy_events[b]); + if (sched->copy_backends[b] != NULL) { + ggml_backend_free(sched->copy_backends[b]); + } } ggml_gallocr_free(sched->galloc); ggml_free(sched->ctx); @@ -2046,6 +2065,9 @@ void ggml_backend_sched_synchronize(ggml_backend_sched_t sched) { GGML_ASSERT(sched); for (int i = 0; i < sched->n_backends; i++) { ggml_backend_synchronize(sched->backends[i]); + if (sched->copy_backends[i] != NULL) { + ggml_backend_synchronize(sched->copy_backends[i]); + } } if (!sched->is_alloc) { // if the graph is not already allocated, always use copy 0 after a synchronization @@ -2064,6 +2086,37 @@ void ggml_backend_sched_set_eval_callback(ggml_backend_sched_t sched, ggml_backe void ggml_backend_sched_set_prefetch_weights(ggml_backend_sched_t sched, bool enabled) { GGML_ASSERT(sched); sched->prefetch_weights = enabled; + + if (enabled) { + // create copy backends and events for backends that support copy_stream + for (int b = 0; b < sched->n_backends; b++) { + if (sched->copy_backends[b] != NULL) { + continue; + } + ggml_backend_dev_t dev = ggml_backend_get_device(sched->backends[b]); + if (dev == NULL) { + continue; + } + struct ggml_backend_dev_props props; + ggml_backend_dev_get_props(dev, &props); + if (props.caps.copy_stream) { + sched->copy_backends[b] = ggml_backend_dev_init(dev, NULL); + sched->copy_events[b] = ggml_backend_event_new(dev); + } + } + } else { + // free copy backends and events + for (int b = 0; b < sched->n_backends; b++) { + if (sched->copy_events[b] != NULL) { + ggml_backend_event_free(sched->copy_events[b]); + sched->copy_events[b] = NULL; + } + if (sched->copy_backends[b] != NULL) { + ggml_backend_free(sched->copy_backends[b]); + sched->copy_backends[b] = NULL; + } + } + } } int ggml_backend_sched_get_n_splits(ggml_backend_sched_t sched) { diff --git a/ggml/src/ggml-blas/ggml-blas.cpp b/ggml/src/ggml-blas/ggml-blas.cpp index 997db1f50c39..9745fa29f5db 100644 --- a/ggml/src/ggml-blas/ggml-blas.cpp +++ b/ggml/src/ggml-blas/ggml-blas.cpp @@ -276,8 +276,6 @@ static struct ggml_backend_i blas_backend_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_blas_guid(void) { diff --git a/ggml/src/ggml-cann/ggml-cann.cpp b/ggml/src/ggml-cann/ggml-cann.cpp index 681628c879c7..5f51ea3bb3c8 100644 --- a/ggml/src/ggml-cann/ggml-cann.cpp +++ b/ggml/src/ggml-cann/ggml-cann.cpp @@ -2758,8 +2758,6 @@ static const ggml_backend_i ggml_backend_cann_interface = { /* .event_record = */ ggml_backend_cann_event_record, /* .event_wait = */ ggml_backend_cann_event_wait, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; /** diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index 8047898a13b0..16cc5116c545 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -207,8 +207,6 @@ static const struct ggml_backend_i ggml_backend_cpu_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_cpu_guid(void) { diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index e84d4cd235ef..d0a61051cfeb 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -1420,9 +1420,6 @@ struct ggml_backend_cuda_context { int device; std::string name; cudaEvent_t copy_event = nullptr; - cudaEvent_t prefetch_sync_event = nullptr; // signals copy stream that compute is done - cudaEvent_t prefetch_event = nullptr; // signals compute stream that prefetch is done - bool prefetch_pending = false; cudaStream_t streams[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = { { nullptr } }; cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES] = {nullptr}; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index c80699ab7305..00de3c47348a 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -711,12 +711,6 @@ ggml_backend_cuda_context::~ggml_backend_cuda_context() { if (copy_event != nullptr) { CUDA_CHECK(cudaEventDestroy(copy_event)); } - if (prefetch_sync_event != nullptr) { - CUDA_CHECK(cudaEventDestroy(prefetch_sync_event)); - } - if (prefetch_event != nullptr) { - CUDA_CHECK(cudaEventDestroy(prefetch_event)); - } for (int i = 0; i < GGML_CUDA_MAX_DEVICES; ++i) { for (int j = 0; j < GGML_CUDA_MAX_STREAMS; ++j) { if (streams[i][j] != nullptr) { @@ -4284,49 +4278,9 @@ static enum ggml_status ggml_backend_cuda_graph_compute(ggml_backend_t backend, ggml_cuda_graph_evaluate_and_capture(cuda_ctx, cgraph, use_cuda_graph, cuda_graph_update_required, graph_key); - // record event on compute stream so the copy stream can wait for compute to finish - // before starting prefetches that may reuse memory from this split's intermediates - if (cuda_ctx->prefetch_sync_event == nullptr) { - CUDA_CHECK(cudaEventCreateWithFlags(&cuda_ctx->prefetch_sync_event, cudaEventDisableTiming)); - } - CUDA_CHECK(cudaEventRecord(cuda_ctx->prefetch_sync_event, cuda_ctx->stream())); - return GGML_STATUS_SUCCESS; } -static void ggml_backend_cuda_prefetch_tensor_async(ggml_backend_t backend, ggml_tensor * tensor, const void * data, size_t offset, size_t size) { - ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context; - - ggml_cuda_set_device(cuda_ctx->device); - - // use stream 1 as dedicated copy stream for overlap with compute - cudaStream_t copy_stream = cuda_ctx->stream(cuda_ctx->device, 1); - - // wait for previous compute to finish before writing to memory that may alias - // with the previous split's intermediate tensors - if (cuda_ctx->prefetch_sync_event != nullptr) { - CUDA_CHECK(cudaStreamWaitEvent(copy_stream, cuda_ctx->prefetch_sync_event, 0)); - } - - CUDA_CHECK(cudaMemcpyAsync((char *)tensor->data + offset, data, size, cudaMemcpyHostToDevice, copy_stream)); - - // record prefetch event so the scheduler can wait for it - if (cuda_ctx->prefetch_event == nullptr) { - CUDA_CHECK(cudaEventCreateWithFlags(&cuda_ctx->prefetch_event, cudaEventDisableTiming)); - } - CUDA_CHECK(cudaEventRecord(cuda_ctx->prefetch_event, copy_stream)); - cuda_ctx->prefetch_pending = true; -} - -static void ggml_backend_cuda_prefetch_event_wait(ggml_backend_t backend) { - ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context; - if (cuda_ctx->prefetch_pending) { - // compute stream waits for copy stream (prefetched data is ready) - CUDA_CHECK(cudaStreamWaitEvent(cuda_ctx->stream(), cuda_ctx->prefetch_event, 0)); - cuda_ctx->prefetch_pending = false; - } -} - static void ggml_backend_cuda_event_record(ggml_backend_t backend, ggml_backend_event_t event) { ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *)backend->context; @@ -4613,8 +4567,6 @@ static const ggml_backend_i ggml_backend_cuda_interface = { /* .event_record = */ ggml_backend_cuda_event_record, /* .event_wait = */ ggml_backend_cuda_event_wait, /* .graph_optimize = */ ggml_backend_cuda_graph_optimize, - /* .prefetch_tensor_async = */ ggml_backend_cuda_prefetch_tensor_async, - /* .prefetch_event_wait = */ ggml_backend_cuda_prefetch_event_wait, }; static ggml_guid_t ggml_backend_cuda_guid() { @@ -4869,6 +4821,7 @@ static void ggml_backend_cuda_device_get_props(ggml_backend_dev_t dev, ggml_back /* .host_buffer = */ host_buffer, /* .buffer_from_host_ptr = */ false, /* .events = */ events, + /* .copy_stream = */ true, }; } diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index ef8decf92bd1..915c8e7b903f 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -7490,8 +7490,6 @@ static ggml_backend_i ggml_backend_opencl_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; ggml_backend_t ggml_backend_opencl_init(void) { diff --git a/ggml/src/ggml-openvino/ggml-openvino.cpp b/ggml/src/ggml-openvino/ggml-openvino.cpp index 046543bac3f3..0e7501fefe38 100644 --- a/ggml/src/ggml-openvino/ggml-openvino.cpp +++ b/ggml/src/ggml-openvino/ggml-openvino.cpp @@ -653,8 +653,6 @@ static const ggml_backend_i ggml_backend_openvino_interface = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; int ggml_backend_openvino_get_device_count() { diff --git a/ggml/src/ggml-rpc/ggml-rpc.cpp b/ggml/src/ggml-rpc/ggml-rpc.cpp index e44a1f4f5e79..17c53a5f049e 100644 --- a/ggml/src/ggml-rpc/ggml-rpc.cpp +++ b/ggml/src/ggml-rpc/ggml-rpc.cpp @@ -758,8 +758,6 @@ static ggml_backend_i ggml_backend_rpc_interface = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; ggml_backend_buffer_type_t ggml_backend_rpc_buffer_type(const char * endpoint, uint32_t device) { diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 16de8364b73e..131bff714f9c 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5570,8 +5570,6 @@ static ggml_backend_i ggml_backend_sycl_interface = { /* .event_record = */ ggml_backend_sycl_event_record, /* .event_wait = */ ggml_backend_sycl_event_wait, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_sycl_guid() { diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index eec6460a4e8a..c001cda7d116 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -3613,8 +3613,6 @@ static ggml_backend_i ggml_backend_webgpu_i = { /* .event_record = */ ggml_backend_webgpu_event_record, /* .event_wait = */ ggml_backend_webgpu_event_wait, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; /* End GGML Backend Interface */ diff --git a/ggml/src/ggml-zdnn/ggml-zdnn.cpp b/ggml/src/ggml-zdnn/ggml-zdnn.cpp index a550feba7ee8..3d67f64b639d 100644 --- a/ggml/src/ggml-zdnn/ggml-zdnn.cpp +++ b/ggml/src/ggml-zdnn/ggml-zdnn.cpp @@ -435,8 +435,6 @@ static ggml_backend_i ggml_backend_zdnn_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_zdnn_guid(void) { diff --git a/ggml/src/ggml-zendnn/ggml-zendnn.cpp b/ggml/src/ggml-zendnn/ggml-zendnn.cpp index dc7d4d4f10c1..e6a9b51b7925 100644 --- a/ggml/src/ggml-zendnn/ggml-zendnn.cpp +++ b/ggml/src/ggml-zendnn/ggml-zendnn.cpp @@ -583,8 +583,6 @@ static struct ggml_backend_i ggml_backend_zendnn_i = { /* .event_record = */ NULL, /* .event_wait = */ NULL, /* .graph_optimize = */ NULL, - /* .prefetch_tensor_async = */ NULL, - /* .prefetch_event_wait = */ NULL, }; static ggml_guid_t ggml_backend_zendnn_guid(void) { From 8f382090a768c0af8b6400b298d76850e5d8a980 Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Wed, 1 Apr 2026 00:20:44 +0800 Subject: [PATCH 3/4] add dedicated events --- ggml/src/ggml-backend.cpp | 26 ++++++++++++++++---------- 1 file changed, 16 insertions(+), 10 deletions(-) diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index e13ed0686d42..cec40c0730d0 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -821,7 +821,8 @@ struct ggml_backend_sched { // prefetch support: copy backends and events for compute/transfer overlap ggml_backend_t copy_backends[GGML_SCHED_MAX_BACKENDS]; - ggml_backend_event_t copy_events[GGML_SCHED_MAX_BACKENDS]; + ggml_backend_event_t copy_events[GGML_SCHED_MAX_BACKENDS]; // copy -> compute sync + ggml_backend_event_t compute_events[GGML_SCHED_MAX_BACKENDS]; // compute -> copy sync int debug; @@ -1677,9 +1678,7 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s // and only for CPU->GPU transfers (host buffers) if (next->backend_id == split_backend_id) { // copy stream waits for previous compute to complete - if (sched->events[split_backend_id][sched->cur_copy] != NULL) { - ggml_backend_event_wait(copy_backend, sched->events[split_backend_id][sched->cur_copy]); - } + ggml_backend_event_wait(copy_backend, sched->compute_events[split_backend_id]); for (int input_id = 0; input_id < next->n_inputs; input_id++) { struct ggml_tensor * input = next->inputs[input_id]; @@ -1868,6 +1867,11 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s } } + // record compute done for copy stream sync + if (sched->compute_events[split_backend_id] != NULL) { + ggml_backend_event_record(sched->compute_events[split_backend_id], split_backend); + } + // record the event of this copy if (split->n_inputs > 0) { if (sched->events[split_backend_id][sched->cur_copy] != NULL) { @@ -1958,6 +1962,7 @@ void ggml_backend_sched_free(ggml_backend_sched_t sched) { ggml_backend_event_free(sched->events[b][c]); } ggml_backend_event_free(sched->copy_events[b]); + ggml_backend_event_free(sched->compute_events[b]); if (sched->copy_backends[b] != NULL) { ggml_backend_free(sched->copy_backends[b]); } @@ -2100,17 +2105,18 @@ void ggml_backend_sched_set_prefetch_weights(ggml_backend_sched_t sched, bool en struct ggml_backend_dev_props props; ggml_backend_dev_get_props(dev, &props); if (props.caps.copy_stream) { - sched->copy_backends[b] = ggml_backend_dev_init(dev, NULL); - sched->copy_events[b] = ggml_backend_event_new(dev); + sched->copy_backends[b] = ggml_backend_dev_init(dev, NULL); + sched->copy_events[b] = ggml_backend_event_new(dev); + sched->compute_events[b] = ggml_backend_event_new(dev); } } } else { // free copy backends and events for (int b = 0; b < sched->n_backends; b++) { - if (sched->copy_events[b] != NULL) { - ggml_backend_event_free(sched->copy_events[b]); - sched->copy_events[b] = NULL; - } + ggml_backend_event_free(sched->copy_events[b]); + sched->copy_events[b] = NULL; + ggml_backend_event_free(sched->compute_events[b]); + sched->compute_events[b] = NULL; if (sched->copy_backends[b] != NULL) { ggml_backend_free(sched->copy_backends[b]); sched->copy_backends[b] = NULL; From 44bff68b9372f177ade029e599ccc531caca7e55 Mon Sep 17 00:00:00 2001 From: Aman Gupta Date: Wed, 1 Apr 2026 00:31:31 +0800 Subject: [PATCH 4/4] copy_stream false to other backends --- ggml/src/ggml-backend-meta.cpp | 1 + ggml/src/ggml-backend.cpp | 7 ++----- ggml/src/ggml-blas/ggml-blas.cpp | 1 + ggml/src/ggml-cann/ggml-cann.cpp | 1 + ggml/src/ggml-cpu/ggml-cpu.cpp | 1 + ggml/src/ggml-hexagon/ggml-hexagon.cpp | 1 + ggml/src/ggml-metal/ggml-metal.cpp | 1 + ggml/src/ggml-opencl/ggml-opencl.cpp | 1 + ggml/src/ggml-openvino/ggml-openvino.cpp | 1 + ggml/src/ggml-rpc/ggml-rpc.cpp | 1 + ggml/src/ggml-sycl/ggml-sycl.cpp | 1 + ggml/src/ggml-virtgpu/ggml-backend-device.cpp | 1 + ggml/src/ggml-vulkan/ggml-vulkan.cpp | 1 + ggml/src/ggml-webgpu/ggml-webgpu.cpp | 1 + ggml/src/ggml-zdnn/ggml-zdnn.cpp | 3 ++- ggml/src/ggml-zendnn/ggml-zendnn.cpp | 3 ++- 16 files changed, 19 insertions(+), 7 deletions(-) diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp index 499c3b767ce3..26f9e3d7fd8d 100644 --- a/ggml/src/ggml-backend-meta.cpp +++ b/ggml/src/ggml-backend-meta.cpp @@ -132,6 +132,7 @@ static void ggml_backend_meta_device_get_props(ggml_backend_dev_t dev, ggml_back /* .host_buffer = */ false, // Not implemented. /* .buffer_from_host_ptr = */ false, // Not implemented. /* .events = */ false, // Not implemented. + /* .copy_stream = */ false, // Not available }; for (ggml_backend_dev_t simple_dev : meta_dev_ctx->simple_devs) { ggml_backend_dev_props tmp_props; diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index cec40c0730d0..31a9e7b39e58 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -1671,11 +1671,10 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s // compute stream waits for previous prefetch to complete ggml_backend_event_wait(split_backend, sched->copy_events[split_backend_id]); - // prefetch next split's weights on the copy stream (overlaps with current split's compute) + // prefetch next split's weights on the copy stream if (split_id + 1 < sched->n_splits) { struct ggml_backend_sched_split * next = &splits[split_id + 1]; - // only prefetch when next split is on the same device (multi-GPU safety) - // and only for CPU->GPU transfers (host buffers) + // only prefetch when next split is on the same device if (next->backend_id == split_backend_id) { // copy stream waits for previous compute to complete ggml_backend_event_wait(copy_backend, sched->compute_events[split_backend_id]); @@ -2093,7 +2092,6 @@ void ggml_backend_sched_set_prefetch_weights(ggml_backend_sched_t sched, bool en sched->prefetch_weights = enabled; if (enabled) { - // create copy backends and events for backends that support copy_stream for (int b = 0; b < sched->n_backends; b++) { if (sched->copy_backends[b] != NULL) { continue; @@ -2111,7 +2109,6 @@ void ggml_backend_sched_set_prefetch_weights(ggml_backend_sched_t sched, bool en } } } else { - // free copy backends and events for (int b = 0; b < sched->n_backends; b++) { ggml_backend_event_free(sched->copy_events[b]); sched->copy_events[b] = NULL; diff --git a/ggml/src/ggml-blas/ggml-blas.cpp b/ggml/src/ggml-blas/ggml-blas.cpp index 9745fa29f5db..d3e91b892713 100644 --- a/ggml/src/ggml-blas/ggml-blas.cpp +++ b/ggml/src/ggml-blas/ggml-blas.cpp @@ -367,6 +367,7 @@ static void ggml_backend_blas_device_get_props(ggml_backend_dev_t dev, struct gg /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ true, /* .events = */ false, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-cann/ggml-cann.cpp b/ggml/src/ggml-cann/ggml-cann.cpp index 5f51ea3bb3c8..106268008070 100644 --- a/ggml/src/ggml-cann/ggml-cann.cpp +++ b/ggml/src/ggml-cann/ggml-cann.cpp @@ -2815,6 +2815,7 @@ static void ggml_backend_cann_device_get_props(ggml_backend_dev_t dev, ggml_back /* .host_buffer = */ host_buffer, /* .buffer_from_host_ptr = */ false, /* .events = */ true, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-cpu/ggml-cpu.cpp b/ggml/src/ggml-cpu/ggml-cpu.cpp index 16cc5116c545..859a4c8676db 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.cpp +++ b/ggml/src/ggml-cpu/ggml-cpu.cpp @@ -397,6 +397,7 @@ static void ggml_backend_cpu_device_get_props(ggml_backend_dev_t dev, struct ggm /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ true, /* .events = */ false, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp index bdb8af0820a3..09be7db8e037 100644 --- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp +++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp @@ -3930,6 +3930,7 @@ static void ggml_backend_hexagon_device_get_props(ggml_backend_dev_t dev, struct /* .host_buffer = */ (bool) opt_hostbuf, /* .buffer_from_host_ptr = */ false, /* .events = */ false, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-metal/ggml-metal.cpp b/ggml/src/ggml-metal/ggml-metal.cpp index a1003b3acff8..9d0ada10bfd9 100644 --- a/ggml/src/ggml-metal/ggml-metal.cpp +++ b/ggml/src/ggml-metal/ggml-metal.cpp @@ -681,6 +681,7 @@ static void ggml_backend_metal_device_get_props(ggml_backend_dev_t dev, ggml_bac /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ true, /* .events = */ true, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 915c8e7b903f..7516905cb232 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -10769,6 +10769,7 @@ static void ggml_backend_opencl_device_get_props(ggml_backend_dev_t dev, struct /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ false, /* .events = */ false, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-openvino/ggml-openvino.cpp b/ggml/src/ggml-openvino/ggml-openvino.cpp index 0e7501fefe38..4ce19d3de1a3 100644 --- a/ggml/src/ggml-openvino/ggml-openvino.cpp +++ b/ggml/src/ggml-openvino/ggml-openvino.cpp @@ -763,6 +763,7 @@ static void ggml_backend_openvino_device_get_props(ggml_backend_dev_t dev, ggml_ /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ false, /* .events = */ false, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-rpc/ggml-rpc.cpp b/ggml/src/ggml-rpc/ggml-rpc.cpp index 17c53a5f049e..340087de1aff 100644 --- a/ggml/src/ggml-rpc/ggml-rpc.cpp +++ b/ggml/src/ggml-rpc/ggml-rpc.cpp @@ -1881,6 +1881,7 @@ static void ggml_backend_rpc_device_get_props(ggml_backend_dev_t dev, struct ggm /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ false, /* .events = */ false, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 131bff714f9c..3b404923ffcc 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5639,6 +5639,7 @@ static void ggml_backend_sycl_device_get_props(ggml_backend_dev_t dev, ggml_back /* .host_buffer = */ host_buffer, /* .buffer_from_host_ptr = */ false, /* .events = */ events, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-virtgpu/ggml-backend-device.cpp b/ggml/src/ggml-virtgpu/ggml-backend-device.cpp index a978812cd908..8c896d1bc76e 100644 --- a/ggml/src/ggml-virtgpu/ggml-backend-device.cpp +++ b/ggml/src/ggml-virtgpu/ggml-backend-device.cpp @@ -70,6 +70,7 @@ static void ggml_backend_remoting_device_get_props(ggml_backend_dev_t dev, ggml_ props->caps.buffer_from_host_ptr = false; props->caps.async = false; props->caps.events = false; + props->caps.copy_stream = false; } ggml_backend_buffer_type_t ggml_backend_remoting_device_get_buffer_type(ggml_backend_dev_t dev) { diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 320ef861e88e..3b40fe738250 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -17858,6 +17858,7 @@ static void ggml_backend_vk_device_get_props(ggml_backend_dev_t dev, struct ggml /* .host_buffer = */ true, /* .buffer_from_host_ptr = */ false, /* .events = */ true, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index c001cda7d116..5ce9a4c59277 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -3954,6 +3954,7 @@ static void ggml_backend_webgpu_device_get_props(ggml_backend_dev_t dev, struct /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ false, /* .events = */ false, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-zdnn/ggml-zdnn.cpp b/ggml/src/ggml-zdnn/ggml-zdnn.cpp index 3d67f64b639d..60ecb22467c2 100644 --- a/ggml/src/ggml-zdnn/ggml-zdnn.cpp +++ b/ggml/src/ggml-zdnn/ggml-zdnn.cpp @@ -487,7 +487,8 @@ static void ggml_backend_zdnn_device_get_props(ggml_backend_dev_t dev, ggml_back /* .async = */ false, /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ false, - /* .events = */ false + /* .events = */ false, + /* .copy_stream = */ false, }; } diff --git a/ggml/src/ggml-zendnn/ggml-zendnn.cpp b/ggml/src/ggml-zendnn/ggml-zendnn.cpp index e6a9b51b7925..330d7d8df9d0 100644 --- a/ggml/src/ggml-zendnn/ggml-zendnn.cpp +++ b/ggml/src/ggml-zendnn/ggml-zendnn.cpp @@ -654,7 +654,8 @@ static void ggml_backend_zendnn_device_get_props(ggml_backend_dev_t dev, struct /* .async = */ false, /* .host_buffer = */ false, /* .buffer_from_host_ptr = */ true, - /* .events = */ false + /* .events = */ false, + /* .copy_stream = */ false, }; }