[Engine] Windows HIP (gfx1100): batched PLE block silently misreads prompts in ~1 of 8 engine starts; STRATA_PLE_BATCH=0 fixes it

Setup

Windows 11, RX 7900 XTX (gfx1100) + RX 7900 XT, i9-13900K, 64 GB DDR5; AMD Adrenalin 26.9.2 (driver 32.0.32015.2008)
Engine 0.1.45.3 built locally with tools\hip\build_windows.bat (STRATA_HIP_ARCHS=gfx1100, STRATA_NO_VISION=1), ROCm 10.2.0a20260930 (TheRock wheels)
Qwen3.8-Flash-Next GSQ-RCO IQ3_S, one card (XTX, "gpu": 0), 64K context, --kv int8, MTP --spec 4, --prefill auto, no hipBLASLt tuning table

Symptom
In some engine starts every prompt longer than ~64 tokens is misread for the whole session: no error, no NaN, the model answers fluently but about something else ("your message is just ##", answers in Chinese, an invented question). Prompts of 64 tokens or fewer are always read correctly. Restarting the engine usually clears it. The log of a bad start is identical to a good one (same slots, free VRAM, chunk size); the only visible difference is that a 176-token prompt is read ~15% faster in bad sessions (1.39-1.44 s vs 1.62-1.70 s).

Reproduction
A script that starts the server, sends two prompts (102 and 176 tokens, answer must contain a given word/number) and stops it, repeated:

Configuration	Bad starts
default	3 / 24
`STRATA_PF_STEP_SYNC=1`	2 / 16
`STRATA_WMMA_GEMM=1`	1 / 24
`STRATA_PLE_BATCH=0`	0 / 24

So the batched PLE block at layer 1 (prefill.cpp, if (l == 1 && ple_on && ple_batch): the BF16 key/value GEMMs + native_ple_postops_batch) is involved. It does not look like a race between prefill steps (step sync does not help) nor hipBLAS alone (WMMA GEMMs still failed once). The bad state seems to be decided once per engine start and then stays for the session. Possibly related: #224 (illegal memory access in the same block) and #541 (HIP stall avoided by STRATA_PLE_BATCH=0), both closed - here it fails silently instead.

Workaround
"env": {"STRATA_PLE_BATCH": "0"} in the run config. Cost on this machine: prompt 810 -> 710 tok/s at 3.5K tokens, 865 -> 814 tok/s at 10.7K; decode unchanged (52-70 tok/s).

Logs of failing starts and the test script are available if useful.

Same symptom on a second gfx1100 card with engine 0.1.47 (Strata 0.2.2, release package strata-engine-0.1.47-windows-x86_64-hip-rdna3), and here STRATA_PLE_BATCH=0 did not prevent it.

Setup: RX 7900 XTX 24 GB, Ryzen 9 9950X, 62.7 GiB RAM, Windows 11 (10.0.26300), driver 32.0.32015.7029. Swift 1.5 Qwen3.8-Flash-Next IQ3_S with --ple-gguf (log: PLE on), --max-context 131072 --kv int8 --spec 4 --spec-min-p 0.5 --mtp <draft> --prefill auto --expert-cache auto, STRATA_HC_Q8=1.

