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