port: Windows native ROCm build for AMD Strix Halo (gfx1151)

[paperlinguist]

Summary

Native Windows build target for DS4 on AMD Strix Halo (Radeon 8060S, gfx1151,
RDNA4) using the ROCm 7.1 HIP SDK.

Tested on: AMD Ryzen AI Max / Strix Halo, 128GB unified RAM, Windows 11, ROCm 7.1
Performance: ~18 t/s prefill / ~16 t/s generation at --ctx 229376 (224k tokens)
Model: DeepSeek-V4-Flash-IQ2XXS-w2Q2K-AProjQ8-SExpQ8-OutQ8-chat-v2-imatrix.gguf (80.76 GiB)

---

What this PR adds

New: compat/win32/ — complete POSIX shim layer

Maps POSIX APIs to Win32 equivalents so all existing source files compile without
modification (except the targeted fixes below):

Header	Key mappings
unistd.h	pread() (thread-safe via ReOpenFile+OVERLAPPED), close() (CRT fd + Winsock socket via GetProcAddress), clock_gettime, usleep, PATH_MAX, off_t=__int64
sys/socket.h	Winsock2 with lazy WSAStartup on first socket(), poll() via WSAPoll
sys/file.h	flock, fcntl no-op, dprintf, sigaction stub, winsize/TIOCGWINSZ
sys/time.h	gettimeofday via GetSystemTimeAsFileTime
sys/mman.h/c	mmap/munmap/msync/mlock via CreateFileMapping/MapViewOfFile (64-bit)
dirent.h	opendir/readdir/closedir via FindFirstFile/FindNextFile
strings.h	strcasecmp / strncasecmp macros to _stricmp / _strnicmp

New: Makefile strix-halo-windows target

All compilation goes through hipcc targeting x86_64-pc-windows-msvc to maintain
a single MSVC ABI throughout (mixing MinGW-compiled objects with MSVC-ABI ROCm DLLs
causes linker failures). Offload arch: gfx1151.

---

Bug fixes required for Windows

ds4.c — 64-bit stat for files > 4 GB

MSVC off_t is 32-bit. fstat() on an 86 GB file returns EOVERFLOW. Fixed by
using _fstat64 / struct _stat64 on _WIN32.

ds4_server.c — WSAStartup never called

The original code guarded out #include <sys/socket.h> on Windows, so the compat
header (which lazily calls WSAStartup) was never included. Every socket() call
returned WSANOTINITIALISED (10093), surfacing as errno 132 / "value too large".

Symptom: ds4-server: failed to listen on 127.0.0.1:8000: value too large

rocm/ds4_rocm_runtime.cuh — 64-bit pread offset

cuda_pread_full cast the file offset to (off_t) which is 32-bit on MSVC.
Offsets beyond 4 GB were silently truncated to negative values, causing every
tensor read past the 4 GB boundary to fail with EIO.

Fix: Cast to (int64_t) instead of (off_t).

rocm/ds4_rocm_runtime.cuh — Q8 cache reserve for Windows

The ROCm Windows driver consumes ~12 GiB of overhead (staging buffers, HipBLAS
workspaces) not present on Linux. The default 5% Q8->F16 cache reserve was too
small: the cache consumed 10+ GiB before inference tensors could be allocated.

Fix: Add +12 GiB to the reserve on _WIN32, leaving ~15 GiB free for the
KV cache.

compat/win32/unistd.h — thread-safe pread

The ROCm streaming engine uses 18 concurrent worker threads all calling pread()
on the same g_model_fd. Concurrent ReadFile on a synchronous Windows HANDLE
races on the shared internal file pointer. Fixed by using ReOpenFile to obtain a
private FILE_FLAG_OVERLAPPED handle per call, with GetOverlappedResult for
async completion.

---

Documentation

WINDOWS_ROCM.md covers full prerequisites (AMD HIP SDK, MSYS2, rocWMMA headers,
import libraries, runtime PATH), build steps, every source change with root cause
and fix, and known remaining issues.

---

Known issues / not in scope

• --ssd-streaming crashes at ~8 GiB: cudaMemcpy(mmap_ptr, HostToDevice) fallback fails on Windows ROCm for demand-paged file memory. Not needed when the full model fits in the 107.87 GiB pool.
• isLargeBar: 0: Resizable BAR not enabled; enabling it may increase max context beyond 232k tokens.
• ds4-agent: uses fork() which is not ported; only ds4, ds4-server, ds4-bench, and ds4-eval are built.