Windows: direct I/O (FILE_FLAG_NO_BUFFERING, 1.47x decode) + VirtualLock pin wiring + select_ctx device cache (#162)

* win: direct I/O via FILE_FLAG_NO_BUFFERING + compat_fsize + VirtualLock primitives compat_open_direct() gives Windows the O_DIRECT twin fd st.h already uses on Linux/macOS: FILE_FLAG_NO_BUFFERING, same 4K-alignment contract as O_DIRECT (the engine's DIRECT=1 path already aligns offset/len and slabs are posix_memalign'd). Measured on GLM-5.2 744B int4, Ryzen 9 9950X3D / 126 GB / PCIe4 NVMe (5.8 GB/s at the engine's 19MBx8T pattern), Windows 11, MinGW GCC 16.1, 32-token greedy runs at --topp 0.7, 40 GB pin, current dev HEAD: buffered: 0.38 tok/s (expert-disk dominates) DIRECT=1: 0.56 tok/s (1.47x) — byte-identical greedy output vs buffered compat_fsize() (GetFileSizeEx): CRT lseek(SEEK_END) returns -1 on NO_BUFFERING fds (measured on UCRT); iobench uses it and gains a NO_BUFFERING branch so disk numbers are comparable across platforms. compat_mlock/compat_munlock: VirtualLock with working-set growth (bare VirtualLock caps at the default working-set minimum, a few hundred KB). Wired into the engine in the next commit. tests/test_compat_direct.c covers the alignment contract, data integrity, fsize on both fd kinds; skips cleanly off Windows. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * win: wire VirtualLock into mem_wire, munlock pairing in expert_host_release MLOCK=1 was a silent no-op on Windows: pinned experts could be paged out by working-set trimming under memory pressure. mem_wire now uses compat_mlock (VirtualLock + working-set growth); expert_host_release unlocks before freeing, mirroring the POSIX branch. Validated: 39.6 GB pin wired in 17s on a 126 GB machine, zero failures; TF oracle 32/32 with MLOCK=1. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> * cuda: thread-local current-device cache in select_ctx cudaSetDevice on every call is expensive when the serial expert loop alternates devices. Measured on RTX 5090 + RTX 4090 (Windows, DLL backend, pre-#68 dispatch): expert-matmul 14.3s -> 25.4s per 32 tokens going from 1 to 2 devices, entirely per-call context switching. The current device is per-thread in the CUDA runtime, so a thread_local cache skips redundant switches; multi-GPU expert serving becomes positive-scaling instead of negative. Kernel suite passes on sm_120 + sm_89; TF oracle 32/32 dual-GPU. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> --------- Co-authored-by: olorin <io@zyphyr.co> Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
2026-07-14 07:12:15 -07:00
parent ec0aadf91a
commit d47b095875
7 changed files with 138 additions and 4 deletions
@@ -115,7 +115,7 @@ else
 PYTHON    ?= python
 endif
 CUDA_OBJ  =
-TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_acc512$(EXE)
+TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE)

 # Windows CUDA DLL path: host links the loader, NOT cudart.
 ifneq ($(IS_WIN),)
@@ -229,6 +229,9 @@ tests/test_idot$(EXE): tests/test_idot.c glm.c st.h json.h tok.h tok_unicode.h c
 tests/test_i4_acc512$(EXE): tests/test_i4_acc512.c
 	$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)

+tests/test_compat_direct$(EXE): tests/test_compat_direct.c compat.h
+	$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
+
 test-c: $(TEST_BINS)
 	$(PYTHON) tools/run_tests.py $(TEST_BINS)

@@ -51,8 +51,18 @@ static DeviceContext *find_ctx(int device) {
    return nullptr;
 }

