diff --git a/external/ggml/include/ggml-cuda.h b/external/ggml/include/ggml-cuda.h index 37c06f73..87fbaa15 100644 --- a/external/ggml/include/ggml-cuda.h +++ b/external/ggml/include/ggml-cuda.h @@ -23,6 +23,7 @@ extern "C" { GGML_BACKEND_API ggml_backend_t ggml_backend_cuda_init(int device); GGML_BACKEND_API bool ggml_backend_is_cuda(ggml_backend_t backend); +GGML_BACKEND_API void ggml_backend_cuda_trim_pools(ggml_backend_t backend); GGML_BACKEND_API void ggml_backend_cuda_clear_graph(ggml_backend_t backend, const struct ggml_cgraph * graph); // device buffer diff --git a/external/ggml/src/ggml-cuda/common.cuh b/external/ggml/src/ggml-cuda/common.cuh index cf39bd3f..79053a82 100644 --- a/external/ggml/src/ggml-cuda/common.cuh +++ b/external/ggml/src/ggml-cuda/common.cuh @@ -1131,6 +1131,9 @@ struct ggml_cuda_pool { virtual void * alloc(size_t size, size_t * actual_size) = 0; virtual void free(void * ptr, size_t size) = 0; + // Release cached device memory back to the driver. Buffers handed out by + // alloc() are unaffected; only idle cached capacity is dropped. + virtual void clear() {} }; template diff --git a/external/ggml/src/ggml-cuda/ggml-cuda.cu b/external/ggml/src/ggml-cuda/ggml-cuda.cu index 44b3cd96..ea67a93f 100644 --- a/external/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/external/ggml/src/ggml-cuda/ggml-cuda.cu @@ -380,6 +380,10 @@ struct ggml_cuda_pool_leg : public ggml_cuda_pool { GGML_ASSERT(pool_size == 0); } + void clear() override { + clear_pool(); + } + void clear_pool() { ggml_cuda_set_device(device); for (int i = 0; i < MAX_BUFFERS; ++i) { @@ -5034,6 +5038,21 @@ bool ggml_backend_is_cuda(ggml_backend_t backend) { return backend != NULL && ggml_guid_matches(backend->guid, ggml_backend_cuda_guid()); } +void ggml_backend_cuda_trim_pools(ggml_backend_t backend) { + if (!ggml_backend_is_cuda(backend)) { + return; + } + ggml_backend_cuda_context * cuda_ctx = (ggml_backend_cuda_context *) backend->context; + CUDA_CHECK(cudaDeviceSynchronize()); + for (int device = 0; device < GGML_CUDA_MAX_DEVICES; ++device) { + for (int stream = 0; stream < GGML_CUDA_MAX_STREAMS; ++stream) { + if (cuda_ctx->pools[device][stream] != nullptr) { + cuda_ctx->pools[device][stream]->clear(); + } + } + } +} + void ggml_backend_cuda_clear_graph(ggml_backend_t backend, const ggml_cgraph * graph) { #ifdef USE_CUDA_GRAPH if (!ggml_backend_is_cuda(backend) || graph == nullptr || graph->n_nodes <= 0) { @@ -5854,6 +5873,12 @@ static void * ggml_backend_cuda_reg_get_proc_address(ggml_backend_reg_t reg, con if (strcmp(name, "ggml_backend_get_features") == 0) { return (void *)ggml_backend_cuda_get_features; } + if (strcmp(name, "ggml_backend_cuda_clear_graph") == 0) { + return (void *)ggml_backend_cuda_clear_graph; + } + if (strcmp(name, "ggml_backend_cuda_trim_pools") == 0) { + return (void *)ggml_backend_cuda_trim_pools; + } return nullptr; } diff --git a/include/engine/framework/core/backend.h b/include/engine/framework/core/backend.h index 584af46f..782667e8 100644 --- a/include/engine/framework/core/backend.h +++ b/include/engine/framework/core/backend.h @@ -44,8 +44,14 @@ bool is_host_backend(ggml_backend_t backend); bool uses_host_graph_plan(BackendType type); bool uses_host_graph_plan(ggml_backend_t backend); bool requested_backend_uses_host_graph_plan(const BackendConfig & config); -void release_backend_graph_resources(ggml_backend_t backend, ggml_cgraph * graph); -void release_backend_graph_resources(BackendType backend_type, ggml_backend_t backend, ggml_cgraph * graph); +// Drop the CUDA/HIP context's cached (idle) pool memory back to the driver. +// No-op on other backends. For use on allocation-failure paths before a retry. +void trim_backend_pools(ggml_backend_t backend); +// evict_cuda_graph_cache=false (the default) is the historical no-op; +// true drops the backend's cached compiled-graph state (CUDA/HIP graph +// cache) for this cgraph at destruction — opt in per family. +void release_backend_graph_resources(ggml_backend_t backend, ggml_cgraph * graph, bool evict_cuda_graph_cache = false); +void release_backend_graph_resources(BackendType backend_type, ggml_backend_t backend, ggml_cgraph * graph, bool evict_cuda_graph_cache = false); void validate_backend_graph_supported(ggml_backend_t backend, ggml_cgraph * graph, const char * label); BackendMemorySnapshot query_backend_memory(ggml_backend_t backend, int device_hint); BackendMemorySnapshot query_backend_memory(const BackendConfig & config); diff --git a/src/framework/core/backend.cpp b/src/framework/core/backend.cpp index 4d70c247..0297b51c 100644 --- a/src/framework/core/backend.cpp +++ b/src/framework/core/backend.cpp @@ -305,11 +305,34 @@ static void cuda_clear_graph(ggml_backend_t backend, ggml_cgraph * graph) { if (fn != nullptr) fn(backend, graph); } -void release_backend_graph_resources(ggml_backend_t backend, ggml_cgraph * graph) { +static void cuda_trim_pools(ggml_backend_t backend) { + if (backend == nullptr) return; + ggml_backend_dev_t device = ggml_backend_get_device(backend); + if (device == nullptr) return; + auto fn = (void (*)(ggml_backend_t)) + ggml_backend_reg_get_proc_address( + ggml_backend_dev_backend_reg(device), + "ggml_backend_cuda_trim_pools"); + if (fn != nullptr) fn(backend); +} + +void trim_backend_pools(ggml_backend_t backend) { + if (is_cuda_backend_handle(backend) || is_hip_backend_handle(backend)) cuda_trim_pools(backend); +} + +// evict_cuda_graph_cache defaults to false, preserving historical behavior +// for existing call sites: before the CUDA backend exported +// ggml_backend_cuda_clear_graph the lookup resolved nothing, and families +// that rebuild same-shape graphs between requests inherit a warm CUDA-graph +// cache from that. Families that prefer bounded memory over the warm +// carry-over opt in with true. +void release_backend_graph_resources(ggml_backend_t backend, ggml_cgraph * graph, bool evict_cuda_graph_cache) { + if (!evict_cuda_graph_cache) return; // existing callsite/behavior unchanged if (is_cuda_backend_handle(backend) || is_hip_backend_handle(backend)) cuda_clear_graph(backend, graph); } -void release_backend_graph_resources(BackendType backend_type, ggml_backend_t backend, ggml_cgraph * graph) { +void release_backend_graph_resources(BackendType backend_type, ggml_backend_t backend, ggml_cgraph * graph, bool evict_cuda_graph_cache) { + if (!evict_cuda_graph_cache) return; // existing callsite/behavior unchanged if (backend_type == BackendType::Cuda || backend_type == BackendType::Hip) cuda_clear_graph(backend, graph); } diff --git a/src/models/qwen3_asr/audio_encoder.cpp b/src/models/qwen3_asr/audio_encoder.cpp index 66870194..f7ca6957 100644 --- a/src/models/qwen3_asr/audio_encoder.cpp +++ b/src/models/qwen3_asr/audio_encoder.cpp @@ -18,6 +18,7 @@ #include #include #include +#include #include namespace engine::models::qwen3_asr { @@ -337,6 +338,14 @@ modules::TransformerEncoderBlockWeights bind_layer(const AudioLayerWeights & wei return block; } +struct GgmlGallocrDeleter { + void operator()(ggml_gallocr_t alloc) const noexcept { + if (alloc != nullptr) { + ggml_gallocr_free(alloc); + } + } +}; + class Qwen3ASRAudioEncoderGraph { public: Qwen3ASRAudioEncoderGraph( @@ -473,8 +482,14 @@ class Qwen3ASRAudioEncoderGraph { ggml_set_output(output_); graph_ = ggml_new_graph_custom(ctx_.get(), 65536, false); ggml_build_forward_expand(graph_, output_); - gallocr_ = ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend_)); - if (gallocr_ == nullptr || !ggml_gallocr_alloc_graph(gallocr_, graph_)) { + // unique_ptr so a throw below frees the partial reservation (a + // throwing constructor runs no destructor) + const auto try_alloc = [&]() { + gallocr_.reset(ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend_))); + return gallocr_ != nullptr && ggml_gallocr_alloc_graph(gallocr_.get(), graph_); + }; + if (!try_alloc() && + (engine::core::trim_backend_pools(backend_), !try_alloc())) { throw std::runtime_error("failed to allocate Qwen3 ASR audio encoder graph"); } ggml_backend_tensor_set(attention_mask_, attention_mask_values.data(), 0, attention_mask_values.size() * sizeof(float)); @@ -483,10 +498,7 @@ class Qwen3ASRAudioEncoderGraph { } ~Qwen3ASRAudioEncoderGraph() { - engine::core::release_backend_graph_resources(backend_, graph_); - if (gallocr_ != nullptr) { - ggml_gallocr_free(gallocr_); - } + engine::core::release_backend_graph_resources(backend_, graph_, true); } bool matches(const Qwen3ASRAudioEncoderWeights & weights, int64_t frames, ggml_backend_t backend, int threads) const { @@ -557,7 +569,7 @@ class Qwen3ASRAudioEncoderGraph { ggml_tensor * attention_mask_ = nullptr; ggml_tensor * output_ = nullptr; ggml_cgraph * graph_ = nullptr; - ggml_gallocr_t gallocr_ = nullptr; + std::unique_ptr, GgmlGallocrDeleter> gallocr_; }; Qwen3ASRAudioEncoderRuntime::Qwen3ASRAudioEncoderRuntime( @@ -588,6 +600,7 @@ Qwen3ASRAudioEmbeddings Qwen3ASRAudioEncoderRuntime::encode(const Qwen3ASRAudioF } const int threads = std::max(1, execution_->config().threads); if (graph_ == nullptr || !graph_->matches(*weights_, features.frames, execution_->backend(), threads)) { + graph_.reset(); graph_ = std::make_unique( assets_, weights_, diff --git a/src/models/qwen3_asr/thinker.cpp b/src/models/qwen3_asr/thinker.cpp index 857b127d..17a4fa31 100644 --- a/src/models/qwen3_asr/thinker.cpp +++ b/src/models/qwen3_asr/thinker.cpp @@ -27,6 +27,7 @@ #include #include #include +#include #include #include @@ -278,6 +279,14 @@ class ThinkerWeightsRuntime { std::shared_ptr weights_; }; +struct GgmlGallocrDeleter { + void operator()(ggml_gallocr_t alloc) const noexcept { + if (alloc != nullptr) { + ggml_gallocr_free(alloc); + } + } +}; + class PrefillGraph { public: PrefillGraph( @@ -353,10 +362,21 @@ class PrefillGraph { for (auto * value : values_) { ggml_build_forward_expand(graph_, value); } - gallocr_ = ggml_gallocr_new(ggml_backend_get_default_buffer_type(runtime_->backend())); - if (gallocr_ == nullptr || - !ggml_gallocr_reserve(gallocr_, graph_) || - !ggml_gallocr_alloc_graph(gallocr_, graph_)) { + // unique_ptr so a throw below (or from later ctor statements) frees + // whatever the allocator already reserved; a throwing constructor + // runs no destructor, and each failed rebuild used to leak its + // partially reserved arena + const auto try_alloc = [&]() { + gallocr_.reset(ggml_gallocr_new(ggml_backend_get_default_buffer_type(runtime_->backend()))); + return gallocr_ != nullptr && + ggml_gallocr_reserve(gallocr_.get(), graph_) && + ggml_gallocr_alloc_graph(gallocr_.get(), graph_); + }; + // On failure the device may simply be full of idle cached pool + // buffers (the legacy pool never shrinks on its own); trim and retry + // once before declaring the size impossible + if (!try_alloc() && + (engine::core::trim_backend_pools(runtime_->backend()), !try_alloc())) { // Size, not a fault: the graph scales with prompt_steps_, which the // caller controls through the transcription prompt and the length of // the audio. Say which, and by how much, so the remedy is obvious. @@ -372,10 +392,7 @@ class PrefillGraph { } ~PrefillGraph() { - engine::core::release_backend_graph_resources(runtime_->backend(), graph_); - if (gallocr_ != nullptr) { - ggml_gallocr_free(gallocr_); - } + engine::core::release_backend_graph_resources(runtime_->backend(), graph_, true); } bool matches(const ThinkerWeightsRuntime & runtime, int64_t prompt_steps, int64_t audio_tokens) const { @@ -457,7 +474,7 @@ class PrefillGraph { std::vector values_; std::vector position_ids_; ggml_cgraph * graph_ = nullptr; - ggml_gallocr_t gallocr_ = nullptr; + std::unique_ptr, GgmlGallocrDeleter> gallocr_; }; class PromptClassificationGraph { @@ -508,10 +525,14 @@ class PromptClassificationGraph { ggml_set_output(token_ids_); graph_ = ggml_new_graph_custom(ctx_.get(), 65536, false); ggml_build_forward_expand(graph_, token_ids_); - gallocr_ = ggml_gallocr_new(ggml_backend_get_default_buffer_type(runtime_->backend())); - if (gallocr_ == nullptr || - !ggml_gallocr_reserve(gallocr_, graph_) || - !ggml_gallocr_alloc_graph(gallocr_, graph_)) { + const auto try_alloc = [&]() { + gallocr_.reset(ggml_gallocr_new(ggml_backend_get_default_buffer_type(runtime_->backend()))); + return gallocr_ != nullptr && + ggml_gallocr_reserve(gallocr_.get(), graph_) && + ggml_gallocr_alloc_graph(gallocr_.get(), graph_); + }; + if (!try_alloc() && + (engine::core::trim_backend_pools(runtime_->backend()), !try_alloc())) { throw std::runtime_error("failed to allocate Qwen3 ASR thinker classification graph"); } position_ids_ = modules::qwen_position_ids(prompt_steps_); @@ -520,10 +541,7 @@ class PromptClassificationGraph { } ~PromptClassificationGraph() { - engine::core::release_backend_graph_resources(runtime_->backend(), graph_); - if (gallocr_ != nullptr) { - ggml_gallocr_free(gallocr_); - } + engine::core::release_backend_graph_resources(runtime_->backend(), graph_, true); } bool matches(const ThinkerWeightsRuntime & runtime, int64_t prompt_steps, int64_t audio_tokens) const { @@ -589,7 +607,7 @@ class PromptClassificationGraph { ggml_tensor * token_ids_ = nullptr; std::vector position_ids_; ggml_cgraph * graph_ = nullptr; - ggml_gallocr_t gallocr_ = nullptr; + std::unique_ptr, GgmlGallocrDeleter> gallocr_; }; class DecodeGraph { @@ -640,6 +658,10 @@ class DecodeGraph { ggml_set_output(logits_); ggml_build_forward_expand(graph_, logits_); buffer_ = ggml_backend_alloc_ctx_tensors(ctx_.get(), runtime_->backend()); + if (buffer_ == nullptr) { + engine::core::trim_backend_pools(runtime_->backend()); + buffer_ = ggml_backend_alloc_ctx_tensors(ctx_.get(), runtime_->backend()); + } if (buffer_ == nullptr) { throw std::runtime_error("failed to allocate Qwen3 ASR thinker decode graph"); } @@ -649,7 +671,7 @@ class DecodeGraph { } ~DecodeGraph() { - engine::core::release_backend_graph_resources(runtime_->backend(), graph_); + engine::core::release_backend_graph_resources(runtime_->backend(), graph_, true); if (buffer_ != nullptr) { ggml_backend_buffer_free(buffer_); } @@ -763,6 +785,10 @@ struct Qwen3ASRThinkerRuntime::Impl { validate_prompt_audio(prompt, audio_embeddings); debug::timing_log_scalar("qwen3_asr.thinker.prompt_prepare_ms", engine::debug::elapsed_ms(timing_start, Clock::now())); if (prefill_graph == nullptr || !prefill_graph->matches(*weights, prompt_steps, audio_embeddings.tokens)) { + // drop the old graph (and its cuda_graphs cache entry) before + // allocating the replacement: assigning over a live unique_ptr + // holds both arenas at once at the rebuild peak + prefill_graph.reset(); prefill_graph = std::make_unique( weights, prompt_steps, @@ -780,6 +806,7 @@ struct Qwen3ASRThinkerRuntime::Impl { debug::timing_log_scalar("qwen3_asr.thinker.prefill_total_ms", engine::debug::elapsed_ms(timing_start, Clock::now())); const int64_t required_cache_steps = prompt_steps + options.max_new_tokens; if (decode_graph == nullptr || !decode_graph->can_run(*weights, required_cache_steps)) { + decode_graph.reset(); decode_graph = std::make_unique(weights, required_cache_steps, decode_graph_arena_bytes); } else { debug::timing_log_scalar("qwen3_asr.thinker.decode.graph.build_ms", 0.0); @@ -818,6 +845,7 @@ struct Qwen3ASRThinkerRuntime::Impl { debug::timing_log_scalar("qwen3_asr.thinker.classify_prompt_prepare_ms", engine::debug::elapsed_ms(timing_start, Clock::now())); if (classification_graph == nullptr || !classification_graph->matches(*weights, prompt_steps, audio_embeddings.tokens)) { + classification_graph.reset(); classification_graph = std::make_unique( weights, prompt_steps,