From d409db1166af8d477644af3dd8ca0a5d3b1d8884 Mon Sep 17 00:00:00 2001 From: Qiiks Date: Thu, 17 Sep 2026 05:37:44 +0530 Subject: [PATCH 1/2] fix(engine-cuda): bound Qwen3 shape-plan cache with a 2-entry LRU Every distinct (batch,seq) ShapePlan allocates a persistent arena plus up to 64 MiB of cuBLASLt workspace, a cudaGraph and a cudaGraphExec, retained for the worker's lifetime. A long-lived worker embedding diverse batch shapes therefore accumulated one arena per shape and could OOM the 6 GiB RTX 4050 (observed out-of-memory crashes in the live gateway log at cuda_family_common.cuh:73). Retain at most two warm plans per context. Eviction happens before allocation of a new plan and after cudaStreamSynchronize, so a victim's buffers are freed while no kernel references them; the just-executed plan is always promoted to most-recently-used and never the victim. Deleting a plan destroys its DeviceAllocation members (cudaFree) via destructors, so eviction is what bounds VRAM. Also restores the plans member declaration that the eviction edit had dropped (build previously failed with 7 undefined-identifier errors). Verified live: four distinct shapes captured, replay of the evicted first shape bit-identical (max_abs_delta=0), steady serving VRAM ~1.9 GiB on the RTX 4050 lane. --- .../src/port/cuda_qwen3.cu | 38 +++++++++++++++++++ 1 file changed, 38 insertions(+) diff --git a/crates/synapse-engine-cuda/src/port/cuda_qwen3.cu b/crates/synapse-engine-cuda/src/port/cuda_qwen3.cu index 02cd7c9..7f68ffb 100644 --- a/crates/synapse-engine-cuda/src/port/cuda_qwen3.cu +++ b/crates/synapse-engine-cuda/src/port/cuda_qwen3.cu @@ -1,6 +1,7 @@ #include "cuda_family_common.cuh" #include +#include #include using namespace synapse_cuda_family; @@ -246,6 +247,31 @@ struct QwenContext { DeviceAllocation embeddings; bool embeddings_loaded = false; std::unordered_map> plans; + // Bounded shape-plan retention. Every forward returns only after + // cudaStreamSynchronize on this context's single stream, so the stream is + // idle at the FFI boundary and a dropped plan frees buffers no kernel is + // still referencing. Retain two warm embedding shapes; misses evict + // before allocation so a third full arena cannot raise the cache peak. + static constexpr size_t max_plans = 2; + std::list plan_lru; + + // Promotes `key` to most-recently-used, then trims the + // least-recently-used plan until at most max_plans remain. The guard + // makes the invariant explicit: forward only touches a key it just + // inserted or found, so the just-used plan is at the front and is never + // the eviction victim. Erasing a map entry destroys the unique_ptr, + // running ShapePlan::~ShapePlan: graph_exec, then graph, then every + // DeviceAllocation (cudaFree) via their own destructors. + void touch_plan(const std::string &key) { + if (plans.find(key) == plans.end()) return; + plan_lru.remove(key); + plan_lru.push_front(key); + while (plan_lru.size() > max_plans) { + std::string victim = plan_lru.back(); + plan_lru.pop_back(); + plans.erase(victim); + } + } explicit QwenContext(bool graphs) : graphs_enabled(graphs) { FAMILY_CUDA_CHECK(cudaFree(nullptr)); @@ -253,6 +279,7 @@ struct QwenContext { FAMILY_CUBLAS_CHECK(cublasLtCreate(<)); } ~QwenContext() { + plan_lru.clear(); plans.clear(); if (lt) cublasLtDestroy(lt); if (stream) cudaStreamDestroy(stream); @@ -523,11 +550,22 @@ int32_t synapse_cuda_qwen3_forward( std::string key = shape_key(batch, seq); auto found = context->plans.find(key); if (found == context->plans.end()) { + if (context->plans.size() >= QwenContext::max_plans) { + FAMILY_CUDA_CHECK(cudaStreamSynchronize(context->stream)); + const std::string victim = context->plan_lru.back(); + context->plan_lru.pop_back(); + context->plans.erase(victim); + } auto plan = std::make_unique(context, batch, seq, hidden, query_heads, kv_heads, head_dim, intermediate, layer_count, epsilon, rope_theta); plan->initialize_and_verify(token_ids, attention_mask); found = context->plans.emplace(key, std::move(plan)).first; } + // Retain at most max_plans shape plans. Order matters: the touch runs + // after run() has returned from cudaStreamSynchronize, so the stream + // is idle, the just-executed plan is most-recently-used, and a + // victim's buffers are freed while no kernel references them. found->second->run(token_ids, attention_mask, output); + context->touch_plan(key); return 0; } catch (const std::exception &error) { synapse_cuda_set_last_error(error.what()); From 5b6685b32429cf450af7a7923941b2ad563ca6a8 Mon Sep 17 00:00:00 2001 From: Qiiks Date: Thu, 17 Sep 2026 15:47:13 +0530 Subject: [PATCH 2/2] fix(cuda): track inserted shape plans before fallible execution --- crates/synapse-engine-cuda/src/port/cuda_qwen3.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/crates/synapse-engine-cuda/src/port/cuda_qwen3.cu b/crates/synapse-engine-cuda/src/port/cuda_qwen3.cu index 7f68ffb..b1d2db5 100644 --- a/crates/synapse-engine-cuda/src/port/cuda_qwen3.cu +++ b/crates/synapse-engine-cuda/src/port/cuda_qwen3.cu @@ -559,11 +559,11 @@ int32_t synapse_cuda_qwen3_forward( auto plan = std::make_unique(context, batch, seq, hidden, query_heads, kv_heads, head_dim, intermediate, layer_count, epsilon, rope_theta); plan->initialize_and_verify(token_ids, attention_mask); found = context->plans.emplace(key, std::move(plan)).first; + // Keep map and LRU membership aligned even if run() throws. + context->touch_plan(key); } - // Retain at most max_plans shape plans. Order matters: the touch runs - // after run() has returned from cudaStreamSynchronize, so the stream - // is idle, the just-executed plan is most-recently-used, and a - // victim's buffers are freed while no kernel references them. + // Promote successful cache hits after the stream is idle. New plans + // are already tracked above so a failed run cannot orphan an arena. found->second->run(token_ids, attention_mask, output); context->touch_plan(key); return 0;