+/* cudaSetDevice on every call doubles expert-matmul time on 2 GPUs when the
+ * serial expert loop alternates devices (measured on RTX 5090 + 4090: 14.3s
+ * -> 25.4s per 32 tokens). The current device is per-thread in the CUDA
+ * runtime, so a thread-local cache skips the redundant switches. */
+static thread_local int g_current_device = -1;
+
 static int select_ctx(DeviceContext *ctx) {
-    return ctx && cuda_ok(cudaSetDevice(ctx->device), "select device");
+    if (!ctx) return 0;
+    if (g_current_device == ctx->device) return 1;
+    if (!cuda_ok(cudaSetDevice(ctx->device), "select device")) return 0;
+    g_current_device = ctx->device;
+    return 1;
 }

 __host__ __device__ static size_t row_bytes(int fmt, int I) {
@@ -57,6 +57,7 @@ static inline int compat_open_direct(const char *path){
 *                                  _read/_lseeki64 which are racy AND
 *                                  corrupt 0x0A bytes in binary files).
 * posix_fadvise -> no-op (advisory only; macOS already no-ops DONTNEED).
+ * mlock         -> compat_mlock  (VirtualLock + crescita working set).
 * posix_memalign->_aligned_malloc(free must be compat_aligned_free).
 * rename        -> compat_rename (MoveFileEx MOVEFILE_REPLACE_EXISTING;
 *                                  CRT rename fails EEXIST if dest exists,
@@ -135,6 +136,25 @@ static inline ssize_t compat_pread(int fd, void *buf, size_t n, off_t off){
 }
 #define pread(fd,buf,n,off) compat_pread(fd,buf,n,off)

+/* --- mlock -> VirtualLock con crescita del working set ---
+ * VirtualLock fallisce oltre il working set MINIMO del processo (default ~qualche
+ * centinaio di KB): prima si allarga il working set di len + margine, poi si blocca.
+ * Best effort come mlock su Linux: -1 su fallimento, il chiamante decide (pin_wire
+ * lo tratta come non-fatale). SeIncreaseWorkingSetPrivilege e' concesso agli utenti
+ * standard di default. */
+static inline int compat_mlock(const void *addr, size_t len){
+    HANDLE p = GetCurrentProcess();
+    SIZE_T mn = 0, mx = 0;
+    if(GetProcessWorkingSetSize(p, &mn, &mx)){
+        SIZE_T need = len + (SIZE_T)(1u<<20);
+        SetProcessWorkingSetSize(p, mn + need, mx + need);   /* best effort */
+    }
+    return VirtualLock((LPVOID)addr, len) ? 0 : -1;
+}
+static inline int compat_munlock(const void *addr, size_t len){
+    return VirtualUnlock((LPVOID)addr, len) ? 0 : -1;
+}
+
 /* --- posix_memalign -> _aligned_malloc ---
 * ATTN: memoria allocata con _aligned_malloc DEVE essere liberata con
 * _aligned_free, NON con free(). Vedi compat_aligned_free sotto.
@@ -215,6 +235,35 @@ static inline ssize_t compat_getline(char **lineptr, size_t *n, FILE *stream){
 }
 #define getline(lineptr,n,stream) compat_getline(lineptr,n,stream)

+/* --- O_DIRECT -> FILE_FLAG_NO_BUFFERING ---
+ * Apre il fd "gemello" senza cache del file system, come il twin O_DIRECT di
+ * st.h su Linux e F_NOCACHE su macOS. Stesso contratto: offset, lunghezza e
+ * buffer del chiamante devono essere allineati a 4K (gli slab expert usano
+ * posix_memalign(4096) e il percorso DIRECT=1 del motore allinea gia' offset
+ * e len); richieste non allineate falliscono con -1, mai dati corrotti.
+ * Il fd si usa con la normale pread() (compat_pread -> ReadFile+OVERLAPPED). */
+static inline int compat_open_direct(const char *path){
+    HANDLE h = CreateFileA(path, GENERIC_READ,
+                           FILE_SHARE_READ|FILE_SHARE_WRITE|FILE_SHARE_DELETE,
+                           NULL, OPEN_EXISTING, FILE_FLAG_NO_BUFFERING, NULL);
+    if(h == INVALID_HANDLE_VALUE) return -1;
+    int fd = _open_osfhandle((intptr_t)h, _O_RDONLY|_O_BINARY);
+    if(fd < 0){ CloseHandle(h); return -1; }
+    return fd;
+}
+
+/* --- dimensione file da fd: GetFileSizeEx ---
+ * La lseek(SEEK_END) del CRT ritorna -1 sui fd NO_BUFFERING (misurato su
+ * UCRT): la dimensione si chiede direttamente al kernel. Funziona su
+ * qualsiasi fd (buffered o direct). -1 su errore. */
+static inline off_t compat_fsize(int fd){
+    intptr_t osfh = _get_osfhandle(fd);
+    if(osfh == -1 || osfh == -2) return -1;
+    LARGE_INTEGER li;
+    if(!GetFileSizeEx((HANDLE)osfh, &li)) return -1;
+    return (off_t)li.QuadPart;
+}
+
 /* --- setenv -> SetEnvironmentVariableA (POSIX setenv assente su Windows) --- */
 static inline int compat_setenv(const char *name, const char *value, int overwrite){
    if(!overwrite && getenv(name)) return 0;
@@ -1457,6 +1457,9 @@ static void expert_host_release(Model *m, ESlot *s){
 #if defined(__APPLE__) || defined(__linux__)
    if(s->slab) munlock(s->slab,(size_t)s->slab_cap);
    if(s->fslab) munlock(s->fslab,(size_t)s->fslab_cap*sizeof(float));
+#elif defined(_WIN32)
+    if(s->slab) compat_munlock(s->slab,(size_t)s->slab_cap);
+    if(s->fslab) compat_munlock(s->fslab,(size_t)s->fslab_cap*sizeof(float));
 #endif
    int64_t bytes=qt_bytes(&s->g)+qt_bytes(&s->u)+qt_bytes(&s->d);
    free(s->slab); free(s->fslab); s->slab=NULL; s->fslab=NULL; s->slab_cap=s->fslab_cap=0;
@@ -3440,6 +3443,8 @@ static int mem_should_wire(void){
 static int mem_wire(void *addr, size_t len){
 #if defined(__APPLE__) || defined(__linux__)
    return mlock(addr, len);
+#elif defined(_WIN32)
+    return compat_mlock(addr, len);   /* VirtualLock + working-set growth */
 #else
    (void)addr; (void)len; return 0;
 #endif
@@ -27,6 +27,10 @@ int main(int argc,char**argv){
    int fd=open(argv[1],O_RDONLY|(direct?O_DIRECT:0));
    if(fd<0 && direct){ fprintf(stderr,"O_DIRECT is unavailable (%s); using buffered I/O\n",strerror(errno));
        direct=0; fd=open(argv[1],O_RDONLY); }
+#elif defined(_WIN32)
+    int fd = direct ? compat_open_direct(argv[1]) : open(argv[1],COMPAT_O_RDONLY);
+    if(fd<0 && direct){ fprintf(stderr,"NO_BUFFERING is unavailable; using buffered I/O\n");
+        direct=0; fd=open(argv[1],COMPAT_O_RDONLY); }
 #else
    int fd=open(argv[1],O_RDONLY);                 /* macOS: F_NOCACHE ~ O_DIRECT */
 #ifdef __APPLE__
@@ -36,7 +40,11 @@ int main(int argc,char**argv){
 #endif
 #endif
    if(fd<0){perror("open");return 1;}
+#ifdef _WIN32
+    off_t sz=compat_fsize(fd);     /* CRT lseek(SEEK_END) ritorna -1 sui fd NO_BUFFERING */
+#else
    off_t sz=lseek(fd,0,SEEK_END);
+#endif
    if(sz<blk*2){fprintf(stderr,"file is too small\n");return 1;}
    /* offset random pre-generati (stessi per ogni configurazione: srand fisso).
     * 30 bit di rand combinati: su Windows RAND_MAX=32767 e un singolo rand()*4096
@@ -81,8 +81,8 @@ static int st_open_fd(shards *S, const char *path) {
    S->paths[S->nfd] = strdup(path); S->fds[S->nfd] = fd;
 #ifdef O_DIRECT
    S->dfds[S->nfd] = open(path, COMPAT_O_RDONLY | O_DIRECT);   /* eager: lookup poi thread-safe */
-#elif defined(__APPLE__)
-    S->dfds[S->nfd] = compat_open_direct(path);          /* macOS: F_NOCACHE ~ O_DIRECT */
+#elif defined(__APPLE__) || defined(_WIN32)
+    S->dfds[S->nfd] = compat_open_direct(path);          /* macOS: F_NOCACHE; Windows: NO_BUFFERING */
 #else
    S->dfds[S->nfd] = -1;                                /* niente equivalente: solo buffered */
 #endif
@@ -0,0 +1,59 @@
+/* test_compat_direct.c — O_DIRECT equivalente su Windows: FILE_FLAG_NO_BUFFERING.
+ * compat_open_direct(path) deve dare un fd che bypassa la cache del file system,
+ * leggibile con pread() a offset/len/buffer allineati a 4K (stesso contratto di
+ * O_DIRECT su Linux: richieste non allineate falliscono, mai dati corrotti). */
+#include <stdio.h>
+
+#ifndef _WIN32
+int main(void){ puts("compat direct tests: skipped (POSIX has native O_DIRECT)"); return 0; }
+#else
+
+#include <stdlib.h>
+#include <string.h>
+#include <stdint.h>
+#include <fcntl.h>
+#include <io.h>
+#include "../compat.h"
+
+static int fail(const char *m){ fprintf(stderr,"compat direct test failed: %s\n",m); return 1; }
+
+#define FSZ (1u<<20)
+#define TMPF "test_direct.tmp"
+
+int main(void){
+    FILE *w=fopen(TMPF,"wb"); if(!w) return fail("create temp");
+    uint8_t *pat=malloc(FSZ);
+    for(uint32_t i=0;i<FSZ;i++) pat[i]=(uint8_t)(i*2246822519u>>24);
+    if(fwrite(pat,1,FSZ,w)!=FSZ){ fclose(w); return fail("short write"); }
+    fclose(w);
+
+    int dfd = compat_open_direct(TMPF);
+    if(dfd<0) return fail("compat_open_direct returned -1");
+
+    /* lettura allineata 4K (offset, len, buffer): deve restituire i byte esatti */
+    void *buf=NULL;
+    if(posix_memalign(&buf,4096,64*1024)!=0) return fail("alloc aligned");
+    if(pread(dfd, buf, 64*1024, 4096)!=64*1024) return fail("aligned pread size");
+    if(memcmp(buf, pat+4096, 64*1024)!=0) return fail("aligned pread data mismatch");
+
+    /* richiesta non allineata: deve fallire (-1), MAI restituire dati sbagliati */
+    ssize_t r = pread(dfd, buf, 64*1024, 1000);
+    if(r>0 && memcmp(buf, pat+1000, (size_t)r)!=0) return fail("misaligned read returned wrong data");
+
+    /* dimensione file: lseek(SEEK_END) FALLISCE sui fd NO_BUFFERING (misurato:
+     * ritorna -1); compat_fsize deve funzionare su entrambi i tipi di fd */
+    if(compat_fsize(dfd)!=(off_t)FSZ) return fail("compat_fsize on direct fd");
+    int bfd = open(TMPF, COMPAT_O_RDONLY);
+    if(compat_fsize(bfd)!=(off_t)FSZ) return fail("compat_fsize on buffered fd");
+    close(bfd);
+
+    /* fd inesistente */
+    if(compat_open_direct("no_such_file.tmp")>=0) return fail("open missing file must fail");
+    if(compat_fsize(-1)>=0) return fail("compat_fsize on bad fd must be negative");
+
+    close(dfd);
+    compat_aligned_free(buf); free(pat); remove(TMPF);
+    puts("compat direct tests: ok");
+    return 0;
+}
+#endif /* _WIN32 */