cuda: make graph-cache eviction reachable; qwen3_asr: survive concurrent-load memory pressure

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

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<typename T>

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;
 }
 

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);

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);
 }
 

src/models/qwen3_asr/audio_encoder.cpp

@@ -18,6 +18,7 @@
 #include <cmath>
 #include <memory>
 #include <stdexcept>
+#include <type_traits>
 #include <utility>
 
 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<std::remove_pointer_t<ggml_gallocr_t>, 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<Qwen3ASRAudioEncoderGraph>(
             assets_,
             weights_,

src/models/qwen3_asr/thinker.cpp

@@ -27,6 +27,7 @@
 #include <memory>
 #include <optional>
 #include <stdexcept>
+#include <type_traits>
 #include <utility>
 #include <vector>
 
@@ -278,6 +279,14 @@ class ThinkerWeightsRuntime {
     std::shared_ptr<const ThinkerWeights> 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<ggml_tensor *> values_;
     std::vector<int32_t> position_ids_;
     ggml_cgraph * graph_ = nullptr;
-    ggml_gallocr_t gallocr_ = nullptr;
+    std::unique_ptr<std::remove_pointer_t<ggml_gallocr_t>, 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<int32_t> position_ids_;
     ggml_cgraph * graph_ = nullptr;
-    ggml_gallocr_t gallocr_ = nullptr;
+    std::unique_ptr<std::remove_pointer_t<ggml_gallocr_t>, 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<PrefillGraph>(
                 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<DecodeGraph>(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<PromptClassificationGraph>(
                 weights,
                 prompt_steps,