HIP: run Qwen3.8-Flash-Next on Windows gfx1200/gfx1201 (port of #247)

Applies jagsan-cyber's Windows HIP work (PR #247, commit aea1ec1) onto
current main (30ec18e), resolving the conflicts by keeping the newer
upstream logic where it supersedes the PR's.

Carried from #247:
- Windows AMD detection through the display-class registry (device id,
  then DriverDesc name), so setup sees the card with no ROCm present.
- ROCm 10 rocm-sdk / HIP SDK resolution on Windows (ROCM_PATH first),
  including the amdclang compiler and per-arch device bitcode, and
  gfx1200 as a build target alongside gfx1100 and gfx1201.
- CMake refuses to mix MSVC host flags with Clang HIP: ROCm's clang++
  compiles the host code too, and the CUDA shim is force-included with
  /FI on MSVC.
- The gfx1201 hipBLAS sticky hipErrorInvalidValue is cleared after a
  successful GEMM, so the next kernel's error check does not exit.
- Multi-shard expert packs (gate/up/down in different shards) and Q8_0
  embedding rows, so a 3-shard Unsloth GGUF packs natively.
- An F32 ple_conv1d is converted to FP16 at startup (the conv kernels
  read FP16).

Kept from main, where the PR's copy was older:
- The STRATA_HIP_ARCHS compile-time device check in device.cu, and the
  arch lists with the validated/unvalidated split, so a binary carried
  to a card it has no code for still names the card and the build's
  architectures instead of failing later.
- The device.cu/gpu-arch reporting that reads STRATA_HIP_ARCHS.
- The gfx1101/gfx1102 dp4a targets in intrinsics.hpp (a superset of the
  PR's list).
- iq_pack.py's tmp.replace() completion mark (#173), which the PR's
  hunk would have dropped.

Fixed in the merge, not in either side:
- native_head: hipHostGetDevicePointer returns the HOST pointer on the
  Windows HIP stack, so the mapped-alias path yields a "device" address
  that is really host memory. The table is only gathered from, so a
  failed or suspicious alias now falls back to a real device copy
  instead of erroring out or reading over PCIe.

Verified: setup.py imports on Windows and detects this PC's RX 9070 XT
as gfx1201 / 15.9 GB. Not yet built or run; no HIP toolchain on this box.

Co-Authored-By: Claude Opus 5.5 (1M context) <[email protected]>

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 wave32 AMD GPUs (gfx1100, gfx1201)" OFF)
+option(STRATA_ENABLE_HIP "Build the HIP targets for wave32 AMD GPUs (gfx1100, gfx1200, gfx1201)" 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")

cmake/hip_backend.cmake

@@ -9,8 +9,14 @@ endif()
 # the maintainers have not run them (community reports: gfx1102 #192, gfx1200 #176).
 set(_strata_hip_validated gfx1100 gfx1201)
 set(_strata_hip_unvalidated gfx1101 gfx1102 gfx1200)
+# CMake hands HIP a ';' list, but ROCm's own Windows tooling (and a -DCMAKE_HIP_ARCHITECTURES
+# typed by hand) uses spaces, which foreach(IN LISTS) would otherwise treat as one element.
+string(REPLACE " " ";" _strata_hip_norm "${CMAKE_HIP_ARCHITECTURES}")
 set(STRATA_HIP_ARCH_LIST "")
-foreach(_arch IN LISTS CMAKE_HIP_ARCHITECTURES)
+foreach(_arch IN LISTS _strata_hip_norm)
+  if(_arch STREQUAL "")
+    continue()
+  endif()
   string(REGEX REPLACE ":.*$" "" _base "${_arch}")      # gfx1100:xnack- -> gfx1100
   if(_base IN_LIST _strata_hip_validated)
   elseif(_base IN_LIST _strata_hip_unvalidated)
@@ -57,11 +63,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 "STRATA_HIP_ARCHS=\"${STRATA_HIP_ARCHS}\"")
 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

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/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

@@ -717,10 +717,20 @@ ROCM_INDEXES = {"gfx1100": "https://rocm.nightlies.amd.com/v2/gfx110X-dgpu/",
                 "gfx1201": "https://rocm.nightlies.amd.com/v2/gfx120X-all/"}
 ROCM_VERSION = os.environ.get("STRATA_ROCM_VERSION", "7.10.0a20251120")   # what Strata's HIP build was tested with
 ROCM_SYSTEM_MIN = (7, 0)       # an older system ROCm is passed over for the wheels (gfx1201 needs ROCm 6.4 or newer)
-AMD_ARCHS = ("gfx1100", "gfx1201")
+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)",   # when sysfs has no product name
+             "gfx1200": "AMD Radeon RX 9060 series (gfx1200)",
              "gfx1201": "AMD Radeon RX 9070 series / AI PRO R9700 (gfx1201)"}
