0.1.39: engine exits (code 1) with "parallel": 4 on UD-IQ4_XS, RTX 3090 24 GB, three concurrent 8K prompts (reproducible; IQ3_S unaffected)

Summary: with "parallel": 4, UD-IQ4_XS on a 24 GB RTX 3090 exits with code 1 (no error line) when three 8K-token prompts are read concurrently. Reproduced 2/2, at the same point both times. IQ3_S with the identical test and "parallel": 4 never failed (2 × 21/21 requests).

Environment

Strata v0.1.39 (6f32ec07), Linux build (Fedora 44, kernel 6.19.10), CUDA 13.3, driver 595.80
RTX 3090 24 GB, Ryzen 9 5950X (AVX2, no AVX-512), 128 GB DDR4 dual-channel
Model: Qwen3.8-Flash-Next UD-IQ4_XS (Unsloth, 3-part GGUF), native pack, MTP draft layer
Engine args: --expert-cache auto --prefill auto --spec 4 --spec-min-p 0.70 --mtp <rt> --max-context 32768 --kv int8 --resident-budget-gib 55 --pcie-frac 0.20 (the second run; the first used defaults --spec-min-p 0.5, no --pcie-frac, same result), "parallel": 4
Startup: --batch: 4 slot sessions on CUDA0 (0.56 GiB each); 15.89 GiB free, then 360 MiB of VRAM free with everything loaded
Host RAM at the time: ~72 GiB available, no swap used; no systemd OOM kill; unit memory peak 48.2 GB

Reproduction
OpenAI streaming chat completions, unique ~8,150-token prompts, max_tokens 512, reasoning on. Concurrency ladder 1 → 2 → 3, three waves per level:

Level 1 (3/3) and level 2 (6/6) passed: ~57 and ~28 tok/s per request.
Level 3: waves 1 and 2 passed (~20 tok/s each). On wave 3 the three slots were reading their prompts (slot 1 takes 8151 tokens, slot 2 takes 8152 tokens, prompt reads at ~1,860–1,885 tok/s). Each had generated 2–4 tokens when the engine exited.

Log at failure

strata serve: prompt 8152 tokens = 0 reused + 8152 read in 4324 ms (1885.1 tok/s), 1 generated in 25 ms (39.6 tok/s), drafts accepted 0 of 0, 1 checkpoints
strata serve: decode expert cache hit rate: 57.9% (271 hits / 468 lookups); 12 more read by the GPU over PCIe or from another GPU (2.5% of all 480 routed)
strata serve: resident RAM: 41.75 GiB of experts in RAM, 13981 exchanged with the VRAM tier, 27846 blob reads from the file
strata serve: expert tiers: GPU 271 hits this request; since the start RAM 1232526 blobs, files 27846 blobs 149546.1 MB read (the GGUF in place)
[strata] thinking: 4 of max 512 tokens, 0.3 tok/s, 19 s
[strata] the engine stopped unexpectedly (exit code 1). The engine exited (code 1). Its last log line: strata serve: expert tiers: GPU 271 hits this request; ...

All three in-flight requests failed (done: 1/2/4 tokens in 23 s (error, cancel=False)). The engine log has no traceback, CUDA error or out-of-memory message.

Possibly related: #511 / #535 (stalls with very little free VRAM). Here 360 MiB was free after loading. I haven't tried a larger --vram-reserve-mib yet; that would be my next test if useful.

Happy to provide full engine/server logs or run a debug build.

About the missing error line: a failed batch window prints its reason only as an ERR line on the engine's stdout, and with all three requests decoding in slots nobody reads it, so the server falls back to the last strata line in the log. #998 makes the server window show that ERR line. It only touches serve/server.py, so you can try it without rebuilding the engine; the line should say which step of the window failed.

We hit what looks like the same failure on a 48 GB card, and the hidden error line was visible in our server log. It is verify: batch instantiate: out of memory.

Setup

Strata v0.1.39 (6f32ec0), Linux, RTX PRO 5000 Blackwell 48 GB (PCIe 5.0 x16), Threadripper PRO 9955WX.
IQ3_S, "parallel": 8, --max-context 131072 --kv-resident 32768 --conversation-cache-mib 8192 --conversation-cache-slots 16 --pool-workers 5, --expert-cache auto, a custom expert profile.
Startup: --batch 8: the slot sessions take 5.82 GiB, expert cache 18124 slots, 34.37 GiB, then 379 MiB of VRAM free with everything loaded. The engine captured the 1- to 6-token windows at startup.

