SYCL: a start refused for lack of VRAM ends with one plain line and a clean exit. A verifier whose init stopped before it made its own queues holds the process's default queue as cs_, sh_cs_ and copy_, and its destructor destroyed that queue; the next verifier's wait died in urQueueFinish (SIGSEGV on an 8 GB A750 with --batch 2 + MTP, where the expert cache gets 0 slots and the verify window refuses to start). The destructor now leaves the default queue alone. The refusal names what is taken and what to lower (--batch / parallel, --max-context, --vram-reserve-mib). verify_refusal_test: two refused verifiers, then the default queue must still run

sycl/CMakeLists.txt

@@ -428,6 +428,8 @@ if(STRATA_SYCL_PARITY)
   target_link_libraries(mmvq_bench PRIVATE strata_kernels)
   add_executable(dflash2_ring_test ${CMAKE_CURRENT_SOURCE_DIR}/src/core/dflash2_ring_test.cpp)   # DF2-SYCL: the ring log: rebuilt K/V rows are the original ones
   target_link_libraries(dflash2_ring_test PRIVATE strata_engine strata_prefill strata_spec)
+  add_executable(verify_refusal_test ${CMAKE_CURRENT_SOURCE_DIR}/src/core/verify_refusal_test.cpp)   # a refused start (no VRAM expert tier) ends cleanly: the verifiers leave the default queue alone
+  target_link_libraries(verify_refusal_test PRIVATE strata_engine strata_prefill strata_spec)
   add_executable(dflash2_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/dflash2_bench.cpp)   # DF2-SYCL: the DFlash2 Q4_0 drafter round on an Arc, kernel checks
   target_link_libraries(dflash2_bench PRIVATE strata_kernels)
   add_executable(dense_bench ${CMAKE_CURRENT_SOURCE_DIR}/src/kernels/dense_bench.cpp)   # the decode window's dense GEMVs, graph, hash

sycl/src/core/verify.cpp