Check per fresh start: tools/needle_bench.py --lengths 4k (three prompts of about 3.3K tokens, greedy, thinking off). Default and STRATA_PLE_BATCH=0 (in the config's env) starts alternated, 10 each:

Setting	Bad starts
default	1 of 10
`STRATA_PLE_BATCH=0`	1 of 10

The variable reached the engine: in every STRATA_PLE_BATCH=0 start the same 3.3K prompts read in 4.26-4.30 s, against 4.07-4.15 s with the default.

A bad start: all three needles missed, and every answer ran to the 40-token cap describing something else ("Based on the visual cues provided in your prompt ..."). Short questions (43-51 tokens) looked normal. Earlier, two full strata_bench.py --long runs had also landed on bad starts: there the 609-token first agent turn ran to its 64-token cap every time instead of a reply of about 12 tokens. Restarting cleared it. The engine log of a bad start has the same lines as a good one, with no errors.

Timing: at 3.3K tokens good and bad starts read at the same speed. In the two bad full runs the 609-token first agent turn read in 2.28-2.34 s (engine prompt time, 6 of 6 turns), against 3.00-3.20 s in six good runs (18 of 18), which matches your "bad sessions read faster".

Also seen once: in an otherwise good start, one 609-token request was misread (it answered with unrelated code; read time normal, 3.10 s) and the requests after it were fine.

The per-start logs and /api/requests dumps are available if useful.

@ahmetem

Thanks for the second data point, @ahmetem, and for the needle_bench check. I opened this issue; here is how my numbers compare, in case it helps narrow things down.

My data (engine 0.1.45.3, built locally for gfx1100, ROCm 10.2.0a20260930, Windows 11, RX 7900 XTX + XT, Adrenalin 26.9.2; Qwen3.8-Flash-Next IQ3_S, single card, 64K context, --kv int8, MTP, --prefill auto). Scripted loop, one engine start per iteration, two prompts per start (102 and 176 tokens, answer must contain a known word), 88 starts in total:

Configuration	Bad starts
default	3 / 24
`STRATA_PF_STEP_SYNC=1`	2 / 16
`STRATA_WMMA_GEMM=1`	1 / 24
`STRATA_PLE_BATCH=0`	0 / 24

Bad starts looked exactly like what you describe: the model answers fluently but about something else (e.g. "Based on the fragments you provided…", "I don't have access to the context needed…"), short prompts are fine, the engine log is identical to a good start, and a restart clears it. Like yours, my bad sessions read the longer prompt faster than good ones (about 15% faster for the 176-token prompt), which makes me suspect that a bad start takes a cheaper code path, or skips some work, rather than computing something wrong in a random way (this is only a guess).

How this fits your result.

If the true rate with STRATA_PLE_BATCH=0 were about 1 in 10 like yours, getting 0 of 24 would happen roughly 8% of the time, so my data does not exclude that the variable only lowers the rate, or that it behaves differently between 0.1.45.3 and 0.1.47. Is there any chance that 0.1.47 changed how the PLE block / prefill is dispatched, so the variable no longer reaches the same path?
My check used short prompts (102 and 176 tokens); yours uses about 3.3K. A longer prompt may expose starts that mine misses.
I have not re-run the 24-start check since I changed my setup (256K context, --resident-experts, vision reader), so I cannot say whether the 0/24 still holds there.

What I can do. I can repeat the loop with needle_bench.py --lengths 4k on 0.1.45.3, 20+ alternating starts with and without STRATA_PLE_BATCH=0, and post the per-start logs (I already have the logs of the failing starts from the first series). If a maintainer knows of any state that is decided once per engine start (kernel or algorithm selection, an allocation, a cache) and then stays fixed for the session, that would explain why a restart clears it and where to look.

Thanks, the 0/24 with short prompts is a useful contrast.

On the dispatch question: I compared the source of the two releases I have here, 0.1.45 and 0.2.2 (engine 0.1.47). The STRATA_PLE_BATCH checks read the same in both: prefill.cpp (ple_batch_env, off when the value is 0) and verify.cpp (the S25 key/value batch). No PLE kernel source differs between the two trees; the only PLE files that changed are ple_fp8_parity.cpp and ple_reader_test.cpp, both tests. So on 0.1.47 the variable should still turn off the same paths, and the timing says it did change something (+4 % prompt time in every start with it). That leaves two readings that 1 of 10 against 0 of 24 cannot separate: a second start-time cause besides batched PLE, or a lower but non-zero rate with the variable off.

One data point on the "cheaper path" idea: at 3.3K tokens my bad starts read at the same speed as the good ones. With the default, the bad start read its second and third needle prompts in 4.11 / 4.11 s, against 4.07-4.15 s in the good starts. With STRATA_PLE_BATCH=0 it was 4.28 / 4.26 s, against 4.26-4.30 s. Only the 609-token first agent turn of the two bad full-suite runs was faster (2.28-2.34 s against 3.00-3.20 s). So the speed difference may depend on prompt length, or show up only on some prompts.

A needle_bench.py --lengths 4k series on 0.1.45.3 with alternating starts would be a good comparison, since it would show whether the longer prompt also finds bad starts with the variable off on your build.

@ahmetem

I ran the series @ahmetem suggested, on my build.

Setup. Engine 0.1.45.3 (built locally for gfx1100, ROCm 10.2.0a20260930), Windows 11, RX 7900 XTX, Adrenalin 26.9.2. Qwen3.8-Flash-Next IQ3_S, single card, --max-context 131072 --kv int8, MTP, --prefill auto, --expert-cache auto, default (full) expert arena, no vision reader, no hipBLASLt table. 24 engine starts, alternating default and STRATA_PLE_BATCH=0 (12 each); after every start, tools/needle_bench.py --lengths 4k --depths 10,50,90 (three prompts of about 4,000 tokens, greedy, thinking off).

Configuration	Bad starts
default	2 / 12 (starts 7 and 13 of 24)
`STRATA_PLE_BATCH=0`	1 / 12 (start 16 of 24)

A bad start missed all three needles (0/3); every good start found all three (3/3). The bad answers were fluent and about something else ("Based on the visual cues from your description, here is an analysis of the situation…", "Based on the analysis of the provided text, here is the solution to complete the sentence…"), the same symptom you describe.

What this says

On my build STRATA_PLE_BATCH=0 does not prevent bad starts: 1 of 12 with it, so the title of this issue was too optimistic. 2/12 against 1/12 is far too few starts to say whether it lowers the rate at all.
My first series (102- and 176-token prompts) gave 0/24 with the variable. The same variable now gives 1/12 with ~4K prompts, so the short prompts were probably not sensitive enough to find these starts. That fits your point that the longer prompt is a better check.
Read time of the three 4K prompts: good starts 4.5-4.8 s with the default and 4.7-5.0 s with the variable; the bad starts were slightly slower, not faster (5.3 / 5.0 / 4.8 s with the default, 5.5 / 5.1 / 5.2 s with the variable). So at 4K the "bad starts read faster" pattern from my short prompts does not show up, which agrees with your 3.3K numbers. Start-up time was 29-35 s in every start.
This is on 0.1.45.3 only; I have not tried 0.2.2 / 0.1.47. It also cannot tell your two readings apart (a second start-time cause besides batched PLE, or a lower but non-zero rate with the variable off).

I have the engine logs of the three bad starts and the per-start CSV; I have not compared the bad logs with good ones yet. If a maintainer has an idea for a variable or a flag worth trying (I can repeat 20+ starts in about 20 minutes), or knows of state that is decided once per engine start and then stays fixed for the session, I'm happy to run it and post the results.

@ahmetem

Update with a root cause, and a fix that works on my build. Short version: this is not the PLE. It is a collision between the address a thread's stack lives at and the GPU-side alias the HIP runtime gives to the pinned expert arena.

What happens

The engine pins the expert arena with cudaHostRegister (PinnedArena, src/core/pinned.cu; on my machine the whole-arena register is refused and it falls back to 47 slices of about 1 GiB). With the Mapped flag the runtime gives every registered range an alias in the GPU's address space (cudaHostGetDevicePointer).
The aliases follow a fixed sequence; the place of a thread's stack changes with every process (ASLR). In about 1 start of 10 an alias covers the address of the stack of the engine thread.
From then on, the runtime treats any pointer inside that alias as the registered memory. A small cudaMemcpy from or to a local variable "succeeds" (no error) but copies nothing, for the rest of the process. Everything that passes scalars and small parameters through the stack (the long-prompt paths) computes with wrong values; short prompts mostly do not use them. That is the whole symptom: fluent answers about something else for prompts longer than about 64 tokens, until restart.

How I found it

Weights on the device are byte-identical between a good and a bad start (hash of every tensor of blocks 0-2). The key GEMM of the PLE block gives about 1e-12 instead of 0.03 in a bad start, and the same wrong result with a fresh GEMM handle and with fresh copies of the inputs and weights. After the first half of layer 0 the residual already differs.
A 16-byte round trip (host to device to host, compare) at several points of the process: fine until the arena is pinned, then it returns zeros in bad starts, from one slice on. Which slice differs by start (5, 8, 17, 19, 25, 29, 33, 35, 36, 40, 43, 44...), and registering the arena again at another address hits the same slice, so it does not depend on the arena's host address or pages.
At the moment of failure, hipPointerGetAttributes on a plain stack variable returns type 1 (host registered) with hostPointer inside the arena, and its address lies inside the alias range of the bad slice. Three cases:

stack variable alias range of the bad slice

0x14D38DB0E0 0x14961D0000 .. 0x14DD34F000

0x480ECDAD60 0x47DC880000 .. 0x48239FF000

0x0C8D52AF60 0x0C5AEC0000 .. 0x0C9D53F000

Bisecting the bad slice shows that a zone of only 45-60 KB (the part of the stack in use) is enough to break it.
That explains what we saw earlier: STRATA_PLE_BATCH=0 and the other variables only changed which copies happened to go through the stack, so they moved the odds without removing the cause; STRATA_WMMA_GEMM=1, STRATA_PREFILL_WARM=0 etc. made no difference (2/16 and 2/11 bad starts).

Fix on my build (Windows only, src/program/generate.cpp)

main now reserves (MEM_RESERVE, nothing committed) all the free address space below 1 TiB, then runs the old main on a new thread. The stacks of that thread, and of the threads it starts, are created above 1 TiB, away from the aliases (I saw them at 0x10003FFF8CC and above); the first thread only waits. STRATA_STACK_ISOLATE=0 turns it off. In pinned.cu a 16-byte round trip after each slice is registered unpins the arena if it ever fails (slower but correct).

// sketch
uint64_t reserve_low_address_space(uint64_t limit) {
    SYSTEM_INFO si; GetSystemInfo(&si);
    uint8_t* at = (uint8_t*) si.lpMinimumApplicationAddress; MEMORY_BASIC_INFORMATION mbi; uint64_t total = 0;
    while ((uint64_t) at < limit && VirtualQuery(at, &mbi, sizeof mbi) == sizeof mbi) {
        uint8_t* end = (uint8_t*) mbi.BaseAddress + mbi.RegionSize;
        if (mbi.State == MEM_FREE) {
            uint64_t b = ((uint64_t) mbi.BaseAddress + 0xFFFF) & ~0xFFFFull, e = std::min<uint64_t>((uint64_t) end, limit) & ~0xFFFFull;
            if (e > b && VirtualAlloc((void*) b, e - b, MEM_RESERVE, PAGE_NOACCESS)) total += e - b;
        }
        if (end <= at) break; at = end;
    }
    return total;
}
// main(): reserve_low_address_space(1ull << 40); CreateThread(nullptr, 64u << 20, engine_thread, &args, STACK_SIZE_PARAM_IS_A_RESERVATION, nullptr);

Results (engine 0.1.45.3, RX 7900 XTX, Qwen3.8-Flash-Next IQ3_S, 128K context, --kv int8, MTP; every start checked with a 1,150-token prompt with a code word, and needle_bench.py --lengths 4k,16k --depths 10,50,90)

Build	Starts	Bad starts
before the fix (series with the instrumented build, same machine)	178	17 (about 10 %)
with the fix, `STRATA_PLE_BATCH` unset (four series; the last one with the instrumentation removed)	210	0

In every one of the 210 starts the arena was pinned in all 47 slices and the probe never failed. Prompt speed at 15K tokens went from 801 tok/s (with STRATA_PLE_BATCH=0) to 830-835 tok/s (unset); write speed is unchanged.

What I have not checked

Only my 0.1.45.3 build with ROCm 10.2.0a20260930; not 0.2.2 / engine 0.1.47, nor Linux (I only saw it on Windows and have no data on Linux).
It would help if @ahmetem could run the same 16-byte round trip after the arena is pinned on his build (a bad start should return zeros). If that is what he sees, the fix above should apply to his build too.
Whether a maintainer would prefer a different place for the fix (e.g. allocating the arena in a way that avoids the alias range, or running the engine thread with a stack allocated by the engine). I'm happy to post the patch as a diff if useful.

I ran a check that needs no rebuild on the prebuilt 0.1.47 (Strata 0.2.2, hip-rdna3 release package). STRATA_ARENA_PIN_GIB=none leaves the expert arena unregistered (log: expert arena: NOT registered with the GPU driver ...; locked 47965 MiB via working-set minimum + VirtualLock), so the arena gets no mapped alias at all.

Same machine and model as before (RX 7900 XTX, Swift 1.5 Flash-Next IQ3_S, 131K context, --kv int8, MTP), fresh starts, needle_bench.py --lengths 4k after each:

Setting	Bad starts
default (arena registered in 46 slices)	1 / 10
`STRATA_PLE_BATCH=0`	1 / 10
`STRATA_ARENA_PIN_GIB=none`	0 / 30

If the rate were still 1 in 10, 0 of 30 would happen about 4 % of the time. So this supports your explanation on 0.1.47 too, but it is not proof on its own. I can't run the 16-byte round trip here, because this is the release binary, not a build.

The workaround has a cost on this PC: the experts reach the GPU by copy, so prompts read noticeably slower. 4K went from 956 to 772-787 tok/s and 16K from 1,207 to 899-918 tok/s (two strata_bench.py --runs 1 runs with the variable, against the 3-run medians of a default start); the 3.3K needle prompts went from 4.07-4.15 s to 4.51-4.76 s. Decode moved by less than the noise (short chat 61-71 tok/s, agent turns 79-80, against 71 and 69). Your stack-isolation fix keeps the arena pinned, so it is the better fix. STRATA_ARENA_PIN_GIB=none is a slower stopgap for anyone on a release binary until a fix ships.

@giacomobord-prog

I have hit the same bug on RDNA4.

Running an ASRock Radeon AI PRO R9700 Creator 32GB (gfx1201) on an Intel platform (Core Ultra 7 270k Plus, ASUS PRO WS880-ACE SE), Windows 11 Pro.

Software:

Strata 0.1.44.1
Model: Qwen 3.8-Flash-next Q5_K_L & Q8_0 (Bartowski, Hugging Face)
AMD Drivers 32.0.32015.2008 and 32.0.31041.5005 (Occurs on both)

When running with kv int8 (tried with f16 as well) see the same markers (KV streaming 100.00% of 0 block reads) and same fluent answers about unrelated subjects on long prompts until restart.

Been at it for a week and 5 in 70 starts have had this issue. Short prompts are unaffected.

@giacomobord-prog, could you open your stack-isolation fix as a PR? It looks like the fix for this issue, and I am happy to test it on gfx1201.

The worst part is that a bad start returns HTTP 200 with fluent text. When an automated setup is used to switch models, it silently produces wrong results! Love the speed, but this is unreliable.

@sylvesterpious

@sylvesterpious thank you for the data point: gfx1201, an official 0.1.44.1 and two driver versions make it very unlikely that this is my build or one card. The marker you quote ("KV streaming 100.00% of 0 block reads") is a good cheap detector of a bad start.

I'd rather not open a PR against a tree I have not built, so here is the fix as a diff, and what I would like to learn from a test on your card.

What the fix does (Windows only, two files, ~110 lines)

generate.cpp: main becomes engine_main. A new main first reserves (MEM_RESERVE, nothing committed) all the free address space below 1 TiB, then runs engine_main on a new thread with a 64 MiB stack reservation. Stacks created after the reservation (that thread's and the ones it starts) land above 1 TiB, away from the GPU-side aliases of the pinned arena. STRATA_STACK_ISOLATE=0 turns it off. The log shows one line: strata: the engine runs on its own thread (1023 GiB of low address space kept free of stacks ...).
pinned.cu: after each slice is registered, a 16-byte host-to-device-to-host round trip; if it does not come back, the arena is unpinned (slower but correct). It never fired in the 210 starts below, because the first change keeps the stacks out of the way. STRATA_PIN_REPAIR=0 turns it off.

Numbers (gfx1100, RX 7900 XTX, 0.1.45.3 sources, ROCm 10.2.0a20260930, Windows 11): 17 bad starts in 178 before; 0 in 210 with the fix (a start is "bad" when a 1,150-token prompt with a code word does not return the word). The pinned arena stayed in 47 slices every time, prompt speed is unchanged, and a 192K-token prompt at 256K context still reads at about 1,220 tok/s.

What I could not do: it is tested on one card only, against the 0.1.45.3 tree (not current main), and I have not ported it to 0.2.3. If a maintainer prefers a PR against main, say so and I will do it.

If you can build from source on gfx1201, an A/B would be the most useful thing: the same start loop you ran (needle_bench.py --lengths 4k after each start, or your marker) with and without STRATA_STACK_ISOLATE=0, 50+ starts each; plus, if you can, the 16-byte probe after the arena is pinned in a bad start (a bad start returns zeros). If you only have the release binary, @ahmetem measured a workaround that needs no rebuild: STRATA_ARENA_PIN_GIB=none (0 bad starts in 30 on 0.1.47, but prompts read about 20% slower).

On your last point: I agree that HTTP 200 with fluent text is the worst part. On my side I added a start check in the server (a 1,150-token prompt with a hidden word, restart the engine if it is missed). It cost about 10 s per start and caught every bad start in 24 of mine, so a setup that switches models automatically can use the same idea until the root cause is fixed upstream.

generate.cpp diff

--- program/generate.cpp.orig_20261022b	2026-10-22 20:28:44.621141500 +0200
+++ program/generate.cpp	2026-10-23 10:59:33.783738800 +0200
@@ -2211,7 +2211,7 @@
     return el > 0 ? (double) done / el : -1.0;
 }
 
-int main(int argc, char** argv) {
+static int engine_main(int argc, char** argv) {
     // **UNBUFFERED, BECAUSE THE INTERESTING OUTPUT IS THE OUTPUT BEFORE A CRASH.**  `stdout` redirected to a
     // pipe or a file is block-buffered, so a program that dies loses every line it had already printed - which
     // turns "it crashed at step 7" into "it crashed somewhere", and the difference is a debugging session.
@@ -16458,3 +16458,61 @@
     cudaFree(arena);
     return 0;
 }
+
+
+// Issue Niko1221/Strata#3431.  On Windows the HIP runtime gives every cudaHostRegister'd range an
+// alias in the GPU's address space, and a pointer whose value falls inside an alias is taken for that registered memory:
+// when a thread's stack lies there (the stack's place changes with every process, ~1 start in 10) the small host<->device
+// copies from local variables stop copying anything for the rest of the process.  So the engine runs on a thread created
+// after all free address space below 1 TiB was reserved (reserved, not committed): its stack, and the stacks of the
+// threads it starts, are placed above, away from the aliases.  The process's first thread only waits.
+// STRATA_STACK_ISOLATE=0 turns it off.
+#if defined(_WIN32)
+namespace {
+struct MainArgs { int argc; char** argv; int rc; };
+DWORD WINAPI engine_thread(LPVOID p) {
+    MainArgs* a = static_cast<MainArgs*>(p);
+    a->rc = engine_main(a->argc, a->argv);
+    return 0;
+}
+uint64_t reserve_low_address_space(uint64_t limit) {
+    SYSTEM_INFO si;
+    GetSystemInfo(&si);
+    uint64_t total = 0;
+    uint8_t* at = static_cast<uint8_t*>(si.lpMinimumApplicationAddress);
+    MEMORY_BASIC_INFORMATION mbi;
+    while (reinterpret_cast<uint64_t>(at) < limit && VirtualQuery(at, &mbi, sizeof mbi) == sizeof mbi) {
+        uint8_t* end = static_cast<uint8_t*>(mbi.BaseAddress) + mbi.RegionSize;
+        if (mbi.State == MEM_FREE) {
+            uint64_t b = (reinterpret_cast<uint64_t>(mbi.BaseAddress) + 0xFFFF) & ~uint64_t(0xFFFF);
+            uint64_t e = std::min<uint64_t>(reinterpret_cast<uint64_t>(end), limit) & ~uint64_t(0xFFFF);
+            if (e > b && VirtualAlloc(reinterpret_cast<void*>(b), (SIZE_T) (e - b), MEM_RESERVE, PAGE_NOACCESS) != nullptr) total += e - b;
+        }
+        if (end <= at) break;
+        at = end;
+    }
+    return total;
+}
+}  // namespace
+#endif
+
+int main(int argc, char** argv) {
+#if defined(_WIN32)
+    const char* iso = std::getenv("STRATA_STACK_ISOLATE");
+    if (iso == nullptr || iso[0] != '0') {
+        const uint64_t reserved = reserve_low_address_space(1ull << 40);
+        MainArgs a{argc, argv, 1};
+        HANDLE h = CreateThread(nullptr, 64u << 20, engine_thread, &a, STACK_SIZE_PARAM_IS_A_RESERVATION, nullptr);
+        if (h != nullptr) {
+            int probe = 0;
+            std::fprintf(stderr, "strata: the engine runs on its own thread (%llu GiB of low address space kept free of stacks; first thread's stack at %p)\n",
+                         (unsigned long long) (reserved >> 30), (void*) &probe);
+            WaitForSingleObject(h, INFINITE);
+            CloseHandle(h);
+            return a.rc;
+        }
+        std::fprintf(stderr, "strata: could not start the engine thread, running on the first one\n");
+    }
+#endif
+    return engine_main(argc, argv);
+}

pinned.cu diff

--- core/pinned.cu.orig_20261022	2026-10-22 20:40:29.067046000 +0200
+++ core/pinned.cu	2026-10-23 10:59:23.890858400 +0200
@@ -394,6 +394,20 @@
 }
 }  // namespace
 
+// 16 bytes host -> device -> host on the current device: true when they come back (see the pin loop's repair)
+static bool pin_copy_health_ok() {
+    uint32_t src[4] = {0x11223344u, 0x55667788u, 0x99aabbccu, 0xddeeff00u}, dst[4] = {0, 0, 0, 0};
+    void* d = nullptr;
+    bool ok = false;
+    if (cudaMalloc(&d, sizeof src) == cudaSuccess) {
+        ok = cudaMemcpy(d, src, sizeof src, cudaMemcpyHostToDevice) == cudaSuccess &&
+             cudaMemcpy(dst, d, sizeof dst, cudaMemcpyDeviceToHost) == cudaSuccess && std::memcmp(src, dst, sizeof src) == 0;
+        cudaFree(d);
+    }
+    (void) cudaGetLastError();
+    return ok;
+}
+
 PinnedArena::PinnedArena(uint64_t bytes, uint64_t slice) : PinnedArena(bytes, uniform_bounds(bytes, slice)) {
     if (slice_bytes) slice_bytes = slice;   // sliced registration: record the uniform size
 }
@@ -438,6 +452,11 @@
             // 64 GB PC pins 30 GiB past that budget without trouble, and capping it at 26 cost ~20% prompt speed.
             if (!capped && env_gib == -2) limited = sliced_pin_limit(limit, limit_why);
 #endif
+            // Issue Niko1221/Strata#3431: on Windows the HIP runtime gives each registered range an alias in the GPU's address
+            // space, and a pointer inside an alias is taken for that registered memory.  If a thread's stack is there, small
+            // host<->device copies copy nothing for the rest of the process.  A 16-byte round trip after each slice tells
+            // (STRATA_PIN_REPAIR=0 turns the check off); main() already keeps the engine's threads clear of the aliases.
+            static const bool repair_on = [] { const char* v = std::getenv("STRATA_PIN_REPAIR"); return v == nullptr || v[0] != '0'; }();
             for (size_t i = 0; i + 1 < bounds.size(); ++i) {
                 const uint64_t off = bounds[i], n = bounds[i + 1] - bounds[i];
                 if (limited && (off > limit || n > limit - off)) break;
@@ -448,6 +467,18 @@
                 slice_starts.push_back(off);
                 registered_bytes = off + n;
                 ++registered_slices;
+                if (repair_on && !pin_copy_health_ok()) {
+                    // The runtime's alias of a registered range covers a pointer the engine uses for small copies (a stack):
+                    // they would copy nothing for the rest of the process.  main() keeps the engine's threads away from
+                    // that (see engine_thread); if it still happens the arena is left unpinned, which is slower but right.
+                    std::fprintf(stderr, "strata pin: small copies stopped working after pinning slice %zu: unpinning the arena\n", i);
+                    for (size_t k2 = slice_starts.size(); k2-- > 0;) (void) cudaHostUnregister((uint8_t*) base + slice_starts[k2]);
+                    (void) cudaGetLastError();
+                    slice_starts.clear();
+                    registered_bytes = 0;
+                    registered_slices = 0;
+                    break;
+                }
             }
             char gib[32];
             std::snprintf(gib, sizeof gib, "%.1f", (double) limit / (double) (1ull << 30));

