From c9613de45a86c36710727d52ccae2b3670e04875 Mon Sep 17 00:00:00 2001 From: derekja Date: Sat, 22 Aug 2026 02:10:17 +0000 Subject: [PATCH 1/3] cuda: export ggml_backend_cuda_clear_graph; qwen3_asr: release graphs before rebuild, exception-safe gallocr Follow-up to #276. Under concurrent /v1/audio/transcriptions load the qwen3_asr family leaks VRAM until allocation fails, after which every request 500s until process restart. Serial load is stable (#276 works); the concurrent-only signature is the tell: ggml's CUDA-graph cache (cuda_graphs) is keyed by cgraph->nodes[0], a host pointer into the graph's ggml_init arena. Every graph destructor already calls engine::core::release_backend_graph_resources, which looks up "ggml_backend_cuda_clear_graph" by proc address -- but the CUDA backend never exported that name, so eviction has been a silent no-op since it was added. Single-threaded servers get away with it: the arena is munmap'd and the next same-size ggml_init reuses the address, so the stale entry is overwritten in place. Concurrent requests (one detached thread per HTTP request) perturb the address space, each rebuild mints a fresh key, and the orphaned entries -- each holding a cudaGraph_t + cudaGraphExec_t -- accumulate until cudaMalloc fails. Measured on a 12 GB H100L vGPU slice, Qwen3-ASR-0.6B Q8, 8 concurrent streams of mixed 1-8 s utterances: VRAM 6.2 -> 8.9 -> 10.5 GB (ceiling) across identical repeated sweeps, then permanent 500s. Identical serial sweeps: byte-stable. nemotron_asr under the same concurrent load: byte-stable (its shapes do not churn), which localized the leak. Three changes: 1. ggml-cuda.cu: export ggml_backend_cuda_clear_graph through get_proc_address, making the existing destructor-side eviction calls effective (docs/build/HIP.md already claims this works; now it does). 2. qwen3_asr thinker/audio_encoder: reset the old graph before constructing its replacement. Assigning make_unique over a live unique_ptr holds both arenas at the rebuild peak, which doubles the transient footprint precisely when interleaved streams force a rebuild on nearly every request. 3. qwen3_asr: hold ggml_gallocr_t in a unique_ptr (the voxtral_realtime pattern) so the CapacityError/runtime_error throws in the graph constructors stop leaking the partially reserved arena -- previously every failed rebuild after an OOM deepened the OOM. Co-Authored-By: Claude Fable 5 --- external/ggml/src/ggml-cuda/ggml-cuda.cu | 3 ++ src/models/qwen3_asr/audio_encoder.cpp | 21 ++++++++---- src/models/qwen3_asr/thinker.cpp | 41 ++++++++++++++++-------- 3 files changed, 45 insertions(+), 20 deletions(-) diff --git a/external/ggml/src/ggml-cuda/ggml-cuda.cu b/external/ggml/src/ggml-cuda/ggml-cuda.cu index 44b3cd96..f6934cb1 100644 --- a/external/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/external/ggml/src/ggml-cuda/ggml-cuda.cu @@ -5854,6 +5854,9 @@ 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; + } return nullptr; } diff --git a/src/models/qwen3_asr/audio_encoder.cpp b/src/models/qwen3_asr/audio_encoder.cpp index 66870194..8a45de7b 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,10 @@ 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) + gallocr_.reset(ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend_))); + if (gallocr_ == nullptr || !ggml_gallocr_alloc_graph(gallocr_.get(), graph_)) { 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)); @@ -484,9 +495,6 @@ class Qwen3ASRAudioEncoderGraph { ~Qwen3ASRAudioEncoderGraph() { engine::core::release_backend_graph_resources(backend_, graph_); - if (gallocr_ != nullptr) { - ggml_gallocr_free(gallocr_); - } } bool matches(const Qwen3ASRAudioEncoderWeights & weights, int64_t frames, ggml_backend_t backend, int threads) const { @@ -557,7 +565,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 +596,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..82d1c6b6 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,14 @@ class PrefillGraph { for (auto * value : values_) { ggml_build_forward_expand(graph_, value); } - gallocr_ = ggml_gallocr_new(ggml_backend_get_default_buffer_type(runtime_->backend())); + // 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 + gallocr_.reset(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_)) { + !ggml_gallocr_reserve(gallocr_.get(), graph_) || + !ggml_gallocr_alloc_graph(gallocr_.get(), graph_)) { // 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. @@ -373,9 +386,6 @@ class PrefillGraph { ~PrefillGraph() { engine::core::release_backend_graph_resources(runtime_->backend(), graph_); - if (gallocr_ != nullptr) { - ggml_gallocr_free(gallocr_); - } } bool matches(const ThinkerWeightsRuntime & runtime, int64_t prompt_steps, int64_t audio_tokens) const { @@ -457,7 +467,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 +518,10 @@ 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())); + gallocr_.reset(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_)) { + !ggml_gallocr_reserve(gallocr_.get(), graph_) || + !ggml_gallocr_alloc_graph(gallocr_.get(), graph_)) { throw std::runtime_error("failed to allocate Qwen3 ASR thinker classification graph"); } position_ids_ = modules::qwen_position_ids(prompt_steps_); @@ -521,9 +531,6 @@ class PromptClassificationGraph { ~PromptClassificationGraph() { engine::core::release_backend_graph_resources(runtime_->backend(), graph_); - if (gallocr_ != nullptr) { - ggml_gallocr_free(gallocr_); - } } bool matches(const ThinkerWeightsRuntime & runtime, int64_t prompt_steps, int64_t audio_tokens) const { @@ -589,7 +596,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 { @@ -763,6 +770,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 +791,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 +830,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, From cb427c924a54535e6158bfb5c8a37edf198f9d81 Mon Sep 17 00:00:00 2001 From: derekja Date: Sat, 22 Aug 2026 02:18:59 +0000 Subject: [PATCH 2/3] cuda: trim idle pool memory and retry when graph allocation fails The legacy CUDA pool (no VMM) caches every buffer it ever allocated and only flushes when ITS OWN cudaMalloc fails; graph arenas allocate through ggml_backend_cuda_buffer_type_alloc_buffer, which does not flush anything. Under concurrent load with CUDA graphs enabled the pool ratchets up to the device ceiling, after which every graph (re)build fails permanently even though gigabytes of idle cached buffers are reclaimable -- the terminal all-500s state. (With GGML_CUDA_DISABLE_GRAPHS=1 the same workload is byte-stable, which is how the pool was isolated as the reservoir.) Adds ggml_cuda_pool::clear() (no-op by default, clear_pool() on the legacy pool), an exported ggml_backend_cuda_trim_pools(), a framework resolver engine::core::trim_backend_pools(), and a trim-and-retry-once on the allocation-failure paths of all four qwen3_asr graphs. A trim costs a device sync and only ever fires on a failure that was previously fatal. Co-Authored-By: Claude Fable 5 --- external/ggml/include/ggml-cuda.h | 1 + external/ggml/src/ggml-cuda/common.cuh | 3 +++ external/ggml/src/ggml-cuda/ggml-cuda.cu | 22 +++++++++++++++++ include/engine/framework/core/backend.h | 3 +++ src/framework/core/backend.cpp | 15 ++++++++++++ src/models/qwen3_asr/audio_encoder.cpp | 8 ++++-- src/models/qwen3_asr/thinker.cpp | 31 ++++++++++++++++++------ 7 files changed, 73 insertions(+), 10 deletions(-) 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 f6934cb1..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) { @@ -5857,6 +5876,9 @@ static void * ggml_backend_cuda_reg_get_proc_address(ggml_backend_reg_t reg, con 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..ba45dd26 100644 --- a/include/engine/framework/core/backend.h +++ b/include/engine/framework/core/backend.h @@ -44,6 +44,9 @@ 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); +// 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); 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); void validate_backend_graph_supported(ggml_backend_t backend, ggml_cgraph * graph, const char * label); diff --git a/src/framework/core/backend.cpp b/src/framework/core/backend.cpp index 4d70c247..867b41f1 100644 --- a/src/framework/core/backend.cpp +++ b/src/framework/core/backend.cpp @@ -305,6 +305,21 @@ static void cuda_clear_graph(ggml_backend_t backend, ggml_cgraph * graph) { if (fn != nullptr) fn(backend, 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); +} + void release_backend_graph_resources(ggml_backend_t backend, ggml_cgraph * graph) { if (is_cuda_backend_handle(backend) || is_hip_backend_handle(backend)) cuda_clear_graph(backend, graph); } diff --git a/src/models/qwen3_asr/audio_encoder.cpp b/src/models/qwen3_asr/audio_encoder.cpp index 8a45de7b..92aa9b50 100644 --- a/src/models/qwen3_asr/audio_encoder.cpp +++ b/src/models/qwen3_asr/audio_encoder.cpp @@ -484,8 +484,12 @@ class Qwen3ASRAudioEncoderGraph { ggml_build_forward_expand(graph_, output_); // unique_ptr so a throw below frees the partial reservation (a // throwing constructor runs no destructor) - gallocr_.reset(ggml_gallocr_new(ggml_backend_get_default_buffer_type(backend_))); - if (gallocr_ == nullptr || !ggml_gallocr_alloc_graph(gallocr_.get(), graph_)) { + 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)); diff --git a/src/models/qwen3_asr/thinker.cpp b/src/models/qwen3_asr/thinker.cpp index 82d1c6b6..6216c891 100644 --- a/src/models/qwen3_asr/thinker.cpp +++ b/src/models/qwen3_asr/thinker.cpp @@ -366,10 +366,17 @@ class PrefillGraph { // whatever the allocator already reserved; a throwing constructor // runs no destructor, and each failed rebuild used to leak its // partially reserved arena - gallocr_.reset(ggml_gallocr_new(ggml_backend_get_default_buffer_type(runtime_->backend()))); - if (gallocr_ == nullptr || - !ggml_gallocr_reserve(gallocr_.get(), graph_) || - !ggml_gallocr_alloc_graph(gallocr_.get(), 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_); + }; + // 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. @@ -518,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_.reset(ggml_gallocr_new(ggml_backend_get_default_buffer_type(runtime_->backend()))); - if (gallocr_ == nullptr || - !ggml_gallocr_reserve(gallocr_.get(), graph_) || - !ggml_gallocr_alloc_graph(gallocr_.get(), 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_); @@ -647,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"); } From 54b50add3b922f8ff3895eb2d8923686b24ced57 Mon Sep 17 00:00:00 2001 From: derekja Date: Sat, 22 Aug 2026 18:50:10 +0000 Subject: [PATCH 3/3] framework: make CUDA graph-cache eviction opt-in (default off), qwen3_asr opts in Per review: families that destroy and rebuild same-shape graphs between requests inherit a warm CUDA-graph cache from the historical no-op (the eviction lookup resolved nothing before the export), and evicting for everyone costs them a re-capture per rebuild (~9% warmed-request on MiniMax-H3). release_backend_graph_resources gains evict_cuda_graph_cache=false so every existing call site keeps the current behavior; qwen3_asr's four graph destructors pass true, keeping the concurrent-load leak fix where it was measured. Other families can opt in after audit. Co-Authored-By: Claude Fable 5 --- include/engine/framework/core/backend.h | 7 +++++-- src/framework/core/backend.cpp | 12 ++++++++++-- src/models/qwen3_asr/audio_encoder.cpp | 2 +- src/models/qwen3_asr/thinker.cpp | 6 +++--- 4 files changed, 19 insertions(+), 8 deletions(-) diff --git a/include/engine/framework/core/backend.h b/include/engine/framework/core/backend.h index ba45dd26..782667e8 100644 --- a/include/engine/framework/core/backend.h +++ b/include/engine/framework/core/backend.h @@ -47,8 +47,11 @@ bool requested_backend_uses_host_graph_plan(const BackendConfig & config); // 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); -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); +// 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 867b41f1..0297b51c 100644 --- a/src/framework/core/backend.cpp +++ b/src/framework/core/backend.cpp @@ -320,11 +320,19 @@ void trim_backend_pools(ggml_backend_t backend) { if (is_cuda_backend_handle(backend) || is_hip_backend_handle(backend)) cuda_trim_pools(backend); } -void release_backend_graph_resources(ggml_backend_t backend, ggml_cgraph * graph) { +// 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 92aa9b50..f7ca6957 100644 --- a/src/models/qwen3_asr/audio_encoder.cpp +++ b/src/models/qwen3_asr/audio_encoder.cpp @@ -498,7 +498,7 @@ class Qwen3ASRAudioEncoderGraph { } ~Qwen3ASRAudioEncoderGraph() { - engine::core::release_backend_graph_resources(backend_, graph_); + engine::core::release_backend_graph_resources(backend_, graph_, true); } bool matches(const Qwen3ASRAudioEncoderWeights & weights, int64_t frames, ggml_backend_t backend, int threads) const { diff --git a/src/models/qwen3_asr/thinker.cpp b/src/models/qwen3_asr/thinker.cpp index 6216c891..17a4fa31 100644 --- a/src/models/qwen3_asr/thinker.cpp +++ b/src/models/qwen3_asr/thinker.cpp @@ -392,7 +392,7 @@ class PrefillGraph { } ~PrefillGraph() { - engine::core::release_backend_graph_resources(runtime_->backend(), graph_); + engine::core::release_backend_graph_resources(runtime_->backend(), graph_, true); } bool matches(const ThinkerWeightsRuntime & runtime, int64_t prompt_steps, int64_t audio_tokens) const { @@ -541,7 +541,7 @@ class PromptClassificationGraph { } ~PromptClassificationGraph() { - engine::core::release_backend_graph_resources(runtime_->backend(), graph_); + engine::core::release_backend_graph_resources(runtime_->backend(), graph_, true); } bool matches(const ThinkerWeightsRuntime & runtime, int64_t prompt_steps, int64_t audio_tokens) const { @@ -671,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_); }