@@ -440,10 +440,12 @@ Verifier::~Verifier() try {
     if (arena_b_) sycl::free(arena_b_, dpct::get_in_order_queue());
     if (h_commitb_) sycl::free(h_commitb_, dpct::get_in_order_queue());
     if (qcnt_) sycl::free(qcnt_, dpct::get_in_order_queue());
-    if (cs_ && cs_ != ext_stream_) dpct::get_current_device().destroy_queue(
+    // (a verifier whose init stopped before it made its own queues still holds the process's default queue in all three (cs_, sh_cs_, copy_): that one is never ours to destroy -
+    // destroying it took the next verifier's wait down with a SIGSEGV in urQueueFinish when a start was refused for lack of VRAM)
+    if (cs_ && cs_ != ext_stream_ && cs_ != &dpct::get_in_order_queue()) dpct::get_current_device().destroy_queue(
         cs_); // set_stream: the stage's stream, shared, not ours
-    if (sh_cs_) dpct::get_current_device().destroy_queue(sh_cs_);
-    if (copy_) { copy_->wait(); dpct::get_current_device().destroy_queue(copy_); }
+    if (sh_cs_ && sh_cs_ != cs_ && sh_cs_ != &dpct::get_in_order_queue()) dpct::get_current_device().destroy_queue(sh_cs_);
+    if (copy_) { copy_->wait(); if (copy_ != &dpct::get_in_order_queue()) dpct::get_current_device().destroy_queue(copy_); }   // (the default one when init never made the copy stream)
     if (commit_done_) dpct::destroy_event(commit_done_);
     if (ev_done_) dpct::destroy_event(ev_done_);
     if (ev_commit_) dpct::destroy_event(ev_commit_);

sycl/src/core/verify_refusal_test.cpp

@@ -0,0 +1,41 @@
+// sycl/src/core/verify_refusal_test.cpp - a start the verifier refuses (no VRAM expert tier: --batch slot sessions and the MTP layer left the expert cache no room on an 8 GB
+// card) ends with the refusal text and a clean exit.  Two verifiers whose init stops at the check never made queues of their own and hold the process's default queue; their
+// destructors used to destroy it (twice), and the second wait died in urQueueFinish (SIGSEGV on the Arc A750 at the end of `--serve --batch 2 --mtp`).  After both are gone the
+// default queue must still run work.
+//   verify_refusal_test        (needs one SYCL GPU)
+#define DPCT_PROFILING_ENABLED
+#include <sycl/sycl.hpp>
+#include <dpct/dpct.hpp>
+#include "strata/core/verify.hpp"
+
+#include <cstdio>
+#include <string>
+
+static int g_fail = 0;
+static void check(bool ok, const char* what) { std::printf("  %-100s %s\n", what, ok ? "ok" : "FAIL"); if (!ok) ++g_fail; }
+
+int main() {
+    using namespace strata::core;
+    sycl::queue* q = &dpct::get_in_order_queue();
+    std::printf("device: %s\n", q->get_device().get_info<sycl::info::device::name>().c_str());
+    {
+        WeightTable wt;
+        ModelGeometry g;
+        SessionState ss;
+        VerifyHits hits;   // no cache: d_res, cache_base null
+        Verifier a, b;
+        std::string ea, eb;
+        const bool oka = a.init(wt, g, ss, hits, nullptr, 4, ea);
+        const bool okb = b.init(wt, g, ss, hits, nullptr, 4, eb);
+        check(!oka && ea.rfind("verify: needs the profile-filled VRAM expert tier", 0) == 0, "first verifier: init refuses, with the expert-tier text");
+        check(!okb && eb.rfind("verify: needs the profile-filled VRAM expert tier", 0) == 0, "second verifier: init refuses, with the expert-tier text");
+    }   // both destructed here
+    float* p = sycl::malloc_device<float>(64, *q);
+    q->fill(p, 2.0f, 64).wait();
+    float h[64] = {};
+    q->memcpy(h, p, sizeof(h)).wait();
+    sycl::free(p, *q);
+    check(h[0] == 2.0f && h[63] == 2.0f, "the process's default queue still runs after both destructors");
+    std::printf("%s\n", g_fail ? "VERIFY-REFUSAL: FAIL" : "VERIFY-REFUSAL: ok");
+    return g_fail ? 1 : 0;
+}

sycl/src/program/generate.cpp

@@ -300,6 +300,16 @@ using Clock = std::chrono::steady_clock;
 
 // --pipeline-windows: every STRATA_PIPELINE_* test and tuning variable (THETA, FORCE_MISS, SWITCH, LOG, TRACE, DOOM_SKIP,
 // LOOKUP_PON, PRESTAGE, AGREE, SNAP_OVERLAP) is read only with STRATA_PIPELINE_DEBUG=1; without it the defaults apply.
+// The one plain line a start ends with when the expert cache got no VRAM at all (set where "no VRAM is left for the expert cache" is found, printed where the
+// verifier then refuses to start): what did not fit and what to lower.
+static std::string g_no_cache_hint;
+static void print_start_refusal(const char* who, const std::string& err) {
+    if (!g_no_cache_hint.empty() && err.rfind("verify: needs the profile-filled", 0) == 0)
+        std::fprintf(stderr, "%s: cannot start: %s\n", who, g_no_cache_hint.c_str());
+    else
+        std::fprintf(stderr, "%s: %s\n", who, err.c_str());
+}
+
 static const char* pipe_dbg_env(const char* name) {
     static const bool on = [] { const char* v = std::getenv("STRATA_PIPELINE_DEBUG"); return v != nullptr && v[0] != 0 && v[0] != '0'; }();
     return on ? std::getenv(name) : nullptr;
@@ -4750,6 +4760,11 @@ int main(int argc, char **argv) try {
                          (long long) min_slots, (long long) ((min_slots * blob) >> 20), (long long) short_mib,
                          (long long) session_mib, (long long) o.max_context, (long long) (mtp_bind >> 20),
                          reserve_tip.c_str());
+            g_no_cache_hint = "the GPU memory is full before the expert cache (it needs at least " + std::to_string(min_slots * blob >> 20) + " MiB; this card is about " +
+                              std::to_string(short_mib) + " MiB short). Taken: " + (o.batch > 1 ? std::to_string(o.batch) + " request slots " + std::to_string(o.batch * session_mib) + " MiB (--batch / parallel), " : std::string()) +
+                              "the context " + std::to_string(session_mib) + " MiB at " + std::to_string(o.max_context) + " tokens (--max-context), the draft head " +
+                              std::to_string(mtp_bind >> 20) + " MiB, the reserve " + std::to_string(at_reserve) + " MiB (--vram-reserve-mib). Lower " +
+                              (o.batch > 1 ? "the number of requests at once (--batch / parallel), " : "") + "--max-context or --vram-reserve-mib";
         }
     } else if (multi_gpu && o.expert_cache > 0) {
         // an explicit cache size leaves room for the prompt path's buffers and the reserve, or the first prompt
@@ -7575,7 +7590,7 @@ int main(int argc, char **argv) try {
                       batch_mtp ? strata::kernels::kVerifyMaxT : std::max(o.spec, o.batch), err) ||
             (use_mtp && !mtp.bind(last_st ? last_st->wt : wt, last_st ? &last_st->head : &native_head,
                                   pipe ? pl_mtp_R : ver.final_R_all(), err))) {
-            std::fprintf(stderr, "strata serve: %s\n", err.c_str());
+            print_start_refusal("strata serve", err);
             return 1;
         }
         if (pipe) mtp.set_source_R(ver.final_R_all());   // the serial decode's rows (the last stage's even verifier)
@@ -13005,7 +13020,7 @@ int main(int argc, char **argv) try {
         vh.n_slots = xcache.slots();
         if (use_df2) ver.set_taps(df2.tap_layers());
         if (!ver.init(wt, g, ss, vh, native_head.loaded() ? &native_head : nullptr, o.spec, err)) {
-            std::fprintf(stderr, "strata generate: %s\n", err.c_str());
+            print_start_refusal("strata generate", err);
             return 1;
         }
         const bool use_mtp = !o.mtp.empty();