Release-binary A/B on a second RX 7900 XTX (gfx1100), Windows 11, as you asked. Qwen3.8-Flash-Next IQ3_XXS, --kv int8, MTP, --kv-resident 32768, the hip-rdna3 packages as published. After each fresh start: a 189-token prompt whose answer is in the text, then a fresh 27K prompt with a known answer. A bad start is the one you describe: "KV streaming ... of 0 block reads", fluent answers about an "image" that was never sent, while the short prompt stays right.

engine	setting	bad starts
0.1.44.1	default	1 / 10
0.1.50	default	1 / 20
0.1.50	`STRATA_ARENA_PIN_GIB=none`	0 / 40
0.1.52	`STRATA_ARENA_PIN_GIB=none`	0 / 20

Over the last ten days 0.1.44.1 had 5 bad starts in 80 here, so 0 in 60 with none fits your explanation. The log line confirms the setting took: expert arena: NOT registered with the GPU driver.

Cost on this PC, production config (131072 context), two rounds each: fresh prompts read 15-20% slower (27K in 26.7-28.0 s instead of 22.5-23.8 s, 9K in 10.7-11.5 s instead of 9.1-9.5 s). Writing (56-64 tok/s) and pictures are unchanged. On 0.1.52, STRATA_KVA=1 wins that back and more: four starts, ten fresh prompts each with a known answer (a sentence hidden at the start, middle or end of 20K to 100K tokens), 40/40 right with it and 40/40 without, and 27K read in 23.1 s instead of 26.8 s, 40K in 30.5 instead of 38.1 s, 100K in 75.9 instead of 95.6 s.

0.1.44.1 does not know none (only a cap), so on that version the workaround means updating. I'm running 0.1.52 with none in production now.

@giacomobord-prog

I tested a third RX 7900 XTX (gfx1100), with engine 0.1.54 (Strata 0.2.8) on Windows 10 19044. PR #4787 has a fix. The fix builds on the stack-isolation diff of @giacomobord-prog. This comment gives four findings that are not in this thread yet. The first finding helps all users of a release binary now.

Setup

RX 7900 XTX 24 GB, Ryzen 9 5900X, 64 GB RAM, driver 32.0.32015.7029.
Swift 1.5 Flash-Next IQ3_XXS, 131072 context, --kv int8, --kv-resident 32768, MTP --spec 4, --expert-cache auto.
A local v0.2.8 build had 4 bad starts in 30. The release package had 3 bad starts in 12 in production.

1. The 0 block reads marker shows on the first request, also on a short one.

In a bad start, every request writes KV streaming: 100.00% of 0 block reads to the log. A 22-token prompt also writes this line, also when its answer is correct.
In a good start, the same 22-token prompt writes of 1176 block reads as the first request.

No start had both kinds of request:

Check	Starts	Bad starts
A 7,118-token code-word prompt, then a short prompt	378	5, every request 0 reads
A short prompt as the first request, then the code-word prompt	39	2, both found
The production log of the release package	12	3, every request 0 reads

So a start check does not need a long prompt. The check sends one short request (8 output tokens, thinking off). Then it reads the new lines of the engine log. The check takes less than 1 s. A long code-word prompt takes about 10 s.

On this PC, a small wrapper script does this check:

The script starts serve/server.py on a private port.
The script sends the short request and reads the engine log.
If the log shows 0 block reads, the script stops the server and starts it again. The script tries at most 5 times.
After a good check, the script opens the client port and forwards it to the server. So no client request reaches a bad start.

A start that passed the check was ready in 28.3 s. Without the script, a start was ready in 27.7 s. In 30 starts, the script found 2 bad starts and started them again. All 30 code-word answers were correct. The script uses only the Python standard library. I can post it if it is useful.

2. A bad start also writes into the pinned arena.

An earlier comment says that the small copies copy nothing. Our measurement shows that the copies go to the arena. A stack variable can lie inside the alias of a slice. The runtime then uses the address "slice host address + (variable − alias)" instead of the variable:

A host-to-device copy from a stack variable reads the arena at that address.
A device-to-host copy into a stack variable writes into the arena at that address.

A standalone program shows this behavior every time:

The program registers a range and reads the alias of the range.
The program unregisters the range and puts a buffer at the alias address.
The program registers the range again. The runtime gives the same alias again.
A 16-byte round trip through the buffer returns zeros. The device-to-host copy changed the registered range.

So a bad start can also change expert data in RAM, in addition to the values on the stack.

3. Why a stack can lie under an alias

The runtime reserves the CPU address range of cudaMalloc memory. VirtualQuery shows MEM_RESERVE for that range.
The runtime does not reserve the address range under a register alias. VirtualQuery shows MEM_FREE for that range. So Windows can put a new stack there.
The aliases were between 8 and 300 GiB.
cudaHostAlloc memory has the same host and device address here. So that memory has no separate alias.

This behavior is probably a defect in the HIP runtime on Windows. A short search found no ROCm report of this defect.

4. Memory from before main can still lie under an alias.

The diff puts every later stack above 1 TiB. But the stack of the first thread stays low. A few memory blocks from before main also stay low. In 7 of 110 starts (about 6 %), an alias covered such memory. I do not know if these starts give wrong answers. The first thread only waits, and the other blocks are not known stacks.

So the PR does two more things:

The PR reserves the address range of each alias.
The PR does not pin a range when its alias covers memory in use. The log gives the reason.

A start with such a range gives correct answers, but it reads prompts 2–9 % slower. The PR had 0 bad starts in 110 starts.

The workaround on this PC

STRATA_ARENA_PIN_GIB=none made prompts 9–12 % slower here. I ran strata_bench.py --long two times for each setting, in alternation. The values are medians of 3 runs:

Median	Pinned, run 1 / run 2	`none`, run 1 / run 2
4K prompt tok/s	840 / 833	740 / 731
16K prompt tok/s	1801 / 1798	1620 / 1610
32K prompt tok/s	1794 / 1793	1639 / 1636
chat decode tok/s	73.5 / 68.8	65.1 / 67.6

Other users in this thread measured 15–25 %. So the cost changes from PC to PC. In the same series, the PR build had the speed of the pinned release package.

@giacomobord-prog

@giacomobord-prog # Strata issue #3431 patch validation - final report (31 October 2026)

On my machine the issue #3431 patch (stack isolation + pinned-slice copy check, Oct 26 comment by giacomobord-prog) removed the startup corruption we see:
13 of 150 unpatched (control) fresh processes were bad starts; 0 of 150 patched processes were. Every control failure had the issue's marker (KV streaming ... 0 block reads) and fluent but unrelated answers on all three needle prompts. No other failure type occurred.

Design (same in all runs)

Build: public Strata v0.1.45 (477ac13) + the exact two-file patch (sha256 26d3353d...), ROCm 10.2.0a20260930, AMD clang 24.0.0git, MSVC 14.44.35207, SDK 10.0.22621.0. One executable for both arms, sha256 d148d079...ad8da, unchanged before and after every run.
Arms differ only in two environment variables: control STRATA_STACK_ISOLATE=0, STRATA_PIN_REPAIR=0; patched both 1. The stack-isolation log line appeared only in patched processes (checked per process by the runner).
Machine: Radeon AI PRO R9700 (gfx1201), driver 32.0.31041.5005, Windows 11 Pro 10.0.26300, Core Ultra 7 270K Plus (24 cores), 253.16 GiB OS-visible RAM (256 GiB installed).
Models: Q5_K_L and Q8_0 Qwen3.8 Flash Next, 128K context, fp16 KV, 32768 resident KV cells, MTP spec 4. Each fresh native process (talking to --serve directly, so no dashboard startup check and no hidden retries) answered three exact 4096-token needle prompts (10/50/90% depth), greedy, thinking off, 40 output tokens. The three prompts are one startup observation. Matched prompts within each pair, AB/BA alternating order, every attempt kept, no discards.
The production startup check/retry mitigation was deliberately not in the loop: these are raw first-start rates.

System information

Item	Value
GPU	AMD Radeon AI PRO R9700, gfx1201, 32 GB (31.7 GiB exposed by WDDM)
GPU driver	32.0.31041.5005 (Adrenalin PRO line); HIP runtime from ROCm 10.2.0a20260930 build, bundled in the package
CPU	Intel Core Ultra 7 270K Plus, 24 cores / 24 threads
Motherboard / BIOS	ASUS Pro WS W880-ACE SE, BIOS 3401
RAM	256 GiB installed DDR5 (256 GiB across the DIMMs; 253.16 GiB OS-visible), running at 5600 MT/s
OS	Windows 11 Pro 10.0.26300
Build toolchain	AMD clang 24.0.0git (4f43f474), MSVC 14.44.35207, Windows SDK 10.0.22621.0; pinned llama.cpp/ggml 3cf03257
Strata source	v0.1.45, commit `477ac13`, plus the issue #3431 patch (sha256 26d3353d...2038)
Binary	strata.exe sha256 d148d079...ad8da; package ZIP sha256 5e16e0e2...bd38; depends on the Microsoft-signed System32 msvcp140_atomic_wait.dll 14.50.35719.0
Strata options	128K context, fp16 KV, --kv-resident 32768, --spec 4, --pool-workers 7 --pool-affinity p-cores, --vram-reserve-mib 2048, prompt and conversation caches 0
Background	GPU keep-alive process running during every run (the machine's standard mitigation for the idle power-down crash); production dashboard stopped

Results: bad-start rate and significance

Fisher exact test, two-sided; Wilson 95% intervals.

Run	Control bad	Patched bad	Fisher p
main-v1 (100 processes)	4/50 (8.0%, 3.2-18.8%)	0/50 (0-7.1%)	0.117 (not significant alone)
main-v2 (200 processes)	9/100 (9.0%, 4.8-16.2%)	0/100 (0-3.7%)	0.003
Combined (300)	13/150 (8.7%, 5.1-14.3%)	0/150 (0-2.5%)	0.0002
Combined Q5	7/75 (9.3%, 4.6-18.0%)	0/75 (0-4.9%)	0.013
Combined Q8	6/75 (8.0%, 3.7-16.4%)	0/75 (0-4.9%)	0.028

All 13 failures were zero-KV starts (Q5 7, Q8 6); 0 answer misses with positive KV, 0 pin-probe fallbacks, 0 device/protocol/cleanup failures. Correction to an earlier chat remark: main-v1 alone is not significant (p = 0.117), not p = 0.04; significance comes from main-v2 and the pooled data. Combined rates are similar across the two runs (8% and 9%), supporting repeatability.

What this does and does not show:

Strong evidence that the patch prevents the startup failure on this setup. 0/150 does not prove the rate is zero (upper 95% bound 2.5%).
Tests assume independent starts; starts were interleaved and alternated but share one machine/driver/runtime, so unmeasured shared conditions cannot be excluded.
The patch bundles two changes (engine-thread stack isolation and the pin round-trip check); this test cannot say which one is responsible. No pin-repair fallback was observed, which suggests the stack isolation does the work here, but that is not proven.
4096-token probes only: not a full-context (128K/1M) or general quality qualification, and not production qualification.

Speed

Passing requests only, medians; Mann-Whitney p.

4096-token prompt, ~10-token answers (300 processes, ~450 requests per arm):

	Control	Patched	Change	p
Prefill (tok/s), Q5	892.5	895.1	+0.3%	0.82
Prefill, Q8	743.0	741.8	-0.2%	0.59
Output (tok/s), Q5	25.9	25.9	+0.1%	0.83
Output, Q8	18.7	18.9	+1.2%	0.92

Output rates here come from ~10-token answers and are indicative only.

Dedicated small speed benchmark (speed-v1: 10 pairs = 20 fresh processes, 4096- and 16384-token prompts, 256 generated tokens each, matched prompts, paired by pair):

Prompt	Metric	Paired patched/control (median)	Patched faster in	Sign-test p
4096	prefill	-1.0%	1/10	0.021
4096	output	-2.0%	2/10	0.109
16384	prefill	-0.1%	3/10	0.344
16384	output	-3.3%	3/10	0.344
all	output	mean -1.5% (sd 4.9%)	5/20	0.041

Per model medians (unpaired, n=5 per cell): Q5 output 37.2 -> 36.7 (4K) and 43.7 -> 42.5 (16K) tok/s; Q8 24.0 -> 23.6 and 27.2 -> 27.0; prefill within 1.2% (Q5 953/942 at 4K, 1308/1303 at 16K; Q8 809/804, 1078/1075). Interpretation: any speed cost is small, of the order of 1-3% on output, with borderline significance in the small paired benchmark (p about 0.04) and no significant effect in the 300-process data. Prefill means are skewed by a few large positive outliers (up to +25%), so medians are quoted. 20 pairs cannot exclude a few-percent effect either way; treat "about 0-3% slower decode" as the working estimate, not a measured constant. Pooled "ALL" medians that mix Q5 and Q8 are not meaningful and are not quoted.

Housekeeping and verification

After the final run: no strata/llama/whirl process, nothing on port 8080, executable hash unchanged, all nine production files match production-before.json, both coordination files restored byte-for-byte after every run (main-v1, main-v2, speed-v1), no native cleanup failures, production never started.
Windows events: between the first snapshot and the final one, 11 new Application-log records appeared: 9 Windows Error Reporting re-filings (all at 07:56 AEDT, in one burst) of old reports (0x116, 0x7E and LiveKernelEvent 141/193 reports whose dump files date from 19-29 Oct), and 2 crashpad_log/setup.exe records at 06:58 AEDT from an unrelated installer. No new WHEA, bugcheck or GPU fault; the WER burst is the known ~4-hourly re-filing, not a new fault. (39 old records dropped out of the query window; the query is capped.)
Not done: the authoritative old-vs-new comparison of the event logs beyond the above; any longer-context or production-config test; deployment of the patched binary.

Files (all under `exports/strata-3431-validation-20261031`)

Raw runs: main-v1/, main-v2/, speed-v1/ (per-start folders with native stdout/stderr, prompts, answers, timings, host/admission snapshots).
Summaries: report-v1/, report-v2/ (per-run), combined-v1-v2.json (pooled statistics), speed-v1/paired-and-unpaired-analysis.txt.
Draft comment (unposted): github-comment-FINAL-DRAFT.txt.
Tools (tools/llm-bench): strata_3431_startup_ab_v1.py, strata_3431_speed_v1.py, strata_3431_combine_v1.py, strata_3431_speed_paired_v1.py, summarize_strata_3431_v1.py, strata_3431_repeat_rounds_v1.ps1 (has a state-field bug: it reads nonexistent fields and stops after the first round; not used for the results).

Navigation Menu

Repository navigation

[Engine] Windows HIP (gfx1100): batched PLE block silently misreads prompts in ~1 of 8 engine starts; STRATA_PLE_BATCH=0 fixes it #3431

Design (same in all runs)

System information

Results: bad-start rate and significance

Speed

Housekeeping and verification

Files (all under `exports/strata-3431-validation-20261031`)

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

stack variable	alias range of the bad slice
0x14D38DB0E0	0x14961D0000 .. 0x14DD34F000
0x480ECDAD60	0x47DC880000 .. 0x48239FF000
0x0C8D52AF60	0x0C5AEC0000 .. 0x0C9D53F000

[Engine] Windows HIP (gfx1100): batched PLE block silently misreads prompts in ~1 of 8 engine starts; STRATA_PLE_BATCH=0 fixes it #3431

Description

Activity

ahmetem commented on Oct 22, 2026

giacomobord-prog commented on Oct 22, 2026

ahmetem commented on Oct 22, 2026

giacomobord-prog commented on Oct 22, 2026

giacomobord-prog commented on Oct 23, 2026

ahmetem commented on Oct 23, 2026

sylvesterpious commented on Oct 26, 2026

giacomobord-prog commented on Oct 26, 2026

jerem91150 commented on Oct 29, 2026

GoldenLeafBird commented on Oct 31, 2026

sylvesterpious commented on Oct 31, 2026

Design (same in all runs)

System information

Results: bad-start rate and significance

Speed

Housekeeping and verification

Files (all under exports/strata-3431-validation-20261031)

Metadata

Metadata

Assignees

Labels

Projects

Milestone

Relationships

Development

Issue actions

Files (all under `exports/strata-3431-validation-20261031`)