What happens
When 7 or 8 slots decode together, the server prints:

[strata] the engine reported an error: verify: batch instantiate: out of memory
[strata] done: 0 tokens in 4 s (0.0 tok/s) (error, cancel=False)
...
[strata] the engine stopped unexpectedly (exit code 1). The engine exited (code 1). Its last log line: strata serve: a batch window failed - ...

Every request in the batch fails, and the engine restarts.

We reproduced it three ways (streaming chat completions):

10 different ~1K-token prompts sent at once, 256 tokens each: 7 of 30 requests failed over 3 rounds.
8 different ~32K-token prompts sent at once, 512 tokens each: 7 of 8 were cut mid-answer in one round.
A multi-turn agent replay with 10 concurrent users: the engine failed twice in 10 minutes.

Workaround that worked: --vram-reserve-mib 1700 (1 GB more than the default 700). After that, no batch-window failure came back: 1K, 32K and 120K prompts with 8 and 10 requests at once, and the 10-user replay. A few 120K requests got HTTP 400 because our client made them longer than the context; that is a separate thing.

So the automatic expert cache seems to leave too little VRAM for the batch-window graphs that are first needed at runtime (7-8 rows here, maybe fewer rows with the larger UD-IQ4_XS buffers in this issue). Two possible fixes: reserve room for a --batch-row window when sizing the cache, or capture every window size up to --batch at startup.

