From d47b0958751205b874e7b65d9f67c432b9ab46d6 Mon Sep 17 00:00:00 2001 From: Tom Olorin <1800017+juicyroots@users.noreply.github.com> Date: Tue, 14 Jul 2026 07:12:15 -0700 Subject: [PATCH] Windows: direct I/O (FILE_FLAG_NO_BUFFERING, 1.47x decode) + VirtualLock pin wiring + select_ctx device cache (#162) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * 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 * 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 * 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 --------- Co-authored-by: olorin Co-authored-by: Claude Fable 5 --- c/Makefile | 5 ++- c/backend_cuda.cu | 12 +++++++- c/compat.h | 49 ++++++++++++++++++++++++++++++ c/glm.c | 5 +++ c/iobench.c | 8 +++++ c/st.h | 4 +-- c/tests/test_compat_direct.c | 59 ++++++++++++++++++++++++++++++++++++ 7 files changed, 138 insertions(+), 4 deletions(-) create mode 100644 c/tests/test_compat_direct.c diff --git a/c/Makefile b/c/Makefile index 797fa1d..e794d4b 100644 --- a/c/Makefile +++ b/c/Makefile @@ -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) diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu index fda663a..d09a9fa 100644 --- a/c/backend_cuda.cu +++ b/c/backend_cuda.cu @@ -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) { diff --git a/c/compat.h b/c/compat.h index bb0b899..721ba8e 100644 --- a/c/compat.h +++ b/c/compat.h @@ -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; diff --git a/c/glm.c b/c/glm.c index 5438e58..a8bcf8d 100644 --- a/c/glm.c +++ b/c/glm.c @@ -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 diff --git a/c/iobench.c b/c/iobench.c index fdc94cb..d3ecb14 100644 --- a/c/iobench.c +++ b/c/iobench.c @@ -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(szpaths[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 diff --git a/c/tests/test_compat_direct.c b/c/tests/test_compat_direct.c new file mode 100644 index 0000000..9b49930 --- /dev/null +++ b/c/tests/test_compat_direct.c @@ -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 + +#ifndef _WIN32 +int main(void){ puts("compat direct tests: skipped (POSIX has native O_DIRECT)"); return 0; } +#else + +#include +#include +#include +#include +#include +#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>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 */