-AMD_CARDS = "the RX 7900 XT / XTX (gfx1100) and the RX 9070 / 9070 XT / Radeon AI PRO R9700 (gfx1201)"
+AMD_CARDS = ("the RX 7900 XT / XTX (gfx1100), the RX 9060 (gfx1200) and the RX 9070 / 9070 XT / Radeon AI PRO "
+             "R9700 (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 rocm_index(arch):
@@ -730,9 +740,11 @@ def rocm_index(arch):
 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:
@@ -777,16 +789,182 @@ def rocm_version(root):
         return None
 
 
-def rocm_root(arch):
-    """ROCm for compiling and running the HIP engine for `arch`: (root, library folders).  A system ROCm 7 with hipcc
-    and hipBLAS, else AMD's TheRock wheels (ROCM_VERSION, from the card family's 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, prefer: list[Path] | None = None) -> list[str]:
+    dirs = []
+    for rel in ("bin", "lib", "lib/llvm/bin", "llvm/bin"):
+        p = root / rel
+        if p.is_dir():
+            dirs.append(str(p))
+    # The wheels' per-GPU-family library folders come first: one of them holds the hipBLASLt with this
+    # card's kernels, and a family that has none would fail at kernel launch.
+    for d in (prefer or []):
+        if (d / "lib").is_dir() and str(d / "lib") not in dirs:
+            dirs.insert(0, str(d / "lib"))
+    for sp in {Path(p) for p in sys.path if p.endswith("site-packages")}:
+        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*")):
-        ver = rocm_version(sysroot)
-        if ver is None or ver >= ROCM_SYSTEM_MIN:
-            return sysroot, [str(sysroot / "lib")]
-        warn(f"the ROCm in {sysroot} is {ver[0]}.{ver[1]}; Strata needs {ROCM_SYSTEM_MIN[0]}.{ROCM_SYSTEM_MIN[1]} or "
-             "newer: using AMD's wheels in .venv instead")
+    if _rocm_usable(sysroot, arch) or (_rocm_has_hipcc(sysroot) and _rocm_has_hipblas(sysroot) and arch == "gfx1100"):
+        return sysroot, [str(sysroot / "lib")]
     index = rocm_index(arch)
     stamp = Path(sys.prefix) / ".strata-rocm.json"
     have = json.loads(stamp.read_text()) if stamp.exists() else {}
@@ -799,19 +977,21 @@ def rocm_root(arch):
         stamp.write_text(json.dumps({"version": ROCM_VERSION, "index": 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")
-    # the card family's libraries only (gfx120X-all -> _rocm_sdk_libraries_gfx120X_all): another family's
-    # libhipblaslt.so first on the path would have no kernels for this card
+    # The card family's libraries only (gfx120X-all -> _rocm_sdk_libraries_gfx120X_all): another family's
+    # libhipblaslt first on the path would have no kernels for this card.
     family = "_rocm_sdk_libraries_" + index.rstrip("/").rsplit("/", 1)[-1].replace("-", "_")
-    dirs = [str(root / "lib")]
+    libs = []
     for sp in {Path(p) for p in sys.path if p.endswith("site-packages")}:
-        libs = sorted(sp.glob("_rocm_sdk_libraries_*"))
-        libs = [d for d in libs if d.name.lower() == family.lower()] or libs
-        dirs += [str(d / "lib") for d in libs if (d / "lib").is_dir()]
+        found = sorted(sp.glob("_rocm_sdk_libraries_*"))
+        match = [d for d in found if d.name.lower() == family.lower()]
+        if not match and arch == "gfx1200":
+            match = [d for d in found if "gfx120x" in d.name.lower()]     # gfx1200 rides with gfx120X
+        libs += [d for d in (match or found) if (d / "lib").is_dir()]
     ok(f"ROCm: {root}")
-    return root, list(dict.fromkeys(dirs))
+    return root, _rocm_lib_dirs(root, libs)
 
 
 def hipblaslt_version(lib_dirs):
@@ -853,31 +1033,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 and             gpu["arch"] in meta.get("archs", []):
+    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(gpu["arch"])
-    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")
+    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() and gpu["arch"] in meta.get("archs", [])
         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}")
@@ -1767,8 +1975,8 @@ def main() -> int:
                          "this mode the experts the GPU does not hold are copied into RAM once when they fit (resident), "
                          "else read through the OS file cache (mmap); resident / mmap force one of the two")
     ap.add_argument("--backend", choices=["cuda", "hip"],
-                    help="cuda = NVIDIA (default), hip = AMD RX 7900 XT/XTX or RX 9070 / AI PRO R9700 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)
@@ -1842,7 +2050,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))
@@ -1858,14 +2066,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", f"the AMD backend runs on {AMD_CARDS} on Linux")
+            fail("no AMD GPU Strata can use",
+                 f"the AMD backend runs on {AMD_CARDS}")
         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:
@@ -1882,7 +2089,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 or RX 9070 / AI PRO R9700: --backend hip" if amd else ""))
+                 + (f"; 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/expert_source.cpp

@@ -1413,9 +1413,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;
@@ -1424,19 +1427,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

@@ -25,6 +25,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()) {
@@ -102,6 +105,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_);
     else if (dev_) cudaFree(const_cast<void*>(dev_));   // the VRAM fallback below
 }
@@ -113,7 +117,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;
@@ -140,11 +144,24 @@ bool NativeEmbed::load(const std::string& path, int64_t n_embd, int64_t n_vocab,
         } else {
             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";
-                return false;
+            if (cudaHostGetDevicePointer(&d, host_, 0) == cudaSuccess && d != nullptr) {
+                dev_ = d;
+            } else {
+                // On the Windows HIP stack hipHostGetDevicePointer hands back the host pointer, so the
+                // kernels would read over PCIe through a "device" address that is really host memory
+                // (and it may be non-null but unmapped). The table is only gathered from, so a real
+                // device copy is the portable answer: it costs its size in VRAM and reads far faster.
+                cudaGetLastError();
+                if (cudaMalloc(&d, bytes_) != cudaSuccess ||
+                    cudaMemcpy(d, gguf.tensor_data(*t), bytes_, cudaMemcpyHostToDevice) != cudaSuccess) {
+                    if (d) cudaFree(d);
+                    cudaGetLastError();
+                    err = "native embedding: no device alias for the mapped table, and no VRAM to copy it into";
+                    return false;
+                }
+                std::fprintf(stderr, "strata: native embedding: host mapping is unusable, kept in VRAM instead\n");
+                dev_ = d;
             }
-            dev_ = d;
         }
         type_ = (int) t->type;
         n_embd_ = n_embd;

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

@@ -54,6 +54,22 @@ __device__ __forceinline__ int2 get_int_from_table_16(const int& q4, const int8_
 #define ggml_cuda_dp4a(a, b, c) STRATA_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;
@@ -312,6 +328,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
@@ -556,6 +574,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
@@ -581,6 +610,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;
     }
@@ -619,6 +649,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;
     }
 }
@@ -670,14 +701,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");
 }
@@ -727,6 +758,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);
@@ -738,6 +771,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

@@ -40,6 +40,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"
@@ -528,18 +529,26 @@ void usage() {
                  "                       does not fit.  Same answers as --mmap-experts for the same placement.\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;
@@ -1829,6 +1838,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" || o.ple_io == "ram" ? strata::kernels::PleIo::Mmap : strata::kernels::PleIo::Direct;
@@ -1878,7 +1888,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

@@ -317,14 +317,17 @@ 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)
-                fo.close()
-                tmp.unlink()
-                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")
+    # The write is only trusted once it is complete: a torn native_experts.txt from an interrupted
+    # run would otherwise be read back as a finished pack.
     tmp.replace(out / "native_experts.txt")
     if a.skip_experts or not a.experts_bin:
         if (out / "experts.bin").exists() and not a.experts_bin: