Activity

1. Kotali-2019

   changed the title [-][Engine] Layer split over two CUDA cards: stage 2 never gets prompt rows, stall at layer 1 of a small tail chunk[/-] [+][Engine] Layer split over two CUDA sm_86 cards: stage 2 never gets prompt rows, stall at layer 1 of a small tail chunk[/+] on Nov 13, 2026
2. JayOfTheKeyboard commented on Nov 15, 2026

   

   JayOfTheKeyboard

   on Nov 15, 2026

   Last edited by JayOfTheKeyboard

   More actions

   Confirming the root cause in this issue on a second sm_86 box, and with it three changes that make a 2-card split serve Edge0-120B Q3_K end to end. Patch against main @ 7b60076 is at the end (4 files, +100/−19). Happy to split it into PRs if you prefer.

   Machine

   2 × RTX 3070 8 GB (sm_86), PCIe 3.0 x8 each, 64 GB RAM, driver 620.70.06, CUDA 13.0, Linux. Model Edge0-120B Q3_K.

   Control

   Unpatched 7b60076, --layer-split 24 --expert-cache 0: the first request (141 tokens) fails with the same error as above:

   verify: layer 24 never rang (an illegal memory access was encountered)
   

   Changes

   1. Per-device latent D/U pair (latent_share.hpp, prefill.cpp, verify.cpp, generate.cpp). Essentially @Kotali-2019's patch from this issue, rebased onto main: latent_shared_d(dev) / latent_shared_u(dev), LatentTab::load_device(), and every pair loaded before any stage sizes its expert cache.

      Added: the three hand-off slots latent_xl_current(), latent_xlseq_current() and latent_act0_current() are now thread_local. Each stage's verifier runs on its own thread, and with process-wide slots one stage's pool call can read the other stage's pointers. latent_act0 is device memory on the other GPU.

   2. Each card holds only its own layers' D/U under an explicit --layer-split (LatentTab::load / load_device). The base pointer is offset by the stage's first layer, so readers keep indexing with the global layer number. This applies only when multi_gpu is set, so --split-device 0 still loads the full pair. It saves 0.66 GiB per card on Edge0-120B, which an 8 GB card needs.

   3. Layer-major prefill refused for latent models (Prefill::lm_prepare). Layer-major is on by default with 2+ stages, but run_layer_major / lm_run have no latent D/U step. Once a card had enough free VRAM for a super-chunk, the request produced non-finite logits (All-logits-NaN degeneration (one repeated token forever) still fires on 0.1.39b / current main — first poison traced to the MoE output rows: garbage bits, varying layer Niko1221/Strata#879) and returned HTTP 503. Reproduced. With this check the request falls back to the chunk-major path.

   Result

   Working config on 2 × 8 GB:

   --layer-split 24 --expert-cache 0 --prefill 512 --spec 4 --vram-reserve-mib 1000 --vram-reserve-later-mib 800 --max-context 262144 --kv int8
   STRATA_STAGE_TRIM=1 STRATA_EMBED_VRAM=0   (no --mtp: the draft layer does not fit on CUDA1)
   

   • Correct answers (checked by content) at prompt lengths from 24 to 246,191 tokens. No OOM at 262144 context. Minimum free VRAM was 813 / 453 MiB.
   • Decode 33 tok/s; after a 95K prompt, 33 tok/s. Prefill 281 tok/s at 94.7K and 271 tok/s at 246K, bound by PCIe 3.0 x8 (RX averaged 5.7–6.2 GB/s on both cards).
   • No stall of the kind in section 2 of this issue at any length. lat_dense_d_tail_test passes on both cards.

   Not tested

   • I could not make a single 8 GB card serve this model, so there was no single-GPU greedy reference. I checked correctness by reading the answers, not by token-for-token comparison.
   • --mtp on a split. The draft layer did not fit, so that path is untouched.
   • --layer-split auto and 3+ stages. The trimming in change 2 is skipped for auto.
   • Only sm_86, CUDA, Linux. Not built for Metal or HIP. The Metal branches are no-ops.
   • I can't say whether the thread_local change is what removes the tail stall in section 2. I never ran this patch without it on a split. @Kotali-2019, it would be worth trying on your CMP 90HX box.

   Related: cuBLAS status 14 on a nearly full split stage

   Separate from this issue, but found on the same setup. Stage sizing can leave a card with a few MiB free. Then the first prompt's BF16 cublasGemmEx (end of Gemm::bf16) fails with cuBLAS status 14 after the engine has reported READY. serve/server.py then restarts the engine into the same failure every 15 s.

   The cause is that cuBLAS loads its kernels on first use. warm_gemm_libraries() would take that memory at boot, but it is compiled for HIP only. Running it on CUDA as well, at the end of Prefill::init, moves the load into boot. There it comes out of the stage's reserve and shows in the "VRAM free with everything loaded" check. It costs about 35 MiB and 130 ms per card.

   Tested by filling CUDA0 to 2 MiB free after boot. The unfixed build fails with status 14. With the warm-up the GEMM passes, and the request then stops at verify: instantiate: out of memory, a clean 400, which I haven't touched. Decode and prefill on the working config were unchanged (within 4%). The patch also prints a hint in ck() when cuBLAS returns 1, 3 or 14 with under 256 MiB free. I can open it as a separate PR if it's wanted.

   Patch: layer-split fixes 1-3 (vs 7b60076)

   diff --git a/edge0-strata/include/strata/core/latent_share.hpp b/edge0-strata/include/strata/core/latent_share.hpp
   index fa8b7d7..dc68419 100644
   --- a/edge0-strata/include/strata/core/latent_share.hpp
   +++ b/edge0-strata/include/strata/core/latent_share.hpp
   @@ -7,12 +7,21 @@
    namespace strata {
    namespace core {
    
   -inline uint8_t*& latent_shared_d() { static uint8_t* p = nullptr; return p; }
   -inline uint8_t*& latent_shared_u() { static uint8_t* p = nullptr; return p; }
   +// One pair per DEVICE.  A layer split runs a stage - and that stage's verifier - on another GPU, and a
   +// pointer allocated on the primary is not valid there (the stage's first layer faults on it).  Every
   +// holder on one card shares that card's copy, so a split pays one 1.85 GiB (Q4) pair per card, not one
   +// per prefill AND one per verifier.
   +inline uint8_t*& latent_shared_d(int dev) { static uint8_t* p[64] = {}; return p[(unsigned) dev & 63u]; }
   +inline uint8_t*& latent_shared_u(int dev) { static uint8_t* p[64] = {}; return p[(unsigned) dev & 63u]; }
   +inline uint8_t*& latent_shared_d() { return latent_shared_d(0); }
   +inline uint8_t*& latent_shared_u() { return latent_shared_u(0); }
    
   -inline const uint8_t*& latent_xl_current() { static const uint8_t* p = nullptr; return p; }
   -inline const uint32_t*& latent_xlseq_current() { static const uint32_t* p = nullptr; return p; }
   -inline const uint8_t*& latent_act0_current() { static const uint8_t* p = nullptr; return p; }
   +// The window's D rows, handed from a verifier to the expert pool call it makes right after (same thread, set then
   +// read).  thread_local: a layer split runs one verifier per stage on its own thread, and a process-wide slot let one
   +// stage's pool call read the OTHER stage's pointers - latent_act0 is device memory of the other GPU.
   +inline const uint8_t*& latent_xl_current() { thread_local const uint8_t* p = nullptr; return p; }
   +inline const uint32_t*& latent_xlseq_current() { thread_local const uint32_t* p = nullptr; return p; }
   +inline const uint8_t*& latent_act0_current() { thread_local const uint8_t* p = nullptr; return p; }
    
    }  // namespace core
    }  // namespace strata
   diff --git a/edge0-strata/src/core/verify.cpp b/edge0-strata/src/core/verify.cpp
   index a919e9d..7cf3cbf 100644
   --- a/edge0-strata/src/core/verify.cpp
   +++ b/edge0-strata/src/core/verify.cpp
   @@ -1065,9 +1065,9 @@ bool Verifier::init(const WeightTable& wt, const ModelGeometry& g, SessionState&
                lat_yl_ = b.take<float>(T * K * LELL);
                // every layer's D and U, resident: the graph covers all layers and touches them per layer
                // (the shared pair from set_shared_latent: no own copy, read-only for every verifier)
   -            if (lat_d_shared_ == nullptr) {   // the session-level pair (LatentTab::load) when nobody passed it in
   -                lat_d_shared_ = strata::core::latent_shared_d();
   -                lat_u_shared_ = strata::core::latent_shared_u();
   +            if (lat_d_shared_ == nullptr) {   // this DEVICE's pair (LatentTab::load_device) when nobody passed one in
   +                lat_d_shared_ = strata::core::latent_shared_d(device_);
   +                lat_u_shared_ = strata::core::latent_shared_u(device_);
                }
                if (lat_d_shared_ == nullptr) {
                    lat_d_ = b.take<uint8_t>((uint64_t) g.n_layers * lat_dense_.d_bytes());
   diff --git a/edge0-strata/src/prefill/prefill.cpp b/edge0-strata/src/prefill/prefill.cpp
   index 41f5cef..5b4845c 100644
   --- a/edge0-strata/src/prefill/prefill.cpp
   +++ b/edge0-strata/src/prefill/prefill.cpp
   @@ -1728,9 +1728,9 @@ bool Prefill::init(const core::WeightTable& wt, const core::ModelGeometry& g, co
            }
            m.owned.push_back(m.lat_labels_dev);
        }
   -    if (m.latent && m.lat_d_dev == nullptr) {   // a session-shared pair when present (latent_share.hpp), else owned below
   -        m.lat_d_dev = strata::core::latent_shared_d();
   -        m.lat_u_dev = strata::core::latent_shared_u();
   +    if (m.latent && m.lat_d_dev == nullptr) {   // this DEVICE's shared pair when present (latent_share.hpp), else owned below
   +        m.lat_d_dev = strata::core::latent_shared_d(m.device);
   +        m.lat_u_dev = strata::core::latent_shared_u(m.device);
        }
        if (m.latent && m.lat_d_dev == nullptr) {   // no shared pair: one owned copy, uploaded once per session
            const uint64_t db = (uint64_t) g.n_layers * m.lat_d.d_bytes();
   @@ -7208,6 +7208,9 @@ int64_t Prefill::lm_prepare(int64_t n, std::string& why, int64_t cap) {
            }
        }
        if (m.fused_bufs) { why = "the fused expert layout"; return 0; }
   +    // latent MoE (Edge0): the experts here are latent-width and need the D/U projections around them; this walk gathers
   +    // hidden-width rows for them, which ends in non-finite logits (#879) once a card has room for a super-chunk
   +    if (m.latent) { why = "latent MoE (no D/U step in the layer-major walk)"; return 0; }
        if (!m.GU || !m.Dm || (n_mmq > 0 && !m.ids_identity)) { why = "the MoE buffers are missing"; return 0; }
        // (an explicit STRATA_PREFILL_LM_TOKENS may be one chunk: the A/B against the default path, experts aside)
        const int64_t min_s = (lm_force_ || lm_force_env()) ? 1 : (lm_tokens_env() > 0 ? m.T : 2 * m.T);   // (lm_force_: run_multi - several prompts, however short)
   diff --git a/edge0-strata/src/program/generate.cpp b/edge0-strata/src/program/generate.cpp
   index 97b0a7b..0c8c5a0 100644
   --- a/edge0-strata/src/program/generate.cpp
   +++ b/edge0-strata/src/program/generate.cpp
   @@ -232,8 +232,11 @@ struct LatentTab {
        int64_t lat_H = 0, lat_ELL = 0, lat_I = 0;
        int lat_G = 8, lat_E = 0;
        bool latent = false;
   +    int alloc_dev = 0;   // the device holding `d`/`u` (the pair is per device; see latent_share.hpp)
   +    int64_t keep_lo = 0, keep_hi = -1;   // layers the primary's device copy holds under an explicit layer split (-1: all)
        bool load(const strata::core::ModelGeometry& g, std::string& err) {
   -        if (d != nullptr) { strata::core::latent_shared_d() = d; strata::core::latent_shared_u() = u; return true; }
   +        // the pair below was allocated on alloc_dev, whatever device is current by the time this is called again
   +        if (d != nullptr) { strata::core::latent_shared_d(alloc_dev) = d; strata::core::latent_shared_u(alloc_dev) = u; return true; }
            const auto& lay0 = strata::kernels::cpu::expert_layout();
            if (!(lay0.native && !lay0.fmt.empty() && lay0.fmt[0].n_embd != (int64_t) g.n_embd)) return true;
    #if defined(STRATA_GPU_METAL)
   @@ -265,22 +268,73 @@ struct LatentTab {
                std::fclose(lf);
                if (got != labels.size()) { err = "latent dense: " + lp + " is too small"; return false; }
            }
   -        const uint64_t db = (uint64_t) g.n_layers * dense.d_bytes(), ub = (uint64_t) g.n_layers * dense.u_bytes();
   -        if (cudaMalloc((void**) &d, db) != cudaSuccess || cudaMalloc((void**) &u, ub) != cudaSuccess) {
   +        int dev_now = 0;
   +        cudaGetDevice(&dev_now);
   +        alloc_dev = dev_now;
   +        // Layer split: this device runs only [keep_lo, keep_hi) (set by main()), so hold just those layers' D/U and
   +        // offset the base by keep_lo layers: every reader indexes base + layer * bytes with the GLOBAL layer number.
   +        // An 8 GB card cannot spare the other half (0.66 GiB of the 1.32 GiB pair).
   +        const int64_t lo = keep_hi > keep_lo ? keep_lo : 0;
   +        const int64_t hi = keep_hi > keep_lo ? std::min<int64_t>(keep_hi, g.n_layers) : g.n_layers;
   +        const uint64_t db = (uint64_t) (hi - lo) * dense.d_bytes(), ub = (uint64_t) (hi - lo) * dense.u_bytes();
   +        uint8_t *da = nullptr, *ua = nullptr;
   +        if (cudaMalloc((void**) &da, db) != cudaSuccess || cudaMalloc((void**) &ua, ub) != cudaSuccess) {
                const cudaError_t ce = cudaGetLastError();
   -            if (d) { cudaFree(d); d = nullptr; }
   -            if (u) { cudaFree(u); u = nullptr; }
   +            if (da) cudaFree(da);
   +            if (ua) cudaFree(ua);
                err = std::string("the latent D/U device buffers could not be had (") + cudaGetErrorString(ce) + ")";
                return false;
            }
   -        for (int64_t l = 0; l < g.n_layers; ++l)
   +        d = da - (size_t) lo * dense.d_bytes();
   +        u = ua - (size_t) lo * dense.u_bytes();
   +        for (int64_t l = lo; l < hi; ++l)
                if (cudaMemcpy(d + (size_t) l * dense.d_bytes(), dense.d(l), dense.d_bytes(), cudaMemcpyHostToDevice) != cudaSuccess ||
                    cudaMemcpy(u + (size_t) l * dense.u_bytes(), dense.u(l), dense.u_bytes(), cudaMemcpyHostToDevice) != cudaSuccess) {
                    err = "the latent D/U upload failed at layer " + std::to_string(l);
                    return false;
                }
   -        strata::core::latent_shared_d() = d;
   -        strata::core::latent_shared_u() = u;
   +        strata::core::latent_shared_d(dev_now) = d;
   +        strata::core::latent_shared_u(dev_now) = u;
   +        return true;
   +#endif
   +    }
   +
   +    // The host table above is shared; the DEVICE copy is per GPU.  A layer split's later stage has its
   +    // own device, and the pair held here was allocated on the primary - the stage's first layer faults
   +    // on it (illegal memory access at layer 24 on a 24/24 split).  Allocate and upload that device's
   +    // copy once; the stage's prefill and its verifier then share it through latent_shared_d(dev).
   +    // lo..hi_: the stage's layers (hi_ <= lo: all of them); only those layers' D/U is held, base offset as in load().
   +    bool load_device(const strata::core::ModelGeometry& g, int dev, int64_t lo_, int64_t hi_, std::string& err) {
   +        if (!latent) return true;   // (always the case on Metal, whose window owns its buffers)
   +#if defined(STRATA_GPU_METAL)
   +        (void) g; (void) dev; (void) lo_; (void) hi_; (void) err;
   +        return true;
   +#else
   +        if (strata::core::latent_shared_d(dev) != nullptr) return true;
   +        const strata::core::OnDevice on(dev);
   +        const int64_t lo = hi_ > lo_ ? lo_ : 0;
   +        const int64_t hi = hi_ > lo_ ? std::min<int64_t>(hi_, g.n_layers) : g.n_layers;
   +        const uint64_t db = (uint64_t) (hi - lo) * dense.d_bytes(), ub = (uint64_t) (hi - lo) * dense.u_bytes();
   +        uint8_t* da = nullptr;
   +        uint8_t* ua = nullptr;
   +        if (cudaMalloc((void**) &da, db) != cudaSuccess || cudaMalloc((void**) &ua, ub) != cudaSuccess) {
   +            const cudaError_t ce = cudaGetLastError();
   +            if (da) cudaFree(da);
   +            if (ua) cudaFree(ua);
   +            err = "latent dense: the D/U device buffers on GPU " + std::to_string(dev) + " could not be had (" +
   +                  cudaGetErrorString(ce) + ")";
   +            return false;
   +        }
   +        uint8_t* dd = da - (size_t) lo * dense.d_bytes();
   +        uint8_t* uu = ua - (size_t) lo * dense.u_bytes();
   +        for (int64_t l = lo; l < hi; ++l)
   +            if (cudaMemcpy(dd + (size_t) l * dense.d_bytes(), dense.d(l), dense.d_bytes(), cudaMemcpyHostToDevice) != cudaSuccess ||
   +                cudaMemcpy(uu + (size_t) l * dense.u_bytes(), dense.u(l), dense.u_bytes(), cudaMemcpyHostToDevice) != cudaSuccess) {
   +                err = "latent dense: the D/U upload on GPU " + std::to_string(dev) + " failed at layer " + std::to_string(l);
   +                return false;
   +            }
   +        strata::core::latent_shared_d(dev) = dd;
   +        strata::core::latent_shared_u(dev) = uu;
            return true;
    #endif
        }
   @@ -6936,6 +6990,8 @@ int main(int argc, char** argv) {
                         (long long) o.prefill_chunk, (long long) mib, (long long) (160 + (o.prefill_chunk * 680) / 1024));
            return mib + 64;   // a margin for the allocator
        };
   +    // an explicit layer split: the primary's D/U device copy needs only its own layers [0, split_at[0]) (LatentTab::load)
   +    if (multi_gpu && !split_auto && !split_at.empty() && lat_tab.d == nullptr) { lat_tab.keep_lo = 0; lat_tab.keep_hi = split_at[0]; }
        if (o.expert_cache < 0) {
            // latent MoE: the session D/U pair is loaded before the free-VRAM read, so auto leaves room for it
            { std::string le; if (!lat_tab.load(g, le)) { std::fprintf(stderr, "strata generate: %s\n", le.c_str()); return 1; } }
   @@ -7483,6 +7539,19 @@ int main(int argc, char** argv) {
            }
        }
    
   +    // Latent MoE on a layer split: the D/U pair is one per DEVICE (latent_share.hpp).  Load the primary's
   +    // and every later stage's BEFORE any stage sizes its expert cache, so each card's free figure already
   +    // excludes its copy, and the stage's prefill and verifier share that one copy instead of two.
   +    {
   +        std::string le;
   +        {   // the primary's copy, on the primary device whatever device happens to be current here
   +            const strata::core::OnDevice on(primary_device);
   +            if (!lat_tab.load(g, le)) { std::fprintf(stderr, "strata generate: %s\n", le.c_str()); return 1; }
   +        }
   +        if (lat_tab.latent)
   +            for (const auto& stp : stages)
   +                if (!lat_tab.load_device(g, stp->dev, split_auto ? 0 : stp->lb, split_auto ? 0 : stp->le, le)) { std::fprintf(stderr, "strata generate: %s\n", le.c_str()); return 1; }
   +    }
        for (auto& stp : stages) {
            GpuStage& st = *stp;
            const auto& lay = strata::kernels::cpu::expert_layout();

   Patch: cuBLAS warm-up on CUDA (separate)

   --- a/edge0-strata/src/prefill/gemm.cu
   +++ b/edge0-strata/src/prefill/gemm.cu
   @@ -147,6 +147,22 @@
    void ck(cublasStatus_t s, const char* what) {
        if (s != CUBLAS_STATUS_SUCCESS) {
            std::fprintf(stderr, "prefill gemm: %s: cuBLAS status %d\n", what, (int) s);
   +#if !defined(__HIPCC__)
   +        // 1 (not initialized), 3 (alloc failed) and 14 (internal error) are what cuBLAS answers when it needs VRAM
   +        // for a kernel or a handle and the card has none: a layer-split stage left with a few MiB free.  The bare
   +        // status read as an engine fault (and the server restarted the engine every 15 s into the same failure), so
   +        // when the card is that full say so, and which knobs give it room.  The exit is the same.
   +        if (s == CUBLAS_STATUS_NOT_INITIALIZED || s == CUBLAS_STATUS_ALLOC_FAILED || s == CUBLAS_STATUS_INTERNAL_ERROR) {
   +            size_t fb = 0, tb = 0;
   +            int dev = 0;
   +            (void) cudaGetDevice(&dev);
   +            if (cudaMemGetInfo(&fb, &tb) == cudaSuccess && fb < (256ull << 20))
   +                std::fprintf(stderr, "prefill gemm: CUDA%d has only %zu MiB of VRAM free: cuBLAS needs some for its kernels, "
   +                                     "so this is VRAM exhaustion, not a bad install.  Give the card room: raise "
   +                                     "--vram-reserve-later-mib (--vram-reserve-mib on the first card), give it fewer "
   +                                     "layers with --layer-split, or lower --max-context.  Exiting.\n", dev, fb >> 20);
   +        }
   +#endif
            std::exit(1);
        }
    }
   --- a/edge0-strata/src/prefill/prefill.cpp
   +++ b/edge0-strata/src/prefill/prefill.cpp
   @@ -1546,6 +1546,12 @@
    // request; generate.cpp starts it on a thread before the expert arena loads, so it costs no start-up time, and
    // Prefill::init waits for it.  The outputs go to scratch buffers of their own, so no later result changes.
    // STRATA_PREFILL_WARM=0 turns it off.
   +//
   +// CUDA: the same call also fixes WHEN cuBLAS takes its VRAM.  Its kernels load on first use, so on a layer split whose
   +// stage had been sized down to a few tens of MiB free the first prompt's GEMM found no VRAM for them and ended with
   +// `cuBLAS status 14` after the engine had said READY (Edge0-AI/edge0#205 follow-up).  Run here, at the end of
   +// Prefill::init, the loads come out of the memory the stage's reserve already holds back, before the "VRAM free with
   +// everything loaded" check and READY, not out of the first request.
    void warm_gemm_libraries(int device) {
        static const bool off = [] { const char* e = std::getenv("STRATA_PREFILL_WARM"); return e != nullptr && e[0] == '0'; }();
        if (off) return;
   @@ -1568,6 +1574,7 @@
            return wanted;
        }();
        if (!arch_ok) return;
   +#endif
        static std::mutex mu;
        static unsigned done_mask = 0;   // per device: each card loads its own code objects
        if (device < 0 || device >= 32) return;
   @@ -1605,9 +1612,10 @@
        if (ok)
            std::fprintf(stderr, "strata prefill: warmed the BF16 and FP16 GEMM libraries on GPU %d in %.0f ms (STRATA_PREFILL_WARM=0 skips)\n",
                         device, std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - t0).count());
   -#else
   -    (void) device;
   -#endif
   +    else
   +        std::fprintf(stderr, "strata prefill: GPU %d: the GEMM warm-up did not run (no VRAM for its buffers); the card is nearly "
   +                             "full and the first prompt may end with a cuBLAS error - raise --vram-reserve-later-mib "
   +                             "(--vram-reserve-mib on the first card) or give it fewer layers\n", device);
    }
    
    bool Prefill::init(const core::WeightTable& wt, const core::ModelGeometry& g, core::SessionState& ss,