HIP: run Qwen3.8-Flash-Next on Windows gfx1201

ROCm 10 host-compiles the wave32 backend for gfx1100, gfx1200, and gfx1201.
Unsloth's split GGUF, Q8_0 experts, an F32 PLE convolution, and a sticky
hipBLAS error on the RX 9070 are handled so the engine can serve the model.

CMakeLists.txt

@@ -33,7 +33,7 @@ if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
 endif()
 
 option(STRATA_ENABLE_CUDA "Build the CUDA targets (Phase 2+)" OFF)
-option(STRATA_ENABLE_HIP "Build the HIP targets for gfx1100 wave32 hardware" OFF)
+option(STRATA_ENABLE_HIP "Build the HIP targets for gfx1100/gfx1200/gfx1201 wave32 hardware" OFF)
 option(STRATA_PREFILL_MMQ "Build the optional GGML MMQ prefill path for HIP" OFF)
 if(STRATA_PREFILL_MMQ AND NOT STRATA_ENABLE_HIP)
   message(FATAL_ERROR "STRATA_PREFILL_MMQ requires STRATA_ENABLE_HIP=ON")

README.md

@@ -84,9 +84,9 @@ one later with `SETUP.bat` (the same as `START-HERE.bat --setup`; on Linux `./se
 For **OrcaRouter's Flash-Next Uncensored IQ3_XXS**, see the [manual compatibility setup](docs/ORCA.md).
 It needs an explicit packing conversion and is not an installer menu option.
 
-An **AMD Radeon RX 7900 XT / XTX on Linux** works too (experimental): `./setup.sh --backend hip`, chosen by itself on
-a PC with no NVIDIA card Strata can use. It installs ROCm without sudo and compiles the engine (one GPU, no images
-yet). Details: [AMD HIP](docs/AMD_HIP.md).
+An **AMD Radeon RX 7900 XT / XTX, RX 9060 or RX 9070** works too (experimental), on Linux or Windows:
+`./setup.sh --backend hip` or `START-HERE.bat --backend hip`, chosen by itself on a PC with no NVIDIA card Strata
+can use. It compiles the engine against ROCm (one GPU, no images yet). Details: [AMD HIP](docs/AMD_HIP.md).
 
 ## Install
 

cmake/hip_backend.cmake

@@ -1,12 +1,18 @@
-# Opt-in HIP configuration. Strata's CUDA-shaped kernels currently target RDNA3 gfx1100 wave32.
+# Opt-in HIP configuration. Wave32 targets: RDNA3 gfx1100 and RDNA4 gfx1200/gfx1201.
 # CMake/compiler discovery stays machine-independent; pass CMAKE_HIP_COMPILER when it is not on PATH.
+set(_strata_hip_archs gfx1100 gfx1200 gfx1201)
 if(NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "")
   set(CMAKE_HIP_ARCHITECTURES gfx1100 CACHE STRING "Strata HIP target architecture")
 endif()
-if(NOT CMAKE_HIP_ARCHITECTURES STREQUAL "gfx1100")
-  message(FATAL_ERROR
-    "Strata HIP currently supports only gfx1100 wave32; CMAKE_HIP_ARCHITECTURES is '${CMAKE_HIP_ARCHITECTURES}'")
-endif()
+set(_strata_hip_req "${CMAKE_HIP_ARCHITECTURES}")
+string(REPLACE " " ";" _strata_hip_req "${_strata_hip_req}")
+foreach(_a IN LISTS _strata_hip_req)
+  if(NOT _a IN_LIST _strata_hip_archs)
+    message(FATAL_ERROR
+      "Strata HIP supports wave32 gfx1100, gfx1200 and gfx1201; "
+      "CMAKE_HIP_ARCHITECTURES is '${CMAKE_HIP_ARCHITECTURES}'")
+  endif()
+endforeach()
 
 enable_language(HIP)
 find_package(hip CONFIG REQUIRED)
@@ -35,11 +41,21 @@ target_include_directories(strata_hip_runtime BEFORE INTERFACE
   "${STRATA_HIP_COMPAT_INCLUDE_DIR}" "${CMAKE_CURRENT_SOURCE_DIR}/include")
 target_compile_definitions(strata_hip_runtime INTERFACE STRATA_USE_HIP=1)
 target_link_libraries(strata_hip_runtime INTERFACE hip::host)
-foreach(_language IN ITEMS CXX HIP)
+# The shim renames the CUDA runtime to HIP. MSVC host compiles take /FI; HIP and
+# other host compilers take -include. The header path uses forward slashes so
+# /FI does not swallow backslashes.
+file(TO_CMAKE_PATH "${STRATA_HIP_COMPAT_INCLUDE_DIR}/cuda_runtime.h" _strata_hip_force)
+if(MSVC)
   target_compile_options(strata_hip_runtime INTERFACE
-    "$<$<COMPILE_LANGUAGE:${_language}>:-include>"
-    "$<$<COMPILE_LANGUAGE:${_language}>:${STRATA_HIP_COMPAT_INCLUDE_DIR}/cuda_runtime.h>")
-endforeach()
+    "$<$<COMPILE_LANGUAGE:CXX>:/FI${_strata_hip_force}>")
+else()
+  target_compile_options(strata_hip_runtime INTERFACE
+    "$<$<COMPILE_LANGUAGE:CXX>:-include>"
+    "$<$<COMPILE_LANGUAGE:CXX>:${_strata_hip_force}>")
+endif()
+target_compile_options(strata_hip_runtime INTERFACE
+  "$<$<COMPILE_LANGUAGE:HIP>:-include>"
+  "$<$<COMPILE_LANGUAGE:HIP>:${_strata_hip_force}>")
 
 # CMake does not infer HIP from Strata's existing CUDA-shaped .cu suffixes.
 file(GLOB_RECURSE _strata_hip_sources CONFIGURE_DEPENDS

docs/AMD_HIP.md

@@ -1,8 +1,9 @@
-# Experimental AMD HIP backend (gfx1100)
+# Experimental AMD HIP backend (gfx1100, gfx1200, gfx1201)
 
-This is a manual Linux source build for the RX 7900 XTX. It is opt-in; the
-NVIDIA installer and CUDA build remain the default. Other AMD architectures,
-wave64, Windows HIP, and mixed AMD/NVIDIA execution are outside this contribution.
+This is an opt-in source build. The NVIDIA installer and CUDA build remain the
+default. Wave32 targets are the RX 7900 XT/XTX (`gfx1100`), the RX 9060
+(`gfx1200`) and the RX 9070 (`gfx1201`). Wave64 and mixed AMD/NVIDIA execution
+are outside this contribution.
 
 The backend maps the CUDA-shaped runtime and BLAS calls to HIP/hipBLAS, uses
 RDNA3's signed integer dot instruction for quantized kernels, and supplies
@@ -14,19 +15,25 @@ This does not claim bit-identical model answers across backends. See
 
 ## Install with setup (recommended)
 
-On Linux with an RX 7900 XT / XTX and the kernel's amdgpu driver (no ROCm install needed):
+On Linux, or on Windows with the HIP SDK or the ROCm 10 Python SDK already installed:
 
 ```sh
-./setup.sh --backend hip
+./setup.sh --backend hip          # Linux
+START-HERE.bat --backend hip      # Windows
 ```
 
-- **Detection:** setup finds the card through the kernel's KFD topology. Integrated Radeon GPUs are listed as not
-  supported. On a PC without an NVIDIA card Strata can use, `--backend hip` is chosen automatically.
-- **ROCm:** installed into `.venv` from AMD's TheRock wheels (~10 GB, no sudo), pinned to the version this backend was
-  tested with (`STRATA_ROCM_VERSION` / `STRATA_ROCM_INDEX` override it). A system ROCm in `/opt/rocm` (or
-  `$ROCM_PATH`) with hipcc and hipBLAS is used instead when present.
-- **Engine:** compiled on your PC (10-20 minutes, once; again after a `git pull` that changes it). This needs a C++
-  compiler and git (`sudo apt install build-essential git`).
+Windows finds the GPU from the display driver and compiles with the ROCm that
+has this card's bitcode. An installed `rocm-sdk` (ROCm 10, `rocm[libraries,devel,device-gfx1201]`)
+is used before the HIP SDK under `C:\Program Files\AMD\ROCm`. Linux still
+detects the card through the kernel's KFD topology.
+
+- **Detection:** Linux uses the kernel's KFD topology. Windows uses the display driver. Integrated Radeon GPUs are
+  listed as not supported. On a PC without an NVIDIA card Strata can use, `--backend hip` is chosen automatically.
+- **ROCm:** an installed SDK that contains this GPU's bitcode is used (`rocm-sdk`, `ROCM_PATH`, `HIP_PATH`, or
+  `C:\Program Files\AMD\ROCm` on Windows; `/opt/rocm` on Linux). Otherwise setup installs TheRock wheels into `.venv`.
+  On Linux the pin is `STRATA_ROCM_VERSION` / `STRATA_ROCM_INDEX`. On Windows the index is `STRATA_ROCM_INDEX_WIN`.
+- **Engine:** compiled on your PC (10-20 minutes, once; again after a `git pull` that changes it). Linux needs a C++
+  compiler and git. Windows needs Visual Studio Build Tools (C++) and git.
 - **Limits for now:** one GPU, no images, no calibration. The Monitor shows no GPU statistics.
 
 The rest of setup is the same as on NVIDIA: the model download, the start script, the server.
@@ -42,7 +49,7 @@ table. VRAM occupancy alone is not a throughput measurement.
 cmake -S . -B build-hip \
   -DCMAKE_BUILD_TYPE=Release \
   -DSTRATA_ENABLE_HIP=ON -DSTRATA_ENABLE_CUDA=OFF \
-  -DCMAKE_HIP_ARCHITECTURES=gfx1100
+  -DCMAKE_HIP_ARCHITECTURES=gfx1201
 cmake --build build-hip --target strata -j2
 ```
 

include/strata/core/native_head.hpp

@@ -53,6 +53,7 @@ public:
 
 private:
     void* host_ = nullptr;
+    void* dev_owned_ = nullptr;
     const void* dev_ = nullptr;
     uint64_t bytes_ = 0;
     size_t row_ = 0;

include/strata/hip_compat/cuda_runtime.h

@@ -1,5 +1,13 @@
 #pragma once
 // Included only by STRATA_ENABLE_HIP builds. CUDA builds use NVIDIA headers.
+#if defined(_WIN32)
+#ifndef WIN32_LEAN_AND_MEAN
+#define WIN32_LEAN_AND_MEAN
+#endif
+#ifndef NOMINMAX
+#define NOMINMAX
+#endif
+#endif
 #include <hip/hip_runtime.h>
 // Do not let HIP's legacy macro corrupt libstdc++ attribute names.
 #ifdef __noinline__

include/strata/hip_compat/intrinsics.hpp

@@ -15,10 +15,10 @@ __device__ __forceinline__ int signed_byte(uint32_t word, int lane) {
 
 // CUDA's signed __dp4a: four signed byte products accumulated modulo 2^32.
 __device__ __forceinline__ int dp4a(int a, int b, int c) {
-#if defined(__gfx1100__) && __has_builtin(__builtin_amdgcn_sudot4)
-    // RDNA3 exposes the signed/unsigned dot4 form. Mark both packed operands
-    // signed to preserve CUDA __dp4a semantics; keep the portable path for
-    // other HIP compilers/targets.
+#if (defined(__gfx1100__) || defined(__gfx1200__) || defined(__gfx1201__)) && __has_builtin(__builtin_amdgcn_sudot4)
+    // RDNA3 and RDNA4 expose the signed/unsigned dot4 form. Mark both packed
+    // operands signed to preserve CUDA __dp4a semantics; keep the portable
+    // path for other HIP compilers/targets.
     return __builtin_amdgcn_sudot4(true, a, true, b, c, false);
 #else
     const uint32_t ua = static_cast<uint32_t>(a);

include/strata/kernels/cpu/expert_layout.hpp

@@ -25,7 +25,7 @@ struct ExpertLayout {
     std::vector<uint64_t> gguf_off;
     /// Per layer, the GGUF file (a name beside the --native shard) that holds its experts when the model's
     /// shards split the layers (Swift's GGUFs: layers 13-47 in shard 2).  Empty = the --native shard itself.
-    std::vector<std::string> gguf_file;
+    std::vector<std::string> gguf_file;   ///< empty, or 3 names per layer (gate, up, down) beside --native
     uint64_t max_blob = BLOB;
     uint64_t total = 0;                   ///< experts.bin size
 

setup.py

@@ -706,16 +706,28 @@ def cuda_lib_dirs():
 # has hipcc and hipBLAS) and the engine is compiled here.  One GPU, no images yet.
 ROCM_INDEX = os.environ.get("STRATA_ROCM_INDEX", "https://rocm.nightlies.amd.com/v2/gfx110X-dgpu/")
 ROCM_VERSION = os.environ.get("STRATA_ROCM_VERSION", "7.10.0a20251120")   # what Strata's HIP build was tested with
-AMD_ARCHS = ("gfx1100",)
-AMD_NAMES = {"gfx1100": "AMD Radeon RX 7900 series (gfx1100)"}   # when sysfs has no product name
+ROCM_INDEX_WIN = os.environ.get("STRATA_ROCM_INDEX_WIN", "https://nightly.repo.amd.com/rocm/whl-next/")
+AMD_ARCHS = ("gfx1100", "gfx1200", "gfx1201")
+AMD_NAMES = {"gfx1100": "AMD Radeon RX 7900 series (gfx1100)",
+             "gfx1200": "AMD Radeon RX 9060 series (gfx1200)",
+             "gfx1201": "AMD Radeon RX 9070 series (gfx1201)"}
+# Windows display-class device ids. Name matching covers the same cards when the id is new.
+_WIN_AMD_DID = {0x744C: "gfx1100", 0x7448: "gfx1100", 0x745E: "gfx1100",
+                0x7590: "gfx1200",
+                0x7550: "gfx1201", 0x7551: "gfx1201"}
+_WIN_AMD_NAME = ((re.compile(r"9070|R9700|R9600", re.I), "gfx1201"),
+                 (re.compile(r"9060", re.I), "gfx1200"),
+                 (re.compile(r"7900|W7900|W7800", re.I), "gfx1100"))
 
 
 def amd_gpus():
     """AMD GPUs from the kernel's KFD topology (the amdgpu driver; no ROCm needed), numbered as HIP numbers them:
     the GPU nodes in order, the CPU nodes skipped.  Integrated GPUs are listed too (not supported)."""
+    if WIN:
+        return _amd_gpus_windows()
     base = Path("/sys/class/kfd/kfd/topology/nodes")
     found = []
-    if WIN or not base.is_dir():
+    if not base.is_dir():
         return found
     for node in sorted((p for p in base.iterdir() if p.name.isdigit()), key=lambda p: int(p.name)):
         try:
@@ -747,16 +759,182 @@ def amd_gpus():
 
 def amd_problem(g):
     if g["arch"] not in AMD_ARCHS:
-        return (f"not supported - Strata's AMD backend runs on the RX 7900 XT / XTX ({', '.join(AMD_ARCHS)}) only, "
-                f"this is {g['arch']}")
+        return (f"not supported - Strata's AMD backend runs on the RX 7900 XT/XTX (gfx1100), "
+                f"RX 9060 (gfx1200) and RX 9070 (gfx1201), this is {g['arch']}")
     return None
 
 
-def rocm_root():
-    """ROCm for compiling and running the HIP engine: (root, library folders).  A system ROCm with hipcc and hipBLAS,
-    else AMD's TheRock wheels (ROCM_VERSION, from ROCM_INDEX) installed into .venv."""
+def _amd_gpus_windows():
+    import winreg
+    found = []
+    try:
+        reg = winreg.OpenKey(winreg.HKEY_LOCAL_MACHINE,
+                             r"SYSTEM\CurrentControlSet\Control\Class\{4d36e968-e325-11ce-bfc1-08002be10318}")
+    except OSError:
+        return found
+    i = 0
+    while True:
+        try:
+            sub = winreg.EnumKey(reg, i)
+        except OSError:
+            break
+        i += 1
+        if not sub.isdigit():
+            continue
+        try:
+            key = winreg.OpenKey(reg, sub)
+        except OSError:
+            continue
+
+        def _get(name, key=key):
+            try:
+                return winreg.QueryValueEx(key, name)[0]
+            except OSError:
+                return None
+
+        name = str(_get("DriverDesc") or "")
+        mid = str(_get("MatchingDeviceId") or "")
+        mem = _get("HardwareInformation.qwMemorySize")
+        did = re.search(r"VEN_1002&DEV_([0-9A-Fa-f]{4})", mid)
+        if not did:
+            continue
+        arch = _WIN_AMD_DID.get(int(did.group(1), 16))
+        if arch is None:
+            arch = next((a for rx, a in _WIN_AMD_NAME if rx.search(name)), None)
+        if arch is None:
+            continue
+        vram = int(mem) / 2 ** 30 if isinstance(mem, int) and mem > 0 else 0.0
+        if name in ("", f"AMD Radeon ({arch})") and arch in AMD_NAMES:
+            name = AMD_NAMES[arch]
+        found.append({"index": len(found), "name": name or AMD_NAMES.get(arch, arch), "vram_gb": vram,
+                      "arch": arch, "driver": "hip", "vendor": "amd"})
+    return found
+
+
+def _rocm_has_hipcc(root: Path) -> bool:
+    return any((root / "bin" / name).exists() for name in ("hipcc", "hipcc.exe", "hipcc.bat"))
+
+
+def _rocm_has_hipblas(root: Path) -> bool:
+    lib = root / "lib"
+    if not lib.is_dir():
+        return False
+    return any(lib.glob(pat) for pat in ("libhipblas.so*", "hipblas.lib", "libhipblas.dll.a"))
+
+
+def _rocm_compiler(root: Path) -> Path | None:
+    for rel in ("llvm/bin/clang++", "llvm/bin/clang++.exe", "lib/llvm/bin/clang++", "lib/llvm/bin/clang++.exe",
+                "lib/llvm/bin/amdclang++.exe", "bin/amdclang++.exe"):
+        p = root / rel
+        if p.exists():
+            return p
+    return None
+
+
+def _rocm_bitcode(root: Path) -> Path | None:
+    for rel in ("lib/llvm/amdgcn/bitcode", "amdgcn/bitcode", "lib/amdgcn/bitcode"):
+        p = root / rel
+        if (p / "oclc_isa_version_1100.bc").exists() or (p / "oclc_isa_version_1201.bc").exists():
+            return p
+    return None
+
+
+def _rocm_usable(root: Path, arch: str) -> bool:
+    bitcode = _rocm_bitcode(root)
+    return bool(root and root.is_dir() and _rocm_has_hipcc(root) and _rocm_has_hipblas(root)
+                and _rocm_compiler(root) and bitcode and (bitcode / f"oclc_isa_version_{arch[3:]}.bc").exists())
+
+
+def _find_rocm_sdk() -> str | None:
+    hit = shutil.which("rocm-sdk")
+    if hit:
+        return hit
+    names = ["rocm-sdk.exe", "rocm-sdk"] if WIN else ["rocm-sdk"]
+    cands = [Path(sys.executable).parent / name for name in names]
+    if WIN:
+        roaming = os.environ.get("APPDATA")
+        if roaming:
+            cands += sorted(Path(roaming).glob("Python/Python*/Scripts/rocm-sdk.exe"))
+    for c in cands:
+        if c.exists():
+            return str(c)
+    return None
+
+
+def _rocm_sdk_root() -> Path | None:
+    sdk = _find_rocm_sdk()
+    if not sdk:
+        return None
+    try:
+        return Path(out([sdk, "path", "--root"]).strip())
+    except (OSError, subprocess.CalledProcessError):
+        return None
+
+
+def _windows_rocm_candidates() -> list[Path]:
+    """ROCM_PATH wins when it is set. Otherwise the ROCm 10 `rocm-sdk` tree comes before the older HIP SDK."""
+    cands = []
+    if os.environ.get("ROCM_PATH"):
+        cands.append(Path(os.environ["ROCM_PATH"]))
+    sdk = _rocm_sdk_root()
+    if sdk:
+        cands.append(sdk)
+    if os.environ.get("HIP_PATH"):
+        cands.append(Path(os.environ["HIP_PATH"]))
+    program = Path(os.environ.get("ProgramFiles", r"C:\Program Files")) / "AMD" / "ROCm"
+    if program.is_dir():
+        cands += sorted((p for p in program.iterdir() if p.is_dir()), reverse=True)
+        cands.append(program)
+    return cands
+
+
+def _rocm_lib_dirs(root: Path) -> list[str]:
+    dirs = []
+    for rel in ("bin", "lib", "lib/llvm/bin", "llvm/bin"):
+        p = root / rel
+        if p.is_dir():
+            dirs.append(str(p))
+    for sp in {Path(p) for p in sys.path if p.endswith("site-packages")}:
+        dirs += [str(d / "lib") for d in sorted(sp.glob("_rocm_sdk_libraries_*")) if (d / "lib").is_dir()]
+        devel = sp / "_rocm_sdk_devel"
+        for rel in ("bin", "lib"):
+            if (devel / rel).is_dir() and str(devel / rel) not in dirs:
+                dirs.append(str(devel / rel))
+    return list(dict.fromkeys(dirs))
+
+
+def rocm_root(arch: str = "gfx1100"):
+    """ROCm for compiling and running the HIP engine: (root, library folders).
+    Windows prefers an installed ROCm 10 SDK or HIP SDK that has this GPU's bitcode.
+    Linux prefers /opt/rocm (or ROCM_PATH), else TheRock wheels pinned to ROCM_VERSION."""
+    if WIN:
+        for root in _windows_rocm_candidates():
+            if _rocm_usable(root, arch):
+                ok(f"ROCm: {root}")
+                return root, _rocm_lib_dirs(root)
+        stamp = Path(sys.prefix) / ".strata-rocm.json"
+        want = {"index": ROCM_INDEX_WIN, "arch": arch, "version": os.environ.get("STRATA_ROCM_VERSION", "")}
+        if not stamp.exists() or json.loads(stamp.read_text()).get("arch") != arch or \
+                json.loads(stamp.read_text()).get("index") != ROCM_INDEX_WIN:
+            spec = f"rocm[libraries,devel,device-{arch}]"
+            if want["version"]:
+                spec += f"=={want['version']}"
+            say(f"  Installing ROCm for {arch} into .venv (AMD's TheRock wheels) ...")
+            run([sys.executable, "-m", "pip", "install", "--disable-pip-version-check", "--index-url",
+                 ROCM_INDEX_WIN, spec])
+            sdk = Path(sys.executable).parent / "rocm-sdk.exe"
+            if sdk.exists():
+                run([str(sdk), "init"])
+            stamp.write_text(json.dumps(want))
+        root = _rocm_sdk_root()
+        if root is None or not _rocm_usable(root, arch):
+            fail(f"ROCm for {arch} is not usable",
+                 "install the HIP SDK or `pip install rocm[libraries,devel,device-" + arch + "]` "
+                 "from https://nightly.repo.amd.com/rocm/whl-next/ and run this again")
+        ok(f"ROCm: {root}")
+        return root, _rocm_lib_dirs(root)
     sysroot = Path(os.environ.get("ROCM_PATH") or "/opt/rocm")
-    if (sysroot / "bin" / "hipcc").exists() and list((sysroot / "lib").glob("libhipblas.so*")):
+    if _rocm_usable(sysroot, arch) or (_rocm_has_hipcc(sysroot) and _rocm_has_hipblas(sysroot) and arch == "gfx1100"):
         return sysroot, [str(sysroot / "lib")]
     stamp = Path(sys.prefix) / ".strata-rocm.json"
     if not stamp.exists() or json.loads(stamp.read_text()).get("version") != ROCM_VERSION:
@@ -766,14 +944,11 @@ def rocm_root():
         stamp.write_text(json.dumps({"version": ROCM_VERSION, "index": ROCM_INDEX}))
     sdk = Path(sys.executable).parent / "rocm-sdk"
     root = Path(out([str(sdk), "path", "--root"]).strip())
-    if not (root / "llvm" / "bin" / "clang++").exists():
+    if _rocm_compiler(root) is None:
         fail(f"ROCm was installed but its compiler is missing ({root})",
-             f"remove {stamp} and run this again; or install ROCm 7 system-wide")
-    dirs = [str(root / "lib")]
-    for sp in {Path(p) for p in sys.path if p.endswith("site-packages")}:
-        dirs += [str(d / "lib") for d in sorted(sp.glob("_rocm_sdk_libraries_*")) if (d / "lib").is_dir()]
+             f"remove {stamp} and run this again; or install ROCm system-wide")
     ok(f"ROCm: {root}")
-    return root, list(dict.fromkeys(dirs))
+    return root, _rocm_lib_dirs(root)
 
 
 def build_engine_hip(gpu, llama) -> Path:
@@ -783,31 +958,59 @@ def build_engine_hip(gpu, llama) -> Path:
     stamp = eng / "BUILD.json"
     meta = json.loads(stamp.read_text()) if stamp.exists() else {}
     src = source_hash(ENGINE_SOURCES)
-    if meta.get("backend") == "hip" and (eng / EXE).exists() and meta.get("src") == src:
+    if meta.get("backend") == "hip" and (eng / EXE).exists() and meta.get("src") == src \
+            and gpu["arch"] in (meta.get("archs") or []):
         ok("engine already built for this PC")
         return eng
-    if not (shutil.which("c++") or shutil.which("g++")) or not shutil.which("git"):
+    if WIN:
+        if find_vcvars() is None or not shutil.which("git"):
+            fail("Visual Studio Build Tools (C++), CMake and git are needed to compile the AMD engine",
+                 "install VS 2019 or 2022 Build Tools with the C++ workload, and Git for Windows")
+    elif not (shutil.which("c++") or shutil.which("g++")) or not shutil.which("git"):
         fail("a C++ compiler and git are needed to compile the AMD engine",
              "Ubuntu/Debian: sudo apt install build-essential git   Fedora: sudo dnf install gcc-c++ git")
-    root, dirs = rocm_root()
-    libs = [str(Path(d).parent) for d in dirs[1:]]
-    bitcode = next((p for p in (root / "lib" / "llvm" / "amdgcn" / "bitcode", root / "amdgcn" / "bitcode") if p.is_dir()),
-                   root / "amdgcn" / "bitcode")
+    root, dirs = rocm_root(gpu["arch"])
+    compiler = _rocm_compiler(root)
+    bitcode = _rocm_bitcode(root)
+    if compiler is None or bitcode is None:
+        fail(f"ROCm at {root} has no HIP compiler or device bitcode for {gpu['arch']}")
+    prefix = [root]
+    for extra in (root / "lib" / "cmake",):
+        if extra.is_dir():
+            prefix.append(extra)
     os.environ.update({"HIP_PLATFORM": "amd", "HIP_COMPILER": "clang", "HIP_RUNTIME": "rocclr", "ROCM_PATH": str(root),
                        "HIP_PATH": str(root)})
-    os.environ["LD_LIBRARY_PATH"] = os.pathsep.join(dirs + [os.environ.get("LD_LIBRARY_PATH", "")]).rstrip(os.pathsep)
-    os.environ["PATH"] = os.pathsep.join([str(root / "bin"), str(root / "llvm" / "bin"), os.environ.get("PATH", "")])
+    if not WIN:
+        os.environ["LD_LIBRARY_PATH"] = os.pathsep.join(dirs + [os.environ.get("LD_LIBRARY_PATH", "")]).rstrip(os.pathsep)
+    os.environ["PATH"] = os.pathsep.join([str(root / "bin"), str(compiler.parent), *dirs,
+                                          os.environ.get("PATH", "")])
     say("  The engine's source changed: compiling it again (only what changed, a few minutes) ..."
         if meta.get("backend") == "hip" and (eng / EXE).exists()
         else "  Compiling the Strata engine for your AMD GPU (10-20 minutes, once) ...")
-    cmake_build(ROOT, ROOT / "build-hip", "strata",
-                ["-DSTRATA_ENABLE_HIP=ON", "-DSTRATA_ENABLE_CUDA=OFF", "-DSTRATA_BUILD_TESTS=OFF",
-                 "-DSTRATA_PREFILL_MMQ=ON", f"-DCMAKE_HIP_ARCHITECTURES={gpu['arch']}",
-                 f"-DCMAKE_HIP_COMPILER={root / 'llvm' / 'bin' / 'clang++'}", f"-DCMAKE_HIP_COMPILER_ROCM_ROOT={root}",
-                 "-DCMAKE_PREFIX_PATH=" + ";".join([str(root), *libs]),
-                 f"-DCMAKE_HIP_FLAGS=--rocm-path={root} --rocm-device-lib-path={bitcode}",
-                 f"-DSTRATA_GGML_DIR={llama}"], None, "")
-    shutil.copy2(ROOT / "build-hip" / EXE, eng / EXE)
+    hip_flags = f"--rocm-path={root.as_posix()} --rocm-device-lib-path={bitcode.as_posix()}"
+    # CMake parses backslashes in these cache values as escapes (\U in \Users).
+    defs = ["-DSTRATA_ENABLE_HIP=ON", "-DSTRATA_ENABLE_CUDA=OFF", "-DSTRATA_BUILD_TESTS=OFF",
+            "-DSTRATA_PREFILL_MMQ=ON", f"-DCMAKE_HIP_ARCHITECTURES={gpu['arch']}",
+            f"-DCMAKE_HIP_COMPILER={compiler.as_posix()}",
+            f"-DCMAKE_HIP_COMPILER_ROCM_ROOT={root.as_posix()}",
+            "-DCMAKE_PREFIX_PATH=" + ";".join(p.as_posix() for p in prefix),
+            f"-DCMAKE_HIP_FLAGS={hip_flags}",
+            f"-DSTRATA_GGML_DIR={Path(llama).as_posix()}"]
+    if WIN:
+        # CMake refuses to mix cl.exe for C/C++ with Clang for HIP. ROCm's clang++
+        # already targets x86_64-pc-windows-msvc, so it compiles the host code too.
+        host_cc = compiler.with_name("clang.exe")
+        defs += [f"-DCMAKE_C_COMPILER={host_cc.as_posix()}", f"-DCMAKE_CXX_COMPILER={compiler.as_posix()}"]
+    cmake_build(ROOT, ROOT / "build-hip", "strata", defs, find_vcvars() if WIN else None,
+                "build-hip.bat" if WIN else "")
+    built = ROOT / "build-hip" / EXE
+    if not built.exists():
+        # Ninja multi-config and some Windows generators put the exe next to a config folder.
+        hits = list((ROOT / "build-hip").rglob(EXE))
+        if not hits:
+            fail("the HIP build finished without producing " + EXE)
+        built = hits[0]
+    shutil.copy2(built, eng / EXE)
     stamp.write_text(json.dumps({"source": "local-hip", "backend": "hip", "version": source_version(),
                                  "archs": [gpu["arch"]], "vision": "none", "lib_dirs": dirs, "src": src}, indent=1))
     ok(f"engine compiled: {eng / EXE}")
@@ -1627,8 +1830,8 @@ def main() -> int:
                     help="map the model's experts from its folder instead of copying them into RAM (for a PC with a big "
                          "GPU and little RAM); auto: when the experts would not fit the RAM")
     ap.add_argument("--backend", choices=["cuda", "hip"],
-                    help="cuda = NVIDIA (default), hip = AMD RX 7900 XT/XTX on Linux (experimental; chosen by itself "
-                         "when the PC has no NVIDIA card Strata can use)")
+                    help="cuda = NVIDIA (default), hip = AMD RX 7900 / 9060 / 9070 on Linux or Windows "
+                         "(experimental; chosen by itself when the PC has no NVIDIA card Strata can use)")
     ap.add_argument("--skip-build", action="store_true", help=argparse.SUPPRESS)
     a = ap.parse_args()
     if a.gpu is not None:                              # --gpu 0,2 means --gpus 0,2 (a user tried it: issue report)
@@ -1702,7 +1905,7 @@ def main() -> int:
     # ---- 1. the PC
     step(1, "checking your PC")
     found = gpus()
-    amd = [] if WIN else amd_gpus()
+    amd = amd_gpus()
     nv_ok = any(gpu_problem(g) is None for g in found)
     amd_ok = [g for g in amd if amd_problem(g) is None]
     hip = a.backend == "hip" or (a.backend is None and not nv_ok and bool(amd_ok))
@@ -1718,14 +1921,13 @@ def main() -> int:
         if a.check and not hip:
             say("  (the AMD card: ./setup.sh --backend hip)")
     if hip:                                            # AMD (experimental): one card, compiled here
-        if WIN:
-            fail("Strata's AMD backend runs on Linux only", "use an NVIDIA RTX 20 series or newer card on Windows")
-        say("  Your AMD GPUs:" if amd else "  No AMD GPU found (the amdgpu driver's KFD topology is empty).")
+        say("  Your AMD GPUs:" if amd else "  No AMD GPU found.")
         for g in amd:
             say(f"    GPU {g['index']}: {g['name']}, {g['vram_gb']:.0f} GB VRAM - " + (amd_problem(g) or "can be used"))
         usable = [g for g in amd if amd_problem(g) is None]
         if not usable:
-            fail("no AMD GPU Strata can use", "the AMD backend runs on the RX 7900 XT / XTX (gfx1100) on Linux")
+            fail("no AMD GPU Strata can use",
+                 "the AMD backend runs on the RX 7900 XT/XTX (gfx1100), RX 9060 (gfx1200) and RX 9070 (gfx1201)")
         if a.gpus:
             fail("several GPUs sharing one model: NVIDIA only for now", "use one AMD card (--gpu N)")
         if a.gpu is not None:
@@ -1742,7 +1944,7 @@ def main() -> int:
         if not found:
             fail("no NVIDIA GPU found (nvidia-smi did not answer)",
                  "install the NVIDIA driver from https://www.nvidia.com/drivers and restart the PC"
-                 + ("; an AMD RX 7900 XT/XTX: --backend hip" if amd else ""))
+                 + ("; an AMD RX 7900 / 9060 / 9070: --backend hip" if amd else ""))
         if len(found) > 1 or gpu_problem(found[0]) is not None:
             gpu_table(found)
         sel = choose_gpus(a, found)                    # asked when two or more cards can share the model

src/core/device.cu

@@ -31,7 +31,7 @@ DeviceInfo device_info(int ordinal) {
     check(cudaGetDeviceCount(&count), "cudaGetDeviceCount");
     if (count == 0) {
 #if defined(STRATA_USE_HIP)
-        throw CudaError("no HIP device is present; this backend targets gfx1100 wave32", -1);
+        throw CudaError("no HIP device is present; this backend targets gfx1100, gfx1200 or gfx1201 (wave32)", -1);
 #else
         throw CudaError("no CUDA device is present; Strata needs an NVIDIA GPU (RTX 20 series or newer)", -1);
 #endif
@@ -64,10 +64,22 @@ DeviceInfo device_info(int ordinal) {
     // fp32-FMA fallback below sm_80, the tensor-core prompt kernels refuse and fall back).  Compiling for a
     // supported arch is enforced by CMake; RUNNING on an older card is caught here, because a binary can be carried
     // to a machine with an older card and would otherwise silently take whatever path the driver chose.  The HIP
-    // backend is validated on gfx1100 (wave32) only.
+    // backend accepts wave32 gfx1100 (RDNA3) and gfx1200/gfx1201 (RDNA4). gcnArchName may carry a
+    // ":sramecc-:xnack-" suffix.
 #if defined(STRATA_USE_HIP)
-    if (std::strncmp(p.gcnArchName, "gfx1100", 7) != 0 || p.warpSize != 32) {
-        throw CudaError("HIP backend requires validated gfx1100 wave32 hardware", -1);
+    auto hip_arch_ok = [](const char* name) {
+        if (name == nullptr) return false;
+        static const char* ok[] = {"gfx1100", "gfx1200", "gfx1201"};
+        for (const char* arch : ok) {
+            const std::size_t n = std::strlen(arch);
+            if (std::strncmp(name, arch, n) == 0 && (name[n] == '\0' || name[n] == ':')) return true;
+        }
+        return false;
+    };
+    if (!hip_arch_ok(p.gcnArchName) || p.warpSize != 32) {
+        throw CudaError(std::string("HIP backend requires wave32 gfx1100, gfx1200 or gfx1201; this device is ") +
+                            p.gcnArchName,
+                        -1);
     }
 #else
     if (d.cc_major * 10 + d.cc_minor < 75) {

src/core/device_main.cpp

@@ -34,7 +34,7 @@ int main(int argc, char** argv) {
         const strata::core::DeviceInfo d = strata::core::device_info(0);
         std::printf("device %d: %s\n", d.ordinal, d.name.c_str());
 #if defined(STRATA_USE_HIP)
-        std::printf("  HIP target          gfx1100 wave32\n");
+        std::printf("  HIP target          wave32 (gfx1100 / gfx1200 / gfx1201)\n");
 #else
         std::printf("  compute capability  %d.%d   (sm_%d%d)\n", d.cc_major, d.cc_minor, d.cc_major, d.cc_minor);
 #endif

src/core/expert_source.cpp

@@ -1099,9 +1099,12 @@ LoadStats load_experts_gguf(const std::string& gguf, uint8_t* dst, const strata:
     // a layer's experts may sit in another shard of the model (native_experts.txt v3): a name beside `gguf`
     const size_t cut = gguf.find_last_of("/\\");
     const std::string dir = cut == std::string::npos ? std::string() : gguf.substr(0, cut + 1);
-    auto file_of = [&](int64_t l) -> std::string {
-        if (lay.gguf_file.empty() || lay.gguf_file[(size_t) l].empty()) return gguf;
-        return dir + lay.gguf_file[(size_t) l];
+    auto file_of = [&](int64_t l, int role) -> std::string {
+        if (lay.gguf_file.size() == (size_t) lay.n_layers * 3) {
+            const std::string& name = lay.gguf_file[(size_t) (3 * l + role)];
+            if (!name.empty()) return dir + name;
+        }
+        return gguf;
     };
     auto worker = [&]() {
         std::ifstream f;
@@ -1110,19 +1113,19 @@ LoadStats load_experts_gguf(const std::string& gguf, uint8_t* dst, const strata:
         for (;;) {
             const int64_t l = next.fetch_add(1);
             if (l >= lay.n_layers || bad) break;
-            const std::string name = file_of(l);
-            if (name != open_name) {
-                f.close();
-                f.clear();
-                f.open(name, std::ios::binary);
-                if (!f) { bad = true; return; }
-                open_name = name;
-            }
             const auto& fm = lay.fmt[(size_t) l];
             const uint64_t blob = lay.bytes[(size_t) l];
             const uint64_t per[3] = {fm.up_off, fm.up_off, blob - fm.down_off};
             const uint64_t at[3] = {0, fm.up_off, fm.down_off};
             for (int r = 0; r < 3; ++r) {
+                const std::string name = file_of(l, r);
+                if (name != open_name) {
+                    f.close();
+                    f.clear();
+                    f.open(name, std::ios::binary);
+                    if (!f) { bad = true; return; }
+                    open_name = name;
+                }
                 const uint64_t src = lay.gguf_off[(size_t) (3 * l + r)];
                 const uint64_t total = per[r] * (uint64_t) lay.n_expert;
                 const uint64_t chunk = per[r] * 16;           // 16 experts per read

src/core/native_head.cpp

@@ -24,6 +24,9 @@ bool NativeHead::load(const std::string& path, int64_t n_in, int64_t n_out, std:
     try {
         strata::GgufFile gguf(path);
         err = strata::check_architecture(gguf);
+        // Unsloth-style splits keep general.architecture on shard 1 only. A later shard that
+        // still carries output.weight is the same model; a present but wrong architecture still fails.
+        if (err == "missing general.architecture") err.clear();
         if (!err.empty()) return false;
         const strata::TensorInfo* tensor = nullptr;
         for (const auto& candidate : gguf.tensors()) {
@@ -101,6 +104,7 @@ void set_native_embed(const NativeEmbed* e) { g_embed = e; }
 const NativeEmbed* native_embed() { return g_embed; }
 
 NativeEmbed::~NativeEmbed() {
+    if (dev_owned_) cudaFree(dev_owned_);
     if (host_) cudaFreeHost(host_);
 }
 
@@ -111,7 +115,7 @@ bool NativeEmbed::load(const std::string& path, int64_t n_embd, int64_t n_vocab,
         for (const auto& c : gguf.tensors())
             if (c.name == "token_embd.weight") t = &c;
         if (!t || t->shape.size() != 2 || t->shape[0] != (uint64_t) n_embd || t->shape[1] != (uint64_t) n_vocab ||
-            !strata::kernels::iq_supported((int) t->type) || n_embd % 256) {
+            !(strata::kernels::iq_supported((int) t->type) || t->type == 8) || n_embd % 256) {
             err = "native embedding: token_embd.weight is absent, of another shape, or of a type without a GPU "
                   "dequantizer";
             return false;
@@ -124,12 +128,18 @@ bool NativeEmbed::load(const std::string& path, int64_t n_embd, int64_t n_vocab,
             return false;
         }
         std::memcpy(host_, gguf.tensor_data(*t), bytes_);
-        void* d = nullptr;
-        if (cudaHostGetDevicePointer(&d, host_, 0) != cudaSuccess) {
-            err = "native embedding: no device alias for the mapped table";
+        // The table is gathered by kernels. A host-mapped alias is not reliable on this
+        // Windows HIP stack, so the rows live in device memory.
+        if (cudaMalloc(&dev_owned_, bytes_) != cudaSuccess) {
+            dev_owned_ = nullptr;
+            err = "native embedding: cannot allocate " + std::to_string(bytes_ >> 20) + " MiB on device";
             return false;
         }
-        dev_ = d;
+        if (cudaMemcpy(dev_owned_, host_, bytes_, cudaMemcpyHostToDevice) != cudaSuccess) {
+            err = "native embedding: device upload failed";
+            return false;
+        }
+        dev_ = dev_owned_;
         type_ = (int) t->type;
         n_embd_ = n_embd;
         n_vocab_ = n_vocab;

src/kernels/cpu/expert_layout.cpp

@@ -122,10 +122,18 @@ bool expert_layout_load(const std::string& pack_dir, int64_t n_layers, int64_t n
             L.gguf_off[(size_t) (3 * l)] = go;
             L.gguf_off[(size_t) (3 * l + 1)] = uo;
             L.gguf_off[(size_t) (3 * l + 2)] = dox;
-            std::string file;             // v3: the shard that holds this layer (a file name beside --native)
-            if (ss >> file) {
-                if (L.gguf_file.empty()) L.gguf_file.assign((size_t) n_layers, std::string());
-                L.gguf_file[(size_t) l] = file;
+            // v3: one shard name for the whole layer. v4: three names, gate then up then down,
+            // when a quant splits one layer across shards.
+            std::string f0, f1, f2;
+            if (ss >> f0) {
+                if (!(ss >> f1 >> f2)) {
+                    f1 = f0;
+                    f2 = f0;
+                }
+                if (L.gguf_file.empty()) L.gguf_file.assign((size_t) (3 * n_layers), std::string());
+                L.gguf_file[(size_t) (3 * l)] = f0;
+                L.gguf_file[(size_t) (3 * l + 1)] = f1;
+                L.gguf_file[(size_t) (3 * l + 2)] = f2;
             }
         }
         L.fmt[(size_t) l] = f;

src/kernels/cuda/iq_kernels.cu

@@ -53,6 +53,22 @@ __device__ __forceinline__ int2 get_int_from_table_16(const int& q4, const int8_
 #define ggml_cuda_dp4a(a, b, c) __dp4a((a), (b), (c))
 
 // ---------------------------------------------------------------- the dot products (vecdotq.cuh)
+// One Q8_0 block (32 values) against the q8_1 block the caller already selected.
+// qs sits 2 bytes into the block, so the int8 lane is loaded as aligned uint16s.
+__device__ __forceinline__ float vec_dot_q8_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
+                                                   const int& kbx, const int&) {
+    const block_q8_0* bq8_0 = (const block_q8_0*) vbq + kbx;
+    const uint16_t* q16 = (const uint16_t*) bq8_0->qs;
+    int sumi = 0;
+#pragma unroll
+    for (int j = 0; j < QK8_0 / 4; ++j) {
+        const int u = (int) q16[2 * j] | ((int) q16[2 * j + 1] << 16);
+        const int v = get_int_b4(bq8_1->qs, j);
+        sumi = ggml_cuda_dp4a(u, v, sumi);
+    }
+    return (float) bq8_0->d * __low2float(bq8_1->ds) * (float) sumi;
+}
+
 __device__ __forceinline__ float vec_dot_q2_0_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
                                                    const int& kbx, const int& iqs) {
     const block_q2_0* bq2_0 = (const block_q2_0*) vbq + kbx;
@@ -311,6 +327,8 @@ template<> struct Fmt<29> { static constexpr int qk = 256, ipb = 8, step = 1;
     __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_iq1_m_q8_1(v, y, kbx, iqs); } };
 template<> struct Fmt<42> { static constexpr int qk = 64, ipb = 2, step = 1;
     __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_q2_0_q8_1(v, y, kbx, iqs); } };
+template<> struct Fmt<8> { static constexpr int qk = 32, ipb = 1, step = 1;
+    __device__ static float dot(const void* v, const block_q8_1* y, int kbx, int iqs) { return vec_dot_q8_0_q8_1(v, y, kbx, iqs); } };
 
 __device__ __forceinline__ float warp_sum(float v) {
 #pragma unroll
@@ -555,6 +573,17 @@ __device__ void dq_iq4_xs(const void* vx, int64_t ibs, dst_t* yy, int tid) {
         y[j + 16] = cvt<dst_t>(d * kvalues_iq4nl[q4[j] >> 4]);
     }
 }
+// Q8_0 (Unsloth token_embd): 256 values = 8 blocks of 32. ggml value is d * int8.
+template<typename dst_t>
+__device__ void dq_q8_0(const void* vx, int64_t ibs, dst_t* yy, int tid) {
+    const block_q8_0* x = (const block_q8_0*) vx + ibs * (QK_K / QK8_0);
+    const int b = tid / 4, part = tid % 4;
+    const float d = (float) x[b].d;
+    for (int j = 0; j < 8; ++j) {
+        const int i = part * 8 + j;
+        yy[b * QK8_0 + i] = cvt<dst_t>(d * (float) x[b].qs[i]);
+    }
+}
 template<typename dst_t>
 __device__ void dq_q2_0(const void* vx, int64_t ibs, dst_t* yy, int tid) {
     // one "superblock" = 256 values = 4 blocks of 64; thread tid writes 8 values
@@ -580,6 +609,7 @@ __device__ __forceinline__ void dq_dispatch(int ty, const void* vx, int64_t ibs,
         case 29: dq_iq1_m(vx, ibs, y, tid); break;
         case 23: dq_iq4_xs(vx, ibs, y, tid); break;
         case 11: dq_q3_k(vx, ibs, y, tid); break;
+        case 8: dq_q8_0(vx, ibs, y, tid); break;
         case 42: dq_q2_0(vx, ibs, y, tid); break;
         default: break;
     }
@@ -618,6 +648,7 @@ size_t iq_row_bytes(int t, int64_t n) noexcept {
         case 23: return (size_t) (n / 256) * sizeof(block_iq4_xs);
         case 11: return (size_t) (n / 256) * sizeof(block_q3_K);
         case 42: return (size_t) (n / 64) * sizeof(block_q2_0);
+        case 8: return (size_t) (n / QK8_0) * sizeof(block_q8_0);
         default: return 0;
     }
 }
@@ -669,14 +700,14 @@ __global__ void embed_rows_kernel(int ty, const uint8_t* __restrict__ table, siz
 void iq_embed_rows(int t, const void* table, size_t row_bytes, const int32_t* tokens, int64_t n_tok, int64_t n_embd,
                    float* out, void* stream) {
     if (n_tok <= 0) return;
-    if (n_embd % 256 != 0 || !is_iq(t)) { std::fprintf(stderr, "iq_embed_rows: bad arguments\n"); std::exit(1); }
+    if (n_embd % 256 != 0 || !(is_iq(t) || t == 8)) { std::fprintf(stderr, "iq_embed_rows: bad arguments\n"); std::exit(1); }
     embed_rows_kernel<<<dim3((unsigned) (n_embd / 256), (unsigned) n_tok), 32, 0, (cudaStream_t) stream>>>(
         t, (const uint8_t*) table, row_bytes, tokens, n_embd, out);
     check("iq_embed_rows");
 }
 
 void iq_dequant_f32(int t, const void* src, int64_t n, float* dst, void* stream) {
-    if (n % 256 != 0 || !is_iq(t)) { std::fprintf(stderr, "iq_dequant_f32: bad arguments\n"); std::exit(1); }
+    if (n % 256 != 0 || !(is_iq(t) || t == 8)) { std::fprintf(stderr, "iq_dequant_f32: bad arguments\n"); std::exit(1); }
     dequant_flat_kernel<float><<<(unsigned) (n / 256), 32, 0, (cudaStream_t) stream>>>(t, src, dst);
     check("iq_dequant_f32");
 }
@@ -726,6 +757,8 @@ void native_expert_grouped(const NativeExpertLayout& L, const unsigned long long
         case 21: native_gu_kernel<21><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 22: native_gu_kernel<22><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 23: native_gu_kernel<23><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
+        case 20: native_gu_kernel<20><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
+        case 8: native_gu_kernel<8><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 29: native_gu_kernel<29><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         case 42: native_gu_kernel<42><<<ggu, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break;
         default: std::fprintf(stderr, "native_expert_grouped: gate/up type %d\n", L.gu_type); std::exit(1);
@@ -737,6 +770,7 @@ void native_expert_grouped(const NativeExpertLayout& L, const unsigned long long
     const dim3 gd((unsigned) ((L.n_embd + 7) / 8), (unsigned) cap_groups);
     switch (L.d_type) {
         case 20: native_down_kernel<20><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
+        case 8: native_down_kernel<8><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
         case 23: native_down_kernel<23><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
         case 42: native_down_kernel<42><<<gd, 256, 0, s>>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break;
         default: std::fprintf(stderr, "native_expert_grouped: down type %d\n", L.d_type); std::exit(1);

src/prefill/gemm.cu

@@ -42,6 +42,20 @@ void ck(cublasStatus_t s, const char* what) {
     }
 }
 
+// gfx1201 hipBLAS returns success and the correct BF16/FP16 product for some
+// shapes (hc up once T>=96, the router), then leaves hipErrorInvalidValue set.
+// The multiply has finished; the next kernel check would otherwise exit.
+void absorb_hipblas_sticky(const char* what) {
+#if defined(__HIPCC__)
+    const hipError_t sticky = hipGetLastError();
+    if (sticky == hipSuccess || sticky == hipErrorInvalidValue) return;
+    std::fprintf(stderr, "prefill gemm: %s left %s\n", what, hipGetErrorString(sticky));
+    std::exit(1);
+#else
+    (void) what;
+#endif
+}
+
 #if defined(__HIPCC__) && defined(STRATA_HIPBLASLT_AVAILABLE)
 struct HipLtCallKey {
     strata::prefill::hipblaslt::InputType type;
@@ -345,6 +359,7 @@ void Gemm::bf16(const uint16_t* X, const uint16_t* W, float* Y, int64_t T, int64
 #if defined(__HIPCC__) && defined(STRATA_HIPBLASLT_AVAILABLE)
     if (try_hipblaslt(hipblaslt_state_, strata::prefill::hipblaslt::InputType::bf16, X, W, Y, T, N, K, ldy,
                       beta, stream_)) {
+        absorb_hipblas_sticky("hipBLASLt bf16");
         return;
     }
 #endif
@@ -353,6 +368,7 @@ void Gemm::bf16(const uint16_t* X, const uint16_t* W, float* Y, int64_t T, int64
                     CUDA_R_16BF, (int) K, X, CUDA_R_16BF, (int) K, &beta, Y, CUDA_R_32F, (int) ldy,
                     CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT),
        "cublasGemmEx");
+    absorb_hipblas_sticky("cublasGemmEx");
 }
 
 void Gemm::f16(const uint16_t* X, const uint16_t* W, float* Y, int64_t T, int64_t N, int64_t K, int64_t ldy,
@@ -363,6 +379,7 @@ void Gemm::f16(const uint16_t* X, const uint16_t* W, float* Y, int64_t T, int64_
 #if defined(__HIPCC__) && defined(STRATA_HIPBLASLT_AVAILABLE)
     if (try_hipblaslt(hipblaslt_state_, strata::prefill::hipblaslt::InputType::f16, X, W, Y, T, N, K, ldy,
                       beta, stream_)) {
+        absorb_hipblas_sticky("hipBLASLt f16");
         return;
     }
 #endif
@@ -370,6 +387,7 @@ void Gemm::f16(const uint16_t* X, const uint16_t* W, float* Y, int64_t T, int64_
                     CUDA_R_16F, (int) K, X, CUDA_R_16F, (int) K, &beta, Y, CUDA_R_32F, (int) ldy,
                     CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT),
        "cublasGemmEx f16");
+    absorb_hipblas_sticky("cublasGemmEx f16");
 }
 
 void Gemm::native(const uint16_t* X, int ggml_type, const void* W_blocks, float* Y, int64_t T, int64_t N, int64_t K,

src/program/generate.cpp

@@ -37,6 +37,7 @@
 #include "strata/kernels/native_qsa_indexer.hpp"
 #include "strata/kernels/native_rope.hpp"
 #include "strata/kernels/mrope.hpp"
+#include "strata/kernels/f16_bits.hpp"
 #include "strata/kernels/kv_q4.hpp"
 #include "strata/kernels/qsa.hpp"
 #include "strata/core/native_head.hpp"
@@ -439,18 +440,26 @@ void usage() {
                  "                       Borrowed GPU-cache entries may read from mmap during prompt prefill.\n");
 }
 
-/// All shards of a split GGUF, from shard 1's path ("...-00001-of-00002.gguf"); just the path when it is not split.
+/// All shards of a split GGUF ("...-0000N-of-0000M.gguf"), from whichever shard path was given.
+/// A path that is not a numbered split is returned as itself.
 std::vector<std::string> model_shards(const std::string& first) {
-    const std::string tag = "-00001-of-";
-    const size_t at = first.rfind(tag);
-    if (at == std::string::npos || first.size() < at + tag.size() + 10) return {first};
-    const int total = std::atoi(first.substr(at + tag.size(), 5).c_str());
+    const size_t of = first.rfind("-of-");
+    if (of == std::string::npos || of < 6 || first.size() < of + 4 + 5 + 5) return {first};
+    const size_t num_at = of - 5;
+    if (first[num_at - 1] != '-') return {first};
+    for (size_t i = 0; i < 5; ++i) {
+        const char c = first[num_at + i];
+        const char t = first[of + 4 + i];
+        if (c < '0' || c > '9' || t < '0' || t > '9') return {first};
+    }
+    const int total = std::atoi(first.substr(of + 4, 5).c_str());
+    if (total < 1 || total > 99) return {first};
     std::vector<std::string> out;
-    for (int i = 1; i <= total && i <= 99; ++i) {
+    for (int i = 1; i <= total; ++i) {
         char num[8];
         std::snprintf(num, sizeof num, "%05d", i);
         std::string p = first;
-        p.replace(at + 1, 5, num);
+        p.replace(num_at, 5, num);
         if (std::ifstream(p, std::ios::binary)) out.push_back(p);
     }
     return out.empty() ? std::vector<std::string>{first} : out;
@@ -1562,6 +1571,7 @@ int main(int argc, char** argv) {
     std::vector<float> ple_emb_host((size_t) strata::kernels::NG_N_EMBD);
     float* ple_emb_dev = nullptr;
     float* ple_scratch = nullptr;
+    uint16_t* ple_conv_f16 = nullptr;
     if (!o.ple_gguf.empty()) {
         strata::kernels::PleIoOptions pio;
         pio.mode = o.ple_io == "mmap" ? strata::kernels::PleIo::Mmap : strata::kernels::PleIo::Direct;
@@ -1605,7 +1615,34 @@ int main(int argc, char** argv) {
         ss.ple.w.norm_key = (const float*) wnk->data;
         ss.ple.w.norm_query = (const float*) wnq->data;
         ss.ple.w.norm_conv = (const float*) wnc->data;
-        ss.ple.w.conv1d_f16 = (const uint16_t*) wc->data;
+        // iq_pack keeps a true F32 ple_conv1d (kind 2, 4 bytes per tap). Both conv kernels
+        // read FP16 in ggml order k + 4*c. An F16 or F16-in-F32 pack is already 2 bytes here.
+        if (wc->kind == strata::core::WeightKind::F32) {
+            const uint64_t n = (uint64_t) wc->ne0 * (uint64_t) (wc->ne1 > 0 ? wc->ne1 : 1);
+            if (wc->data == nullptr || n == 0 || wc->bytes != n * sizeof(float)) {
+                std::fprintf(stderr, "strata generate: ple_conv1d is F32 but its span is not %llu floats\n",
+                             (unsigned long long) n);
+                return 1;
+            }
+            std::vector<float> host((size_t) n);
+            std::vector<uint16_t> half((size_t) n);
+            if (cudaMemcpy(host.data(), wc->data, (size_t) wc->bytes, cudaMemcpyDeviceToHost) != cudaSuccess) {
+                std::fprintf(stderr, "strata generate: ple_conv1d F32 readback failed\n");
+                return 1;
+            }
+            for (uint64_t i = 0; i < n; ++i) half[(size_t) i] = strata::kernels::f16_from_f32(host[(size_t) i]);
+            if (cudaMalloc((void**) &ple_conv_f16, (size_t) n * sizeof(uint16_t)) != cudaSuccess ||
+                cudaMemcpy(ple_conv_f16, half.data(), (size_t) n * sizeof(uint16_t), cudaMemcpyHostToDevice) !=
+                    cudaSuccess) {
+                std::fprintf(stderr, "strata generate: ple_conv1d F16 upload failed\n");
+                return 1;
+            }
+            ss.ple.w.conv1d_f16 = ple_conv_f16;
+            std::fprintf(stderr, "strata generate: ple_conv1d converted F32 to F16 (%llu values)\n",
+                         (unsigned long long) n);
+        } else {
+            ss.ple.w.conv1d_f16 = (const uint16_t*) wc->data;
+        }
         ss.ple.consts = strata::kernels::ple_artifact_consts();
         if (o.ple_delay_us > 0) ple_table.set_injected_delay_us(o.ple_delay_us);
         ss.ple.table = &ple_table;

tools/iq_pack.py

@@ -314,12 +314,15 @@ def main() -> int:
                  % (n_expert, offset, src.name))
         for l, gt, dt, off, blob, ts in layout:
             ws = [model.where[t.name] for t in ts]
-            if len({w[3] for w in ws}) != 1:
-                print("layer %d: its gate/up/down tensors are in different shards" % l)
-                return 1
-            gg, shard = ws[0][0], ws[0][3]
-            line = "%d %d %d %d %d %d %d %d" % (l, gt, dt, off, blob, *[gg.data_start + t.offset for t in ts])
-            fo.write(line + ("" if shard == src else " " + shard.name) + "\n")
+            names = [w[3].name for w in ws]
+            line = "%d %d %d %d %d %d %d %d" % (
+                l, gt, dt, off, blob, *[w[0].data_start + t.offset for w, t in zip(ws, ts)])
+            if len(set(names)) == 1:
+                if names[0] != src.name:
+                    line += " " + names[0]
+            else:
+                line += " " + " ".join(names)
+            fo.write(line + "\n")
     if a.skip_experts or not a.experts_bin:
         if (out / "experts.bin").exists() and not a.experts_bin:
             print("note: %s/experts.bin exists; the engine reads it instead of the GGUF" % out)