serve: an engine that exits after an unread ERR line says why (#997 #890

OPUS-PORT-Allinea il fork ai 7 commit di upstream dopo v0.1.40

@Jackwwg83

@Jackwwg83 thanks, that ERR line was the key. Each set of slots that decodes together keeps its own graph pair (20-30 MiB each on an L40S), so with 8 slots the cache's last few hundred MiB run out. #1185 frees the least recently used ones when an instantiate hits OOM; on an L40S with ~440 MiB free it went from 8 engine restarts in 7 minutes to none.

serve: an engine that exits after an unread ERR line says why (Niko12…

Root cause, as far as we can tell, and a fix.

Every slot layout of a batch window (which slots are in it, how many rows) captures its own CUDA graph pair (exec_bm_ / commit_bm_ in src/core/verify.cpp). The pairs are never freed: set_batch_graph_limit(64) is set only with --batch-mtp ("0.1.39 keeps all"). As users start and end, new layouts keep coming (8 slots: up to 255), each pair holds VRAM, and the next cudaGraphInstantiate fails: verify: batch instantiate: out of memory, and the engine exits.

Measured on an RTX PRO 5000 48 GB (PCIe 5.0), IQ3_S, "parallel": 8, --vram-reserve-mib 1700, a 10-minute replay of a multi-user agent workload:

VRAM used grows about 1.1 GB in 2-5 minutes; the crash comes at 82-89 captured layouts (about 13 MiB a pair), with 5 users and with 8.
v0.1.40 as released: the engine stopped 1 time at 5 users, 2 at 8, 1 at 10 (8 failed requests), and at 12 users the requests then hung (0.1.39 serve: after an engine restart, waiting requests hang forever or fail with "list.remove(x): x not in list" #1012, fixed in 0.1.40.1).
A larger --vram-reserve-mib only delays it.

With 4 slots there are only 15 layouts, which is why it is rare; on a card with little VRAM left after the expert cache (like the 24 GB card in this issue) it is enough.

Our fix (verify.cpp, verify.hpp, generate.cpp):

Every stage's verifier keeps at most 64 layouts, the least recently used going first: the --batch-mtp limit, for every run. STRATA_BATCH_GRAPHS overrides it (0: no limit).
When an instantiation fails with cudaErrorMemoryAllocation, the least recently used half of the pairs is freed (never the layout being captured; cs_ is synchronized first) and the instantiation is tried again.

With the fix, the same replays at 5, 8, 10, 12 and 15 users ran 10 minutes each with 0 engine stops. With no limit (STRATA_BATCH_GRAPHS=0), 8 users for 5 minutes: the fallback freed graphs 2 times and the engine went on.

The patch (against our 0.1.40-based tree; the hunks touch only the batch graph cache)

diff --git a/include/strata/core/verify.hpp b/include/strata/core/verify.hpp
index a252fa6..ab24078 100644
--- a/include/strata/core/verify.hpp
+++ b/include/strata/core/verify.hpp
@@ -155,7 +155,8 @@ public:
     /// Commit an accepted prefix in each contiguous slot group of the last window. `keep` has
     /// one entry per slot (indexed by slot ID), each in 1..that slot's group length.
     bool commit_slot_prefixes(const int* keep, std::string& err);
-    /// --batch-mtp: slot rotation makes many row layouts, so bound the captured batch graph pairs (LRU). 0 = unbounded.
+    /// Bound the captured batch graph pairs (LRU; 0 = unbounded): every slot layout of a window has its own pair, and
+    /// slots that start and end make many layouts (#997: unbounded, ~13 MiB of VRAM each filled a 48 GB card).
     void set_batch_graph_limit(size_t n) { batch_graph_limit_ = n; }
 
     // ---- The stages of a layer split as a PIPELINE.  A batch window over the slot GROUP
@@ -273,7 +274,9 @@ private:
     int brow_[8] = {};                     ///< ... and row t is slot brow_[t]
     bool last_batch_ = false;              ///< the last run was a batch window (set_plan_slot: one group)
     std::map<std::vector<int>, cudaGraphExec_t> exec_bm_, commit_bm_;   ///< full row layout avoids slot-ID collisions
-    std::map<std::vector<int>, uint64_t> bm_used_;   ///< last use of each captured layout (LRU, only with a graph limit)
+    std::map<std::vector<int>, uint64_t> bm_used_;   ///< last use of each captured layout (LRU)
+    /// the VRAM ran short for a batch graph: frees the older half of the pairs (never `keep`'s) -> how many went
+    size_t drop_batch_graphs(const std::vector<int>& keep);
     uint64_t bm_tick_ = 0;
     size_t batch_graph_limit_ = 0;         ///< 0: keep every captured batch graph (0.1.39); N: LRU-evict beyond N layouts
     int last_rows_[8] = {};                ///< the slots of the last batch window's rows
diff --git a/src/core/verify.cpp b/src/core/verify.cpp
index cd9a642..fc1f490 100644
--- a/src/core/verify.cpp
+++ b/src/core/verify.cpp
@@ -2162,9 +2162,18 @@ bool Verifier::capture_batch(const int* rows, int S, int hbase, std::string& err
         err = !ok ? rerr : std::string("verify: end batch capture: ") + cudaGetErrorString(ce);
         return false;
     }
-    const cudaError_t ie = cudaGraphInstantiate(&ex, graph, 0);
+    cudaError_t ie = cudaGraphInstantiate(&ex, graph, 0);
+    if (ie == cudaErrorMemoryAllocation && drop_batch_graphs(bkey(rows, S, hbase)) > 0) {   // #997
+        (void) cudaGetLastError();
+        ex = nullptr;
+        ie = cudaGraphInstantiate(&ex, graph, 0);
+    }
     cudaGraphDestroy(graph);
-    if (ie != cudaSuccess) { err = std::string("verify: batch instantiate: ") + cudaGetErrorString(ie); return false; }
+    if (ie != cudaSuccess) {
+        ex = nullptr;
+        err = std::string("verify: batch instantiate: ") + cudaGetErrorString(ie);
+        return false;
+    }
     cudaGraphUpload(ex, cs_);
     cudaStreamSynchronize(cs_);
     std::string list;
@@ -2255,15 +2264,54 @@ bool Verifier::capture_commit_batch(const int* rows, int S, int hbase, std::stri
         if (graph) cudaGraphDestroy(graph);
         return false;
     }
-    if (ce != cudaSuccess || cudaGraphInstantiate(&cex, graph, 0) != cudaSuccess) {
+    cudaError_t ie = ce;
+    if (ce == cudaSuccess) {
+        ie = cudaGraphInstantiate(&cex, graph, 0);
+        if (ie == cudaErrorMemoryAllocation && drop_batch_graphs(bkey(rows, S, hbase)) > 0) {   // #997
+            (void) cudaGetLastError();
+            cex = nullptr;
+            ie = cudaGraphInstantiate(&cex, graph, 0);
+        }
+    }
+    if (ie != cudaSuccess) {
         if (graph) cudaGraphDestroy(graph);
-        err = std::string("verify: batch commit capture: ") + cudaGetErrorString(ce);
+        cex = nullptr;
+        err = std::string("verify: batch commit capture: ") + cudaGetErrorString(ie);
         return false;
     }
     cudaGraphDestroy(graph);
     return true;
 }
 
+// #997: every slot layout of a batch window has its graph pair, ~13 MiB of VRAM on a PRO 5000.  When an instantiation
+// finds no VRAM, the least recently used half goes (`keep`, the layout being captured, stays) and it is tried again.
+size_t Verifier::drop_batch_graphs(const std::vector<int>& keep) {
+    if (cudaStreamSynchronize(cs_) != cudaSuccess) return 0;
+    std::vector<std::pair<uint64_t, std::vector<int>>> old;
+    for (const auto& kv : exec_bm_)
+        if (kv.first != keep) {
+            const auto u = bm_used_.find(kv.first);
+            old.emplace_back(u == bm_used_.end() ? 0 : u->second, kv.first);
+        }
+    std::sort(old.begin(), old.end());
+    const size_t n = (old.size() + 1) / 2;
+    for (size_t i = 0; i < n; ++i) {
+        const std::vector<int>& k = old[i].second;
+        const auto e = exec_bm_.find(k);
+        if (e->second) cudaGraphExecDestroy(e->second);
+        exec_bm_.erase(e);
+        const auto c = commit_bm_.find(k);
+        if (c != commit_bm_.end()) {
+            if (c->second) cudaGraphExecDestroy(c->second);
+            commit_bm_.erase(c);
+        }
+        bm_used_.erase(k);
+    }
+    std::fprintf(stderr, "strata verify: no VRAM for a batch graph: freed %zu older slot layouts' graphs, %zu kept (#997)\n",
+                 n, exec_bm_.size());
+    return n;
+}
+
 bool Verifier::stage_batch(const int* rows, int S, int hbase, const int32_t* tokens, const int64_t* pos,
                            std::string& err) {
     using namespace strata::kernels;
@@ -2299,34 +2347,32 @@ bool Verifier::stage_batch(const int* rows, int S, int hbase, const int32_t* tok
             return false;
         }
     refresh_ar();
-    // --batch-mtp only (limit 0 = 0.1.39: no eviction): slot rotation creates new layouts; bound the captured graph
-    // pairs, evicting the least recently used layout.
-    if (batch_graph_limit_ > 0) {
-        const auto key = bkey(rows, S, hbase);
-        if (exec_bm_.find(key) == exec_bm_.end() && exec_bm_.size() >= batch_graph_limit_) {
-            if (cudaStreamSynchronize(cs_) != cudaSuccess) {
-                err = "verify: synchronizing before batch graph eviction failed";
-                return false;
-            }
-            auto old = exec_bm_.begin();
-            uint64_t oldest = UINT64_MAX;
-            for (auto it = exec_bm_.begin(); it != exec_bm_.end(); ++it) {
-                const auto u = bm_used_.find(it->first);
-                const uint64_t t = u == bm_used_.end() ? 0 : u->second;
-                if (t < oldest) { oldest = t; old = it; }
-            }
-            const auto old_key = old->first;
-            if (old->second) cudaGraphExecDestroy(old->second);
-            exec_bm_.erase(old);
-            bm_used_.erase(old_key);
-            auto commit_old = commit_bm_.find(old_key);
-            if (commit_old != commit_bm_.end()) {
-                if (commit_old->second) cudaGraphExecDestroy(commit_old->second);
-                commit_bm_.erase(commit_old);
-            }
+    // slots that start and end create new layouts: bound the captured graph pairs (limit 0 = 0.1.39: no eviction),
+    // evicting the least recently used layout.  The last use is kept either way (drop_batch_graphs).
+    const auto key = bkey(rows, S, hbase);
+    if (batch_graph_limit_ > 0 && exec_bm_.find(key) == exec_bm_.end() && exec_bm_.size() >= batch_graph_limit_) {
+        if (cudaStreamSynchronize(cs_) != cudaSuccess) {
+            err = "verify: synchronizing before batch graph eviction failed";
+            return false;
+        }
+        auto old = exec_bm_.begin();
+        uint64_t oldest = UINT64_MAX;
+        for (auto it = exec_bm_.begin(); it != exec_bm_.end(); ++it) {
+            const auto u = bm_used_.find(it->first);
+            const uint64_t t = u == bm_used_.end() ? 0 : u->second;
+            if (t < oldest) { oldest = t; old = it; }
+        }
+        const auto old_key = old->first;
+        if (old->second) cudaGraphExecDestroy(old->second);
+        exec_bm_.erase(old);
+        bm_used_.erase(old_key);
+        auto commit_old = commit_bm_.find(old_key);
+        if (commit_old != commit_bm_.end()) {
+            if (commit_old->second) cudaGraphExecDestroy(commit_old->second);
+            commit_bm_.erase(commit_old);
         }
-        bm_used_[key] = ++bm_tick_;
     }
+    bm_used_[key] = ++bm_tick_;
     if (!capture_batch(rows, S, hbase, err) || !capture_commit_batch(rows, S, hbase, err)) return false;
     const Clock::time_point t0 = Clock::now();
     const QsaShapes s = shapes_of(g);
diff --git a/src/program/generate.cpp b/src/program/generate.cpp
index 34b037d..3a186b8 100644
--- a/src/program/generate.cpp
+++ b/src/program/generate.cpp
@@ -6214,6 +6214,10 @@ int main(int argc, char** argv) {
                 for (auto& u : bslot_ss[k]) ptrs.push_back(u.get());
                 strata::core::Verifier& vk = k == 0 ? ver : stages[k - 1]->ver;
                 const strata::core::OnDevice on_k(k == 0 ? 0 : stages[k - 1]->dev);
+                // #997: a graph pair per slot layout, never freed, filled the VRAM.  STRATA_BATCH_GRAPHS=0: unbounded
+                // (the out-of-VRAM fallback then frees the older half: drop_batch_graphs)
+                const char* bg = std::getenv("STRATA_BATCH_GRAPHS");
+                vk.set_batch_graph_limit(bg != nullptr ? (size_t) std::strtoull(bg, nullptr, 10) : 64);
                 if (!vk.init_slots(ptrs, err)) {
                     std::fprintf(stderr, "strata serve: --batch, stage %zu: %s\n", k + 1, err.c_str());
                     return 1;

verify: a window graph that finds no VRAM frees older batch slot layo…

A window graph that finds no VRAM now frees the graphs of the least recently used batch slot layouts and tries again, instead of exiting (PR #1185 by ischencheng, in 0.1.40.2 with authorship kept); the start prints how many were freed. An exit with code 1 also names its reason since 0.1.40.1 (#998). Please retest parallel 4 to 8 on your L40S and the smaller card. If it still runs out, try a bigger --vram-reserve-mib and send the log.

serve: an engine that exits after an unread ERR line says why (Niko12…

verify: a window graph that finds no VRAM frees older batch slot layo…

Thanks for the report. To keep the tracker current we are closing issues older than three days: 0.1.44 is out with a long list of fixes, and many older reports are covered by it or by 0.1.43.1.

If this still happens on 0.1.44, please leave a comment here with your engine log (and the version line it prints), and we will reopen it. A new issue that links this one works too.

Fork 0.1.40-nvfp4.2: upstream 0.1.40.1 (its Niko1221#1058 gate replac…

feat(serve): engine-fault ledger with the row-recovery crossover remi…

Navigation Menu

Repository navigation

0.1.39: engine exits (code 1) with "parallel": 4 on UD-IQ4_XS, RTX 3090 24 GB, three concurrent 8K prompts (reproducible; IQ3_S unaffected) #997

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

0.1.39: engine exits (code 1) with "parallel": 4 on UD-IQ4_XS, RTX 3090 24 GB, three concurrent 8K prompts (reproducible; IQ3_S unaffected) #997

Description

Activity

ischencheng commented on Oct 5, 2026

Jackwwg83 commented on Oct 6, 2026

ischencheng commented on Oct 6, 2026

Jackwwg83 commented on Oct 7, 2026

Niko1221 commented on Oct 7, 2026

Niko1221 commented on Oct 13, 2026

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions