Merge main (0.2.1.2 Strix Halo engine hotfix)

CMakeLists.txt

@@ -736,6 +736,13 @@ if(STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP)
   strata_utf8_manifest(strata)
   target_include_directories(strata PRIVATE ${_strata_gpu_include_directories})
   target_link_libraries(strata PRIVATE strata_engine strata_prefill strata_spec ${_strata_gpu_runtime_target})
+  if(WIN32 AND STRATA_ENABLE_HIP)
+    # 0.2.1.2: the ROCm math libraries live in <package>\rocm\bin, which is on PATH only when the launcher starts the engine; load them on first
+    # use (main() adds that folder to the DLL search first) so `strata.exe --version` works from the package folder.  amdhip64_7.dll and the
+    # C++ runtime sit beside the exe (tools/hip/package_windows.py) and stay normal imports.
+    target_link_libraries(strata PRIVATE delayimp)
+    target_link_options(strata PRIVATE "LINKER:/DELAYLOAD:hipblas.dll" "LINKER:/DELAYLOAD:libhipblaslt.dll" "LINKER:/DELAYLOAD:rocblas.dll")
+  endif()
 
   # NEW-MODELS-PLAN.md T1: Qwen3.6-35B-A3B's reference path (one GPU, the model resident, plain kernels), checked
   # against llama.cpp by tools/qwen35/; the engine itself does not run the family yet

docs/STRIX_HALO.md

@@ -8,6 +8,8 @@ Ubuntu 26.04, kernel 7.0, ROCm 7.14.1). Setup (`setup.sh`, `START-HERE.bat`) rec
 named as what it is and not supported. On Linux setup compiles the engine for gfx1151 itself (the steps below are the same, by hand). The 0.2.0-beta release also has ready-made packages: `Strata-0.2.0-beta-linux-amd-strixhalo.AppImage` (engine `halo`, gfx1151, ROCm 7.14.1) and, on Windows, `Strata-0.2.0-beta-windows-hip-rdna3.exe` (engine `rdna3`: gfx1100, gfx1101, gfx1102 and gfx1151, ROCm 10.2.0a20260930); any app installer plus **Update engine** (or `strata update`) fetches the same package. What was run: both packages were built, linked and scanned, and the Linux package refuses to start on a card of another family (tried on an RX 9000 card); the maintainers have no Strix Halo and ran neither on one. The measurements on this page come from source builds on Linux (the maintainers' earlier box and community machines). Nobody has run a Strix Halo on Windows. The packaged engines do not carry the hipBLASLt table of section 3 yet (planned for 0.2.1), so their prompt reading is slower than the numbers below until then. A Windows zip of the 0.1.x START-HERE route without gfx1151 code is refused with a message, never used. Setup counts the GPU's memory as the BIOS carve-out plus the shared memory
 (the GTT pool; on Windows half the RAM) and recommends UD-IQ4_XS from about 80 GB of memory; it changes no BIOS or kernel setting.
 
+**Windows start-up (0.2.1.2).** On a Strix Halo under Windows the engine no longer registers the expert arena with the GPU driver: registering it took all free VRAM (91 GiB before, 0 after, on a 64 GB carve-out + 32 GB shared PC), the expert cache got no slots and the engine stopped with "the embedding buffer failed" (0.2.1 and earlier). The arena stays in RAM, held by the working-set lock, and the cache gets its normal slots. `STRATA_ARENA_PIN_GIB=N` (or `0`) registers it as before, `STRATA_UMA_REGISTER=1` too. Linux is unchanged.
+
 What the engine does differently on this chip:
 
 - **Unified memory.** The GPU's "VRAM" is the GTT pool of the same RAM. `device_free_bytes()` (the figure `--expert-cache auto` and

include/strata/core/arch_defaults.hpp

@@ -34,7 +34,8 @@ bool igpu_prompt_set_enabled(const char* gcn_arch, bool windows);
 
 /// 021-minipc: UMA mode (the expert arena stays unregistered).  Exact (no bit changes) but NOT a default: measured on the 780M it was
 /// 1-4% slower and did not stop the resets (4 of 20 runs against 1 of 10), so kIgpuUmaArenaDefaultOn is false.  When true it applies
-/// to the Radeon 780M only (gfx1103, integrated, Linux); gfx1150 / gfx1151 are not covered (no hardware).  STRATA_UMA_REGISTER=1
+/// to the Radeon 780M only (gfx1103, integrated, Linux).  EXCEPT a Strix Halo on Windows (gfx1151, integrated): there it is always on (0.2.1.2,
+/// registering the arena used up all free VRAM and the engine did not start).  STRATA_UMA_REGISTER=1
 /// registers the arena as before; STRATA_ARENA_PIN_GIB=none selects UMA mode per run.
 inline constexpr bool kIgpuUmaArenaDefaultOn = false;
 bool igpu_uma_arena_default(const char* gcn_arch, bool integrated, bool windows, bool on = kIgpuUmaArenaDefaultOn);

src/core/arch_defaults.cpp

@@ -74,9 +74,15 @@ bool igpu_prompt_set_enabled(const char* gcn_arch, bool windows) {
 }
 
 bool igpu_uma_arena_default(const char* gcn_arch, bool integrated, bool windows, bool on) {
-    if (!on || windows || !integrated || !strata::kernels::gfx_arch_is(gcn_arch, "gfx1103")) return false;
+    if (!integrated) return false;
     const char* reg = std::getenv("STRATA_UMA_REGISTER");
-    return !(reg != nullptr && reg[0] == '1');
+    if (reg != nullptr && reg[0] == '1') return false;
+    // 0.2.1.2: a Strix Halo (gfx1151) on Windows.  Registering the 47 GiB arena with the driver took ALL the free VRAM (91 GiB before,
+    // 0 after; BIOS carve-out 64 GB + 32 GB shared): the expert cache got 0 slots and the engine exited ("the embedding buffer
+    // failed"), 3 of 3 starts (notes-022-strix.md).  Unregistered, the arena is held by the working-set lock and the cache gets
+    // 24,576 slots.  Independent of `on` (that is the 780M's opt-in).  STRATA_ARENA_PIN_GIB=N / 0 still registers it.
+    if (windows) return strata::kernels::gfx_arch_is(gcn_arch, "gfx1151");
+    return on && strata::kernels::gfx_arch_is(gcn_arch, "gfx1103");
 }
 
 std::vector<std::pair<std::string, std::string>> igpu_default_env(const char* gcn_arch, bool integrated) {

src/core/arch_defaults_test.cpp

@@ -190,6 +190,17 @@ int main() {
     CHECK(!igpu_uma_arena_default("gfx1103", false, false, true));    // a gfx1103 that is not integrated
     for (const char* g : {"gfx1150", "gfx1151", "gfx1152", "gfx1100", "gfx1101", "gfx1200", "gfx1201", "gfx1030", "gfx11030", "", "gfx90a"})
         CHECK(!igpu_uma_arena_default(g, true, false, true));
+    // 0.2.1.2: Strix Halo on Windows is unregistered by default (on or off), nothing else on Windows; Linux gfx1151 unchanged
+    CHECK(igpu_uma_arena_default("gfx1151", true, true));
+    CHECK(igpu_uma_arena_default("gfx1151:sramecc+:xnack-", true, true));
+    CHECK(!igpu_uma_arena_default("gfx1151", false, true));           // not integrated
+    CHECK(!igpu_uma_arena_default("gfx1151", true, false));           // Linux
+    for (const char* g : {"gfx1150", "gfx1152", "gfx1103", "gfx1100", "gfx1101", "gfx1102", "gfx1200", "gfx1201", "gfx1030", "", "gfx90a"})
+        CHECK(!igpu_uma_arena_default(g, true, true));
+    CHECK(!igpu_uma_arena_default(nullptr, true, true));
+    put("STRATA_UMA_REGISTER", "1");
+    CHECK(!igpu_uma_arena_default("gfx1151", true, true));
+    unset("STRATA_UMA_REGISTER");
     CHECK(!igpu_uma_arena_default(nullptr, true, false, true));
     put("STRATA_UMA_REGISTER", "1");
     CHECK(!igpu_uma_arena_default("gfx1103", true, false, true));     // the user registers the arena as before

src/program/generate.cpp

@@ -2271,6 +2271,23 @@ int main(int argc, char** argv) {
     // 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.
     std::setvbuf(stdout, nullptr, _IONBF, 0);
+#if defined(STRATA_USE_HIP) && defined(_WIN32)
+    {   // 0.2.1.2: hipblas / libhipblaslt / rocblas are delay-loaded (CMakeLists.txt) and live in <exe folder>\rocm\bin, which only the
+        // launcher puts on PATH; make the folder part of this process's DLL search, so `strata.exe --version` run straight from the
+        // package folder works.  The HIP runtime and the C++ runtime sit beside the exe and are found without it.
+        wchar_t exe[MAX_PATH * 2] = {};
+        const DWORD n = GetModuleFileNameW(nullptr, exe, (DWORD) (sizeof(exe) / sizeof(exe[0])));
+        if (n > 0 && n < sizeof(exe) / sizeof(exe[0])) {
+            std::wstring dir(exe, n);
+            const size_t cut = dir.find_last_of(L"\\/");
+            if (cut != std::wstring::npos) {
+                dir.resize(cut);
+                dir += L"\\rocm\\bin";
+                if (GetFileAttributesW(dir.c_str()) != INVALID_FILE_ATTRIBUTES) SetDllDirectoryW(dir.c_str());
+            }
+        }
+    }
+#endif
     // `strata --version` / `-V` (also `strata generate --version`): the version and nothing else, before any GPU is touched
     for (int i = 1; i < argc && i < 3; ++i) {
         if (std::strcmp(argv[i], "--version") == 0 || std::strcmp(argv[i], "-V") == 0) {

tools/hip/package_windows.py

@@ -62,7 +62,12 @@ vcruntime140*.dll) are redistributable under the Visual Studio license (Distribu
 
 def imports(objdump: Path, dll: Path) -> list[str]:
     text = subprocess.run([str(objdump), "-p", str(dll)], capture_output=True, text=True, check=True).stdout
-    return [m.group(1) for m in re.finditer(r"DLL Name:\s*(\S+)", text)]
+    names = [m.group(1) for m in re.finditer(r"DLL Name:\s*(\S+)", text)]
+    # 0.2.1.2: strata.exe delay-loads hipblas / libhipblaslt / rocblas, which objdump -p does not list; dll_closure reads the delay-import table too
+    sys.path.insert(0, str(Path(__file__).resolve().parent))
+    from dll_closure import pe_imports
+    seen = {n.lower() for n in names}
+    return names + [n for n in pe_imports(dll) if n.lower() not in seen]
 
 
 def crt_dir() -> Path | None: