perf: znver5 + LTO + OMP 32T + GPU DENSE/ASYNC (Strix Halo optimiert)

Initial: Colibri upstream + HIP/ROCm port (Strix Halo / Radeon 8060S)
README: GB10 (DGX Spark) row (#136 ) — 0.50 tok/s warm, matmul-bound, unified memory; confirms #76 fix on aarch64 (TF 32/32 + greedy 20/20 on NEON and CUDA sm_121)
2026-07-14 23:56:35 +02:00 · 2026-07-14 23:53:20 +02:00 · 2026-07-13 21:13:40 +02:00 · 2026-07-13 21:00:50 +02:00 · 2026-07-13 20:54:30 +02:00 · 2026-07-13 20:47:23 +02:00
39 changed files with 1353 additions and 95 deletions
@@ -1,51 +1 @@
-# ambienti python / cache
+.coli_kv\n*.o\n*.log\n__pycache__/\n*.pyc\n/tmp/
 c/mio_env/
 __pycache__/
 **/__pycache__/
 *.pyc
 web/node_modules/
 web/dist/
 web/*.tsbuildinfo
 desktop/src-tauri/target/
 desktop/src-tauri/gen/
 # binari compilati (si rigenerano con make / coli build)
 c/glm
 c/olmoe
 c/iobench
 c/tok_test
 c/backend_cuda.o
 c/backend_cuda_test
 c/tests/test_json
 c/tests/test_st
 c/tests/test_tier
 c/tests/test_grammar
 # oracoli tiny generati (make_glm_oracle.py) e dati benchmark scaricati
 c/glm_tiny/
 c/glm_tiny_i2/
 c/glm_tiny_i4/
 c/glm_tiny_mix/
 c/glm_bench_medium/
 c/bench/
 # pesi modello / artefatti di run
 *.safetensors
 stats*.txt
 *.log
 # script di ricerca stadio-0 (esperimenti locali, non parte del motore)
 /engine.py
 /run.py
 /validate_ref.py
 /s0_costmodel.py
 /s1_trace_hitrate.py
 /s2_research.py
 /quant_test.py
 /profile_run.py
 /sweep.py
 c/backend_metal.o
 c/backend_metal_test
 c/tests/test_tier
 c/mio_env/
 c/bench/
@@ -160,6 +160,12 @@ make iobench.exe        # disk I/O benchmark
 make test-c             # run C tests
 make test-python        # run Python tests (requires python)
 # AVX-VNNI: Intel Alder Lake+ (and Meteor Lake+) CPUs have a 128-bit int8
 # dot-product instruction (VPDPBUSD) the engine can use for ~1.3x faster
 # quantized matmul. The x86-64-v3 default (portable AVX2) compiles it out;
 # build for THIS machine to enable it:
 make glm.exe ARCH=native                       # banner prints "idot: avx-vnni"
 # Verify (tiny model, 2.4 MB):
 pip install torch transformers safetensors huggingface_hub
 python tools/make_glm_oracle.py                # generate tiny oracle
@@ -171,9 +177,49 @@ python coli chat --model D:\glm52_i4            # interactive chat
 python coli serve --model D:\glm52_i4            # OpenAI-compatible API
 ```
-**Status:** Phase 1 complete (compiles, correct, static-linked). O_DIRECT (Phase 2),
+**Warmup (overnight cache priming):** the engine's expert cache learns from
-GPU via `LoadLibrary` on `coli_cuda.dll` (Phases G0–G2), and full-model validation
+your workload. The included `warmup.ps1` script runs `coli run` in a loop with
-are separate workstreams. See `PORT_WINDOWS_PLAN.md` for the full plan.
+diverse prompts to build the `.coli_usage` histogram unattended, so the next
 real session starts with a large, accurate hot-expert pin. Each run saves usage
 atomically on clean completion.
 ```powershell
 .\warmup.ps1 -Rounds 1 -Ngen 32               # ~60-90 min, durable progress
 ```
 **NVIDIA GPU (optional, via runtime DLL):** on Windows the engine is built with
 MinGW gcc but CUDA kernels require MSVC + nvcc. The split is clean: build the
 CUDA backend into a standalone `coli_cuda.dll` (nvcc + MSVC), then the host
 `glm.exe` loads it at runtime via `LoadLibrary` (`c/backend_loader.c`). The host
 never links cudart directly; if the DLL is absent the engine falls back to CPU
 without error.
 ```powershell
 # Prerequisites: CUDA Toolkit + MSVC Build Tools (cl.exe) + nvcc on PATH.
 # Build the DLL from a shell with the MSVC environment set (vcvars64.bat or
 # "x64 Native Tools Command Prompt for VS"):
 make cuda-dll CUDA_HOME="C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.8" CUDA_ARCH=sm_120
 # Build the host with the runtime loader (CUDA_DLL=1 adds -DCOLI_CUDA and
 # links backend_loader.o instead of cudart):
 make glm.exe CUDA_DLL=1 ARCH=native
 # Run with the GPU expert tier (8 GB VRAM budget here; scale to your free VRAM):
 $env:COLI_CUDA="1"; $env:COLI_GPU="0"; $env:CUDA_EXPERT_GB="8"
 python coli chat --model D:\glm52_i4 --topp 0.7
 ```
 The DLL exports 11 `extern "C"` symbols (`coli_cuda_init`, `coli_cuda_matmul`,
 etc.); `backend_loader.c` resolves them via `GetProcAddress` on first use.
 `ColiCudaTensor*` is opaque to the host (stored, never dereferenced), so the
 MSVC-allocated struct is safe across the ABI boundary. `CUDA_ARCH` must match
 your GPU's compute capability (e.g. `sm_120` for Blackwell / RTX 50-series,
 `sm_89` for Ada / RTX 40-series).
 **Status:** Phase 1 complete (compiles, correct, static-linked). The Windows
 GPU tier (runtime `coli_cuda.dll` via `LoadLibrary`) is implemented and
 verified on RTX 50-series (sm_120). O_DIRECT (Phase 2) and full-model
 validation against the transformers oracle remain separate workstreams.
 ### OpenAI-compatible API
@@ -496,6 +542,10 @@ Real numbers from real machines, stock build (`setup.sh`, gcc 13), greedy decodi
 | Ryzen 7 9800X3D (16T) · WSL2 · 70 GB RAM · Samsung 9100 PRO PCIe 5.0 · RTX 5090 ([#101](https://github.com/JustVugg/colibri/issues/101)) | **10.51 GB/s** O_DIRECT | MTP off · learned pin 24 GB · hit 54% · OMP hot-team on | **0.41 tok/s** · disk-bound (36.5 s disk vs 24.0 s matmul) · **CUDA expert tier ≈ 0%** (AVX-512 CPU matches the 5090) · `--topp 0.7` → **0.52 tok/s** |
 | EPYC 7443 (24C/48T, Zen3 AVX2) · Linux · **430 GB RAM** · NVMe RAID-Z1 via TrueNAS VM ([#104](https://github.com/JustVugg/colibri/issues/104)) | ~1 GB/s (VM overhead) | 77.5 GB pin · cap auto-raised to 194/layer · MTP off | **1.00 tok/s** · **hit 98%** · disk eliminated → **RAM-bandwidth + matmul bound** (no AVX-512/VNNI on Zen3) |
 | Intel i5-12600K (10C/16T, AVX2) · **native Windows 11, no WSL** · 32 GB · MinGW GCC 16.1 ([#113](https://github.com/JustVugg/colibri/issues/113)) | buffered (no O_DIRECT on MinGW) | int8 MTP head · cold, small-RAM (cap ~2/layer) | **0.08 tok/s** · hit 3.7% · **MTP 57% acceptance** — first native-Windows datapoint, port validated |
 | Ryzen 9 9950X3D2 (16C/32T, avx512-vnni) · native Linux · 121 GB · Samsung 9100 PRO **PCIe Gen5** · RTX 5090 (28 GB expert tier, 1475 pinned) ([#120](https://github.com/JustVugg/colibri/issues/120)) | **11.48 GB/s** O_DIRECT | `MTP=0 DIRECT=1 PIPE_WORKERS=16 PREFETCH=1` | **1.23 tok/s** · MTP-off wins disk-bound · fastest x86 datapoint yet |
 | Ryzen AI Max+ 395 (Strix Halo, 16C/32T Zen5, avx512-vnni) · Arch Linux · 128 GB unified LPDDR5x · SK hynix P41 PCIe 4.0 ([#124](https://github.com/JustVugg/colibri/issues/124)) | — | `DIRECT=1 PIPE=1 --topp 0.7` · auto-pin | 0.06 cold → **1.10 tok/s** sustained · first Strix Halo / gfx1151 datapoint (unified memory: no discrete VRAM tier) |
 | Intel Core Ultra 9 185H (16C/22T, avx-vnni) · **native Windows 11, no WSL** · 32 GB · Crucial P3 QLC NTFS · RTX 5070 Ti (unused) ([#128](https://github.com/JustVugg/colibri/issues/128)) | — | int8 MTP head · **with [#131](https://github.com/JustVugg/colibri/pull/131) (pipe + RAM fixes), warm cache, no GPU** | 0.03 cold → **0.5 tok/s** warm (~7-prompt warmup) · cache-warming on native Windows once the portability blockers are fixed — stock main hung on the `\r\n` READY sentinel before #131 |
 | Dell Pro Max GB10 (DGX Spark: Grace 10×X925 + 10×A725, **aarch64 i8mm/sve2**) · Linux · 121 GB unified LPDDR5x · Dell OEM 4 TB NVMe · GB10 sm_121 ([#136](https://github.com/JustVugg/colibri/issues/136)) | **5.58 GB/s** O_DIRECT (NVIDIA-OEM unit in #76 was 10.74 — same platform, different SSD) | int8 MTP head · warm cache | 0.21 cold → **0.50 tok/s** warm · hit 83% · MTP 73% (3.20 tok/fw) · **matmul-bound** (matmul 130 s vs disk 58 s) — unified memory, CUDA placement tier neutral; the lever here is an i8mm compute kernel, not placement |
 Takeaways: with 24 GB of RAM the engine auto-caps the expert cache to 2 slots/layer, so decode stays cold even on a disk 2–2.7× faster than the dev box — **on small-RAM machines the RAM cap, not the disk, is the binding constraint**, exactly as the table above predicts; `--topp 0.7` alone bought a clean 1.6× end-to-end speedup. The M5 Max datapoint lands right on the table's second row: **~1 tok/s of a 744B model on a laptop SSD** — and its 14 GB/s disk shifts the bottleneck back to RAM budget and kernels. The Framework 13 rows are the cache thesis proven end-to-end on one machine: 0.29 → 0.37 tok/s (hit 28% → 66%, speculation finally engaging at 52% acceptance) just by giving the cache its RAM — int8 MTP head + a bigger cap + the learned pin. The cap part is now automatic (cap auto-raise, 2026-07-10). The 9950X pair is the cleanest bottleneck experiment yet — same machine, same history, only the disk swapped: ×5.8 disk bandwidth bought ×2.9 tokens, and the profile **flipped from 66% disk to 57% matmul**. But the crossover depends on the CPU kernel: the 9800X3D row ([#101](https://github.com/JustVugg/colibri/issues/101)) shows that with the OMP hot-team tuning on, the AVX-512 CPU matmul is fast enough that even a **10 GB/s NVMe stays disk-bound** — and there the **CUDA expert tier buys ≈ 0%**, because the CPU already matches the 5090 on expert matmul. The GPU tier earns its VRAM only when the CPU is the weak link, not by default. (Honest correction from #101: an earlier version of that report ran with the OMP tuning off, which manufactured a false matmul-bound crossover and a false +14% for CUDA — neither survived a clean re-run.)
@@ -1,7 +1,18 @@
 # --- OS detection ---
 # On Windows the OS env var is 'Windows_NT' in EVERY shell (cmd, PowerShell,
 # MSYS2, Git-Bash), and GNU Make imports it. Detect Windows from it FIRST so
 # `make` works from a native PowerShell/CMD shell — there `uname` is not on
 # PATH, so the old uname-only check returned empty and fell through to the
 # Linux branch (no .exe, no -static). We only call `uname` when NOT on Windows.
 ifeq ($(OS),Windows_NT)
 IS_WIN   := 1
 UNAME_S  :=
 UNAME_M  :=
 else
 IS_WIN   :=
 UNAME_S  := $(shell uname -s)
-MINGW    := $(findstring MINGW,$(UNAME_S))
+UNAME_M  := $(shell uname -m)
-MSYS     := $(findstring MSYS,$(UNAME_S))
+endif
 IS_WIN   := $(MINGW)$(MSYS)
 ifeq ($(UNAME_S),Darwin)
 # --- macOS / Apple Silicon ---
@@ -24,9 +35,12 @@ EXE     =
 else ifneq ($(IS_WIN),)
 # --- Windows 11 x86-64 (MinGW-w64 / MSYS2) ---
 # GCC + libgomp + winpthreads: pthread, OpenMP, clock_gettime, opendir/readdir,
-# AVX2 intrinsics — tutto gratis, nessun porting.
+# AVX2 intrinsics - tutto gratis, nessun porting.
-# ARCH default = x86-64-v3 (binario portabile con AVX2). Per max velocita'
+# ARCH default = x86-64-v3 (portable binary with AVX2). For max speed on THIS
-# su QUESTA macchina: make ARCH=native
+# machine use ARCH=native: on AVX-VNNI CPUs (Intel Alder Lake+, Meteor Lake+)
 # it also unlocks the 128-bit VPDPBUSD int8/int4 dot kernel (dot_i8i8/dot_i4i8),
 # which the x86-64-v3 baseline does not define. The #ifdef guards in glm.c mean
 # a v3 build simply compiles out the VNNI path - safe on any x86-64.
 CC      = gcc
 ARCH   ?= x86-64-v3
 CFLAGS  = -D_FILE_OFFSET_BITS=64 -O3 -march=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function
@@ -59,7 +73,17 @@ endif
 # CUDA=1 adds an opt-in backend for resident tensors. The default build remains
 # pure C and keeps the original zero-dependency runtime.
 #
 # Two paths:
 #   - Linux/macOS: CUDA=1 links backend_cuda.o directly (cudart via -l).
 #   - Windows:      CUDA_DLL=1 builds a standalone coli_cuda.dll (nvcc+MSVC),
 #                   then the host glm.exe loads it at runtime via backend_loader.c
 #                   (LoadLibrary/GetProcAddress). MinGW gcc cannot compile .cu
 #                   (nvcc needs cl.exe), and cross-linking MSVC objects into a
 #                   gcc binary is fragile — the DLL split keeps the toolchains
 #                   clean. See backend_loader.c and README "cuda-dll" below.
 CUDA      ?= 0
 CUDA_DLL  ?= 0
 CUDA_HOME ?= /usr/local/cuda
 NVCC      ?= $(CUDA_HOME)/bin/nvcc
 CUDA_ARCH ?= native
@@ -67,13 +91,23 @@ NVCCFLAGS ?= -O3 -std=c++17 -arch=$(CUDA_ARCH) -Xcompiler=-Wall,-Wextra
 PYTHON    ?= python3
 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)
 # Windows CUDA DLL path: host links the loader, NOT cudart.
 ifneq ($(IS_WIN),)
 ifeq ($(CUDA_DLL),1)
 CFLAGS  += -DCOLI_CUDA
 CUDA_OBJ = backend_loader.o
 endif
 endif
 # Linux CUDA direct-link path (unchanged).
 ifeq ($(CUDA),1)
 ifeq ($(UNAME_S),Darwin)
 $(error CUDA=1 is supported only on Linux)
 endif
 ifneq ($(IS_WIN),)
-# GPU: stub only in Phase 1 (G0). G1 builds coli_cuda.dll with MSVC+nvcc.
+# On Windows use CUDA_DLL=1 (runtime DLL), not CUDA=1 (direct link).
-$(error CUDA=1 on Windows requires G1: build coli_cuda.dll with MSVC+nvcc (see PORT_WINDOWS_PLAN.md §8))
+$(error On Windows use: make CUDA_DLL=1 cuda-dll  (see backend_loader.c))
 endif
 CFLAGS   += -DCOLI_CUDA
 LDFLAGS  += -L$(CUDA_HOME)/lib64 -Wl,-rpath,$(CUDA_HOME)/lib64 -lcudart -lstdc++
@@ -102,6 +136,22 @@ glm: glm$(EXE)
 glm$(EXE): glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h $(CUDA_OBJ) $(METAL_OBJ)
 	$(CC) $(CFLAGS) glm.c $(CUDA_OBJ) $(METAL_OBJ) -o glm$(EXE) $(LDFLAGS)
 # Windows runtime loader object: resolves coli_cuda_* from coli_cuda.dll.
 backend_loader.o: backend_loader.c backend_cuda.h compat.h
 	$(CC) $(CFLAGS) -c backend_loader.c -o $@
 # Windows CUDA DLL: compile backend_cuda.cu with nvcc (+MSVC cl.exe as host
 # compiler, required by nvcc on Windows) into coli_cuda.dll. Run this from a
 # shell that has the MSVC environment set (e.g. after vcvars64.bat, or from a
 # "x64 Native Tools Command Prompt"). COLI_CUDA_BUILDING_DLL enables
 # __declspec(dllexport) so the 11 API symbols are exported.
 cuda-dll: backend_cuda.cu backend_cuda.h
 	@command -v $(NVCC) >/dev/null 2>&1 || { echo "nvcc not found: set CUDA_HOME or NVCC" >&2; exit 1; }
 	@command -v cl >/dev/null 2>&1 || { echo "cl.exe (MSVC) not in PATH — run vcvars64.bat first" >&2; exit 1; }
 	$(NVCC) $(NVCCFLAGS) -shared -DCOLI_CUDA_BUILDING_DLL \
 		-L"$(CUDA_HOME)/lib/x64" -lcudart \
 		backend_cuda.cu -o coli_cuda.dll
 backend_cuda.o: backend_cuda.cu backend_cuda.h
 	@command -v $(NVCC) >/dev/null 2>&1 || { echo "nvcc not found: set CUDA_HOME or NVCC" >&2; exit 1; }
 	$(NVCC) $(NVCCFLAGS) -c backend_cuda.cu -o $@
@@ -169,7 +219,7 @@ check:
 	$(MAKE) test
 clean:
-	rm -f olmoe$(EXE) glm$(EXE) iobench$(EXE) backend_cuda.o backend_cuda_test$(EXE) backend_cuda_bench$(EXE) backend_metal.o backend_metal_test $(TEST_BINS)
+	rm -f olmoe$(EXE) glm$(EXE) iobench$(EXE) backend_cuda.o backend_loader.o backend_cuda_test$(EXE) backend_cuda_bench$(EXE) backend_metal.o backend_metal_test coli_cuda.dll coli_cuda.lib coli_cuda.exp $(TEST_BINS)
 	rm -rf tests/__pycache__
-.PHONY: all glm cuda-test cuda-bench portable test-c test-python test check clean
+.PHONY: all glm cuda-test cuda-bench cuda-dll portable test-c test-python test check clean
@@ -4,6 +4,16 @@
 #include <stddef.h>
 #include <stdint.h>
 /* COLI_CUDA_DLLEXPORT marks functions exported from coli_cuda.dll on Windows.
 * Define COLI_CUDA_BUILDING_DLL when compiling the .cu into the DLL (so the
 * functions are __declspec(dllexport)); the host loader does NOT include this
 * header's declarations — it resolves symbols at runtime via GetProcAddress. */
 #if defined(_WIN32) && defined(COLI_CUDA_BUILDING_DLL)
 #define COLI_CUDA_DLLEXPORT __declspec(dllexport)
 #else
 #define COLI_CUDA_DLLEXPORT
 #endif
 #ifdef __cplusplus
 extern "C" {
 #endif
@@ -14,18 +24,18 @@ extern "C" {
 typedef struct ColiCudaTensor ColiCudaTensor;
 /* Devices are CUDA ordinals, not positions in the input list. */
-int coli_cuda_init(const int *devices, int count);
+COLI_CUDA_DLLEXPORT int coli_cuda_init(const int *devices, int count);
-void coli_cuda_shutdown(void);
+COLI_CUDA_DLLEXPORT void coli_cuda_shutdown(void);
-int coli_cuda_device_count(void);
+COLI_CUDA_DLLEXPORT int coli_cuda_device_count(void);
-int coli_cuda_device_at(int index);
+COLI_CUDA_DLLEXPORT int coli_cuda_device_at(int index);
-int coli_cuda_mem_info(int device, size_t *free_bytes, size_t *total_bytes);
+COLI_CUDA_DLLEXPORT int coli_cuda_mem_info(int device, size_t *free_bytes, size_t *total_bytes);
 /* device < 0 returns aggregate statistics for all configured devices. */
-void coli_cuda_stats(int device, size_t *tensor_count, size_t *tensor_bytes);
+COLI_CUDA_DLLEXPORT void coli_cuda_stats(int device, size_t *tensor_count, size_t *tensor_bytes);
-void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64_t *rows,
+COLI_CUDA_DLLEXPORT void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64_t *rows,
                           double *h2d_ms, double *kernel_ms, double *d2h_ms);
 /* Upload without executing, so capacity failures happen during model startup. */
-int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
+COLI_CUDA_DLLEXPORT int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
                            const void *weights, const float *scales,
                            int fmt, int I, int O, int device);
@@ -35,7 +45,7 @@ int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
 * The first successful call uploads W and its row scales; later calls reuse it.
 * Returns 1 on success and 0 when CUDA is not initialized or the format is invalid.
 */
-int coli_cuda_matmul(ColiCudaTensor **tensor,
+COLI_CUDA_DLLEXPORT int coli_cuda_matmul(ColiCudaTensor **tensor,
                     float *y, const float *x,
                     const void *weights, const float *scales,
                     int fmt, int S, int I, int O, int device);
@@ -43,25 +53,25 @@ int coli_cuda_matmul(ColiCudaTensor **tensor,
 /* Fused expert pipeline: y = down(silu(gate(x)) * up(x)).  All three tensors
 * must already be resident on one device.  Activations cross PCIe once in
 * each direction instead of once per matrix. */
-int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
+COLI_CUDA_DLLEXPORT int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
                         ColiCudaTensor *down, float *y, const float *x, int S);
 /* Packed group of same-shaped experts. Inputs and outputs contain sum(rows)
 * consecutive [D] rows in call order. */
-int coli_cuda_expert_group(ColiCudaTensor *const *gates,
+COLI_CUDA_DLLEXPORT int coli_cuda_expert_group(ColiCudaTensor *const *gates,
                           ColiCudaTensor *const *ups,
                           ColiCudaTensor *const *downs,
                           const int *rows, int count,
                           float *y, const float *x);
 /* Decode-only MLA weight-absorption core for one token. kv_b is [H*(Q+V),K]. */
-int coli_cuda_attention_absorb(ColiCudaTensor *kv_b,float *ctx,const float *q,
+COLI_CUDA_DLLEXPORT int coli_cuda_attention_absorb(ColiCudaTensor *kv_b,float *ctx,const float *q,
                               const float *latent,const float *rope,int H,int Q,
                               int R,int V,int K,int T,float attention_scale);
-void coli_cuda_tensor_free(ColiCudaTensor *tensor);
+COLI_CUDA_DLLEXPORT void coli_cuda_tensor_free(ColiCudaTensor *tensor);
-size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor);
+COLI_CUDA_DLLEXPORT size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor);
-int coli_cuda_tensor_device(const ColiCudaTensor *tensor);
+COLI_CUDA_DLLEXPORT int coli_cuda_tensor_device(const ColiCudaTensor *tensor);
 #ifdef __cplusplus
 }
@@ -0,0 +1,565 @@
 #include "backend_cuda.h"
 #include <hip/hip_runtime.h>
 // HIP: mma.h not available, TC_INT4 disabled via env
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <mutex>
 struct ColiCudaTensor {
    void *weights;
    float *scales;
    size_t weight_bytes;
    int fmt, I, O, device;
    int tracked;
 };
 typedef struct {
    int device;
    float *x, *y, *gate, *up;
    size_t x_cap, y_cap, gate_cap, up_cap;
    uint8_t *qx; float *qscale;
    size_t qx_cap, qscale_cap;
    float *host_x,*host_y; size_t host_x_cap,host_y_cap;
    float *aq,*al,*ar,*ac; size_t aq_cap,al_cap,ar_cap,ac_cap;
    hipStream_t stream;
    void *group_desc; size_t group_desc_cap;
    size_t tensor_count, tensor_bytes;
 } DeviceContext;
 typedef struct {
    const void *g,*u,*d; const float *gs,*us,*ds;
    int gf,uf,df,rows,offset;
 } GroupDesc;
 static DeviceContext g_ctx[COLI_CUDA_MAX_DEVICES];
 static int g_nctx;
 static uint64_t g_group_calls,g_group_experts,g_group_rows;
 static double g_group_h2d_ms,g_group_kernel_ms,g_group_d2h_ms;
 static std::mutex g_group_stats_mu;
 static int cuda_ok(hipError_t err, const char *what) {
    if (err == hipSuccess) return 1;
    std::fprintf(stderr, "[CUDA] %s: %s\n", what, hipGetErrorString(err));
    return 0;
 }
 static DeviceContext *find_ctx(int device) {
    for (int i = 0; i < g_nctx; i++) if (g_ctx[i].device == device) return &g_ctx[i];
    return nullptr;
 }
 static int select_ctx(DeviceContext *ctx) {
    return ctx && cuda_ok(hipSetDevice(ctx->device), "select device");
 }
 __host__ __device__ static size_t row_bytes(int fmt, int I) {
    if (fmt == 0) return (size_t)I * sizeof(float);
    if (fmt == 1) return (size_t)I;
    if (fmt == 2) return (size_t)(I + 1) / 2;
    if (fmt == 3) return (size_t)(I + 3) / 4;
    return 0;
 }
 __device__ static float weight_at(const void *weights, int fmt, size_t row, int i) {
    const uint8_t *base = static_cast<const uint8_t *>(weights) + row;
    if (fmt == 0) return reinterpret_cast<const float *>(base)[i];
    if (fmt == 1) return static_cast<float>(reinterpret_cast<const int8_t *>(base)[i]);
    const uint8_t *q = base;
    if (fmt == 2) {
        uint8_t v = q[i >> 1];
        int n=(i&1)?(v>>4):(v&15); return static_cast<float>(n&8?n-16:n);
    }
    uint8_t v = q[i >> 2];
    return static_cast<float>(((v >> ((i & 3) * 2)) & 3) - 2);
 }
 __global__ static void offset_to_signed_s4(uint8_t *q,size_t n){
    size_t i=(size_t)blockIdx.x*blockDim.x+threadIdx.x;if(i<n)q[i]^=0x88;
 }
 __global__ static void quant_matmul(float *y, const float *x, const void *weights,
                                    const float *scales, int fmt, int S, int I, int O,
                                    size_t rb) {
    int o = blockIdx.x;
    int s = blockIdx.y;
    float sum = 0.0f;
    size_t row = (size_t)o * rb;
    const float *xs = x + (size_t)s * I;
    for (int i = threadIdx.x; i < I; i += blockDim.x)
        sum += xs[i] * weight_at(weights, fmt, row, i);
    __shared__ float partial[256];
    partial[threadIdx.x] = sum;
    __syncthreads();
    for (int n = blockDim.x >> 1; n; n >>= 1) {
        if (threadIdx.x < n) partial[threadIdx.x] += partial[threadIdx.x + n];
        __syncthreads();
    }
    if (!threadIdx.x)
        y[(size_t)s * O + o] = partial[0] * (fmt ? scales[o] : 1.0f);
 }
 __global__ static void silu_mul(float *gate, const float *up, size_t n) {
    size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        float v = gate[i];
        gate[i] = (v / (1.0f + expf(-v))) * up[i];
    }
 }
 __global__ static void quantize_s4_rows(uint8_t *q,float *scale,const float *x,int S,int K){
    int s=blockIdx.x; if(s>=S)return; const float *xs=x+(size_t)s*K;
    float v=0; for(int i=threadIdx.x;i<K;i+=blockDim.x)v=fmaxf(v,fabsf(xs[i]));
    __shared__ float m[256]; m[threadIdx.x]=v; __syncthreads();
    for(int n=128;n;n>>=1){if(threadIdx.x<n)m[threadIdx.x]=fmaxf(m[threadIdx.x],m[threadIdx.x+n]);__syncthreads();}
    float sc=m[0]>0?m[0]/7.f:1.f; if(!threadIdx.x)scale[s]=sc;
    uint8_t *dst=q+(size_t)s*((K+1)/2);
    for(int b=threadIdx.x;b<(K+1)/2;b+=blockDim.x){
        int i=b*2,a=__float2int_rn(xs[i]/sc),c=i+1<K?__float2int_rn(xs[i+1]/sc):0;
        a=max(-8,min(7,a)); c=max(-8,min(7,c)); dst[b]=(uint8_t)((a&15)|((c&15)<<4));
    }
 }
 __global__ static void grouped_s4_wmma(float *y,const uint8_t *x,const float *xscale,
                                        const GroupDesc *desc,int K,int O,int which){
 #if __CUDA_ARCH__ >= 750
    // HIP: nvcuda removed
    int warp=threadIdx.x/32,lane=threadIdx.x%32,tile=blockIdx.x*8+warp,c=blockIdx.y;
    if(tile*8>=O)return; GroupDesc d=desc[c];
    const void *w=which==0?d.g:(which==1?d.u:d.d);
    const float *ws=which==0?d.gs:(which==1?d.us:d.ds);
    int fmt=which==0?d.gf:(which==1?d.uf:d.df);
    if(fmt!=2)return;
    wmma::fragment<wmma::accumulator,8,8,32,int> acc; wmma::fill_fragment(acc,0);
    const uint8_t *a=x+(size_t)d.offset*((K+1)/2);
    const uint8_t *b=(const uint8_t*)w+(size_t)(tile*8)*((K+1)/2);
    for(int k=0;k<K;k+=32){
        wmma::fragment<wmma::matrix_a,8,8,32,wmma::experimental::precision::s4,wmma::row_major> af;
        wmma::fragment<wmma::matrix_b,8,8,32,wmma::experimental::precision::s4,wmma::col_major> bf;
        wmma::load_matrix_sync(af,a+k/2,K);
        wmma::load_matrix_sync(bf,b+k/2,K);
        wmma::mma_sync(acc,af,bf,acc);
    }
    __shared__ int out[8][64]; wmma::store_matrix_sync(out[warp],acc,8,wmma::mem_row_major);
    for(int i=lane;i<64;i+=32){int s=i/8,o=tile*8+i%8;
        if(s<d.rows&&o<O)y[(size_t)(d.offset+s)*O+o]=(float)out[warp][i]*xscale[d.offset+s]*ws[o];}
 #endif
 }
 __global__ static void grouped_hidden(float *y,const float *x,const GroupDesc *desc,
                                      int I,int D,int which){
    int o=blockIdx.x,s=blockIdx.y,c=blockIdx.z; GroupDesc d=desc[c];
    if(s>=d.rows) return;
    const void *w=which?d.u:d.g; const float *sc=which?d.us:d.gs; int fmt=which?d.uf:d.gf;
    size_t rb=row_bytes(fmt,D),row=(size_t)o*rb; const float *xs=x+(size_t)(d.offset+s)*D;
    float sum=0; for(int i=threadIdx.x;i<D;i+=blockDim.x) sum+=xs[i]*weight_at(w,fmt,row,i);
    __shared__ float p[256]; p[threadIdx.x]=sum; __syncthreads();
    for(int n=128;n;n>>=1){ if(threadIdx.x<n)p[threadIdx.x]+=p[threadIdx.x+n]; __syncthreads(); }
    if(!threadIdx.x) y[(size_t)(d.offset+s)*I+o]=p[0]*(fmt?sc[o]:1.f);
 }
 __global__ static void grouped_down(float *y,const float *x,const GroupDesc *desc,int D,int I){
    int o=blockIdx.x,s=blockIdx.y,c=blockIdx.z; GroupDesc d=desc[c];
    if(s>=d.rows) return;
    size_t rb=row_bytes(d.df,I),row=(size_t)o*rb; const float *xs=x+(size_t)(d.offset+s)*I;
    float sum=0; for(int i=threadIdx.x;i<I;i+=blockDim.x) sum+=xs[i]*weight_at(d.d,d.df,row,i);
    __shared__ float p[256]; p[threadIdx.x]=sum; __syncthreads();
    for(int n=128;n;n>>=1){ if(threadIdx.x<n)p[threadIdx.x]+=p[threadIdx.x+n]; __syncthreads(); }
    if(!threadIdx.x) y[(size_t)(d.offset+s)*D+o]=p[0]*(d.df?d.ds[o]:1.f);
 }
 __device__ static void unpack_s4(uint8_t v,float *lo,float *hi){
    int a=v&15,b=v>>4; *lo=(float)(a&8?a-16:a); *hi=(float)(b&8?b-16:b);
 }
 /* Exact low-row W4A32 path. It consumes each packed weight byte once instead
 * of routing both nibbles through weight_at(), preserving FP32 activations. */
 __global__ static void grouped_hidden_w4(float *y,const float *x,const GroupDesc *desc,
                                         int I,int D,int which){
    int o=blockIdx.x,s=blockIdx.y,c=blockIdx.z;GroupDesc d=desc[c];if(s>=d.rows)return;
    const uint8_t *w=(const uint8_t*)(which?d.u:d.g);const float *sc=which?d.us:d.gs;
    const uint8_t *row=w+(size_t)o*((D+1)/2);const float *xs=x+(size_t)(d.offset+s)*D;
    float sum=0;for(int b=threadIdx.x;b<(D+1)/2;b+=blockDim.x){float a,z;unpack_s4(row[b],&a,&z);
        int i=b*2;sum+=xs[i]*a;if(i+1<D)sum+=xs[i+1]*z;}
    __shared__ float p[256];p[threadIdx.x]=sum;__syncthreads();
    for(int n=128;n;n>>=1){if(threadIdx.x<n)p[threadIdx.x]+=p[threadIdx.x+n];__syncthreads();}
    if(!threadIdx.x)y[(size_t)(d.offset+s)*I+o]=p[0]*sc[o];
 }
 __global__ static void grouped_hidden_w4_dual(float *gate,float *up,const float *x,
                                               const GroupDesc *desc,int I,int D){
    int o=blockIdx.x,s=blockIdx.y,c=blockIdx.z;GroupDesc d=desc[c];if(s>=d.rows)return;
    const uint8_t *gr=(const uint8_t*)d.g+(size_t)o*((D+1)/2);
    const uint8_t *ur=(const uint8_t*)d.u+(size_t)o*((D+1)/2);
    const float *xs=x+(size_t)(d.offset+s)*D;float ga=0,ua=0;
    for(int b=threadIdx.x;b<(D+1)/2;b+=blockDim.x){float g0,g1,u0,u1;unpack_s4(gr[b],&g0,&g1);unpack_s4(ur[b],&u0,&u1);
        int i=b*2;ga+=xs[i]*g0;ua+=xs[i]*u0;if(i+1<D){ga+=xs[i+1]*g1;ua+=xs[i+1]*u1;}}
    __shared__ float gp[256],upv[256];gp[threadIdx.x]=ga;upv[threadIdx.x]=ua;__syncthreads();
    for(int n=128;n;n>>=1){if(threadIdx.x<n){gp[threadIdx.x]+=gp[threadIdx.x+n];upv[threadIdx.x]+=upv[threadIdx.x+n];}__syncthreads();}
    if(!threadIdx.x){size_t z=(size_t)(d.offset+s)*I+o;gate[z]=gp[0]*d.gs[o];up[z]=upv[0]*d.us[o];}
 }
 __global__ static void grouped_down_w4(float *y,const float *x,const GroupDesc *desc,int D,int I){
    int o=blockIdx.x,s=blockIdx.y,c=blockIdx.z;GroupDesc d=desc[c];if(s>=d.rows)return;
    const uint8_t *row=(const uint8_t*)d.d+(size_t)o*((I+1)/2);
    const float *xs=x+(size_t)(d.offset+s)*I;float sum=0;
    for(int b=threadIdx.x;b<(I+1)/2;b+=blockDim.x){float a,z;unpack_s4(row[b],&a,&z);
        int i=b*2;sum+=xs[i]*a;if(i+1<I)sum+=xs[i+1]*z;}
    __shared__ float p[256];p[threadIdx.x]=sum;__syncthreads();
    for(int n=128;n;n>>=1){if(threadIdx.x<n)p[threadIdx.x]+=p[threadIdx.x+n];__syncthreads();}
    if(!threadIdx.x)y[(size_t)(d.offset+s)*D+o]=p[0]*d.ds[o];
 }
 __global__ static void attention_absorb_kernel(float *ctx,const float *q,const float *latent,
                                                const float *rope,const void *weights,const float *wscale,
                                                int fmt,int H,int Q,int R,int V,int K,int T,float scale){
    int h=blockIdx.x,tid=threadIdx.x,rbase=h*(Q+V);extern __shared__ float sm[];
    float *qa=sm,*cl=qa+K,*scores=cl+K;
    for(int k=tid;k<K;k+=blockDim.x){float a=0;for(int d=0;d<Q;d++)
        a+=q[(size_t)h*(Q+R)+d]*weight_at(weights,fmt,(size_t)(rbase+d)*row_bytes(fmt,K),k)*(fmt?wscale[rbase+d]:1.f);qa[k]=a;}
    __syncthreads();
    for(int t=tid;t<T;t+=blockDim.x){float a=0;const float *lt=latent+(size_t)t*K,*rt=rope+(size_t)t*R;
        for(int k=0;k<K;k++)a+=qa[k]*lt[k];for(int d=0;d<R;d++)a+=q[(size_t)h*(Q+R)+Q+d]*rt[d];scores[t]=a*scale;}
    __syncthreads();
    if(!tid){float mx=scores[0];for(int t=1;t<T;t++)mx=fmaxf(mx,scores[t]);float z=0;
        for(int t=0;t<T;t++){scores[t]=expf(scores[t]-mx);z+=scores[t];}for(int t=0;t<T;t++)scores[t]/=z;}
    __syncthreads();
    for(int k=tid;k<K;k+=blockDim.x){float a=0;for(int t=0;t<T;t++)a+=scores[t]*latent[(size_t)t*K+k];cl[k]=a;}
    __syncthreads();
    for(int v=tid;v<V;v+=blockDim.x){int row=rbase+Q+v;float a=0;size_t rb=row_bytes(fmt,K);
        for(int k=0;k<K;k++)a+=cl[k]*weight_at(weights,fmt,(size_t)row*rb,k);ctx[(size_t)h*V+v]=a*(fmt?wscale[row]:1.f);}
 }
 static int reserve(float **ptr, size_t *cap, size_t bytes) {
    if (*cap >= bytes) return 1;
    if (*ptr) hipFree(*ptr);
    *ptr = nullptr;
    *cap = 0;
    if (!cuda_ok(hipMalloc(ptr, bytes), "scratch allocation")) return 0;
    *cap = bytes;
    return 1;
 }
 static int reserve_bytes(void **ptr,size_t *cap,size_t bytes){
    if(*cap>=bytes) return 1; if(*ptr) hipFree(*ptr); *ptr=nullptr; *cap=0;
    if(!cuda_ok(hipMalloc(ptr,bytes),"descriptor allocation")) return 0; *cap=bytes; return 1;
 }
 static int reserve_pinned(float **ptr,size_t *cap,size_t bytes){
    if(*cap>=bytes)return 1;if(*ptr)hipFreeHost(*ptr);*ptr=nullptr;*cap=0;
    if(!cuda_ok(hipMallocHost((void**)ptr,bytes),"pinned staging allocation"))return 0;*cap=bytes;return 1;
 }
 extern "C" int coli_cuda_init(const int *devices, int count) {
    int available = 0;
    if (!devices || count < 1 || count > COLI_CUDA_MAX_DEVICES) return 0;
    if (!cuda_ok(hipGetDeviceCount(&available), "device discovery")) return 0;
    g_nctx = 0;
    for (int i = 0; i < count; i++) {
        int device = devices[i];
        if (device < 0 || device >= available) {
            std::fprintf(stderr, "[CUDA] invalid device %d (available: 0..%d)\n", device, available - 1);
            g_nctx = 0;
            return 0;
        }
        if (find_ctx(device)) {
            std::fprintf(stderr, "[CUDA] duplicate device %d\n", device);
            g_nctx = 0;
            return 0;
        }
        DeviceContext *ctx = &g_ctx[g_nctx];
        *ctx = {};
        ctx->device = device;
        if (!select_ctx(ctx)) { g_nctx = 0; return 0; }
        hipDeviceProp_t prop{};
        if (!cuda_ok(hipGetDeviceProperties(&prop, device), "device properties")) { g_nctx = 0; return 0; }
        if(!cuda_ok(hipStreamCreateWithFlags(&ctx->stream,hipStreamNonBlocking),"stream creation")){
            g_nctx=0;return 0;
        }
        g_nctx++;
        std::fprintf(stderr, "[CUDA] device %d: %s, %.1f GB VRAM, sm_%d%d\n",
                     device, prop.name, prop.totalGlobalMem / 1e9, prop.major, prop.minor);
    }
    return 1;
 }
 extern "C" void coli_cuda_shutdown(void) {
    for (int i = 0; i < g_nctx; i++) {
        DeviceContext *ctx = &g_ctx[i];
        if (!select_ctx(ctx)) continue;
        if (ctx->x) hipFree(ctx->x);
        if (ctx->y) hipFree(ctx->y);
        if (ctx->gate) hipFree(ctx->gate);
        if (ctx->up) hipFree(ctx->up);
        if (ctx->qx) hipFree(ctx->qx);
        if (ctx->qscale) hipFree(ctx->qscale);
        if(ctx->aq)hipFree(ctx->aq);if(ctx->al)hipFree(ctx->al);if(ctx->ar)hipFree(ctx->ar);if(ctx->ac)hipFree(ctx->ac);
        if (ctx->host_x) hipFreeHost(ctx->host_x);
        if (ctx->host_y) hipFreeHost(ctx->host_y);
        if (ctx->stream) hipStreamDestroy(ctx->stream);
        if (ctx->group_desc) hipFree(ctx->group_desc);
        ctx->x = ctx->y = ctx->gate = ctx->up = nullptr;
        ctx->qx=nullptr; ctx->qscale=nullptr;
        ctx->aq=ctx->al=ctx->ar=ctx->ac=nullptr;
        ctx->host_x=ctx->host_y=nullptr;ctx->stream=nullptr;
        ctx->x_cap = ctx->y_cap = ctx->gate_cap = ctx->up_cap = 0;
        ctx->qx_cap=ctx->qscale_cap=0;
        ctx->aq_cap=ctx->al_cap=ctx->ar_cap=ctx->ac_cap=0;
        ctx->host_x_cap=ctx->host_y_cap=0;
        ctx->group_desc=nullptr; ctx->group_desc_cap=0;
    }
    g_nctx = 0;
 }
 extern "C" int coli_cuda_device_count(void) { return g_nctx; }
 extern "C" int coli_cuda_device_at(int index) {
    return index >= 0 && index < g_nctx ? g_ctx[index].device : -1;
 }
 extern "C" int coli_cuda_mem_info(int device, size_t *free_bytes, size_t *total_bytes) {
    DeviceContext *ctx = find_ctx(device);
    if (!free_bytes || !total_bytes || !select_ctx(ctx)) return 0;
    return cuda_ok(hipMemGetInfo(free_bytes, total_bytes), "memory info");
 }
 extern "C" void coli_cuda_stats(int device, size_t *tensor_count, size_t *tensor_bytes) {
    size_t count = 0, bytes = 0;
    for (int i = 0; i < g_nctx; i++) if (device < 0 || g_ctx[i].device == device) {
        count += g_ctx[i].tensor_count;
        bytes += g_ctx[i].tensor_bytes;
    }
    if (tensor_count) *tensor_count = count;
    if (tensor_bytes) *tensor_bytes = bytes;
 }
 extern "C" void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64_t *rows,
                                        double *h2d_ms, double *kernel_ms, double *d2h_ms) {
    if(calls) *calls=g_group_calls; if(experts) *experts=g_group_experts; if(rows) *rows=g_group_rows;
    if(h2d_ms) *h2d_ms=g_group_h2d_ms; if(kernel_ms) *kernel_ms=g_group_kernel_ms;
    if(d2h_ms) *d2h_ms=g_group_d2h_ms;
 }
 extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
                                        const void *weights, const float *scales,
                                        int fmt, int I, int O, int device) {
    DeviceContext *ctx = find_ctx(device);
    if (!tensor || !weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0;
    size_t rb = row_bytes(fmt, I);
    if (!rb || (fmt && !scales)) return 0;
    if (*tensor) {
        ColiCudaTensor *t = *tensor;
        return t->fmt == fmt && t->I == I && t->O == O && t->device == device;
    }
    ColiCudaTensor *t = static_cast<ColiCudaTensor *>(std::calloc(1, sizeof(*t)));
    if (!t) return 0;
    t->fmt = fmt; t->I = I; t->O = O; t->device = device; t->weight_bytes = rb * (size_t)O;
    if (!cuda_ok(hipMalloc(&t->weights, t->weight_bytes), "tensor allocation") ||
        !cuda_ok(hipMemcpy(t->weights, weights, t->weight_bytes, hipMemcpyHostToDevice), "tensor upload")) {
        coli_cuda_tensor_free(t);
        return 0;
    }
    if(fmt==2){offset_to_signed_s4<<<(unsigned)((t->weight_bytes+255)/256),256>>>((uint8_t*)t->weights,t->weight_bytes);
        if(!cuda_ok(hipGetLastError(),"int4 weight conversion")){coli_cuda_tensor_free(t);return 0;}}
    if (fmt) {
        if (!cuda_ok(hipMalloc(&t->scales, (size_t)O * sizeof(float)), "scale allocation") ||
            !cuda_ok(hipMemcpy(t->scales, scales, (size_t)O * sizeof(float), hipMemcpyHostToDevice), "scale upload")) {
            coli_cuda_tensor_free(t);
            return 0;
        }
    }
    t->tracked = 1;
    ctx->tensor_count++;
    ctx->tensor_bytes += t->weight_bytes + (fmt ? (size_t)O * sizeof(float) : 0);
    *tensor = t;
    return 1;
 }
 extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor,
                                 float *y, const float *x,
                                 const void *weights, const float *scales,
                                 int fmt, int S, int I, int O, int device) {
    if (S < 1 || !coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device)) return 0;
    ColiCudaTensor *t = *tensor;
    DeviceContext *ctx = find_ctx(t->device);
    if (!select_ctx(ctx)) return 0;
    size_t rb = row_bytes(fmt, I);
    size_t xb = (size_t)S * I * sizeof(float), yb = (size_t)S * O * sizeof(float);
    if (!reserve(&ctx->x, &ctx->x_cap, xb) || !reserve(&ctx->y, &ctx->y_cap, yb)) return 0;
    if (!cuda_ok(hipMemcpy(ctx->x, x, xb, hipMemcpyHostToDevice), "input upload")) return 0;
    dim3 grid((unsigned)O, (unsigned)S);
    quant_matmul<<<grid, 256>>>(ctx->y, ctx->x, t->weights, t->scales, fmt, S, I, O, rb);
    if (!cuda_ok(hipGetLastError(), "matmul launch") ||
        !cuda_ok(hipMemcpy(y, ctx->y, yb, hipMemcpyDeviceToHost), "output download")) return 0;
    return 1;
 }
 extern "C" int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
                                      ColiCudaTensor *down, float *y,
                                      const float *x, int S) {
    if (!gate || !up || !down || !x || !y || S < 1 ||
        gate->device != up->device || gate->device != down->device ||
        gate->I != up->I || gate->O != up->O ||
        down->I != gate->O || down->O != gate->I) return 0;
    DeviceContext *ctx = find_ctx(gate->device);
    if (!select_ctx(ctx)) return 0;
    int D = gate->I, I = gate->O;
    size_t xb=(size_t)S*D*sizeof(float), ib=(size_t)S*I*sizeof(float);
    size_t yb=(size_t)S*D*sizeof(float);
    if (!reserve(&ctx->x,&ctx->x_cap,xb) || !reserve(&ctx->y,&ctx->y_cap,yb) ||
        !reserve(&ctx->gate,&ctx->gate_cap,ib) || !reserve(&ctx->up,&ctx->up_cap,ib)) return 0;
    if (!cuda_ok(hipMemcpy(ctx->x,x,xb,hipMemcpyHostToDevice),"expert input upload")) return 0;
    dim3 hidden_grid((unsigned)I,(unsigned)S), output_grid((unsigned)D,(unsigned)S);
    quant_matmul<<<hidden_grid,256>>>(ctx->gate,ctx->x,gate->weights,gate->scales,
        gate->fmt,S,D,I,row_bytes(gate->fmt,D));
    quant_matmul<<<hidden_grid,256>>>(ctx->up,ctx->x,up->weights,up->scales,
        up->fmt,S,D,I,row_bytes(up->fmt,D));
    size_t n=(size_t)S*I;
    silu_mul<<<(unsigned)((n+255)/256),256>>>(ctx->gate,ctx->up,n);
    quant_matmul<<<output_grid,256>>>(ctx->y,ctx->gate,down->weights,down->scales,
        down->fmt,S,I,D,row_bytes(down->fmt,I));
    if (!cuda_ok(hipGetLastError(),"expert MLP launch") ||
        !cuda_ok(hipMemcpy(y,ctx->y,yb,hipMemcpyDeviceToHost),"expert output download")) return 0;
    return 1;
 }
 extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates,
                                        ColiCudaTensor *const *ups,
                                        ColiCudaTensor *const *downs,
                                        const int *rows, int count,
                                        float *y, const float *x) {
    if (!gates || !ups || !downs || !rows || !x || !y || count < 1) return 0;
    ColiCudaTensor *first=gates[0];
    if (!first) return 0;
    int device=first->device,D=first->I,I=first->O,total=0,max_rows=0;
    GroupDesc host[64]; if(count>64) return 0;
    int all_s4=1;
    for(int c=0;c<count;c++){
        ColiCudaTensor *g=gates[c],*u=ups[c],*d=downs[c];
        if(!g||!u||!d||rows[c]<1||g->device!=device||u->device!=device||d->device!=device||
           g->I!=D||u->I!=D||g->O!=I||u->O!=I||d->I!=I||d->O!=D) return 0;
        host[c]={g->weights,u->weights,d->weights,g->scales,u->scales,d->scales,
                 g->fmt,u->fmt,d->fmt,rows[c],total};
        all_s4&=g->fmt==2&&u->fmt==2&&d->fmt==2;
        total+=rows[c]; if(rows[c]>max_rows) max_rows=rows[c];
    }
    DeviceContext *ctx=find_ctx(device); if(!select_ctx(ctx)) return 0;
    size_t xb=(size_t)total*D*sizeof(float), ib=(size_t)total*I*sizeof(float);
    if(!reserve(&ctx->x,&ctx->x_cap,xb)||!reserve(&ctx->y,&ctx->y_cap,xb)||
       !reserve(&ctx->gate,&ctx->gate_cap,ib)||!reserve(&ctx->up,&ctx->up_cap,ib)||
       !reserve_bytes(&ctx->group_desc,&ctx->group_desc_cap,(size_t)count*sizeof(GroupDesc))) return 0;
    int async=!getenv("COLI_CUDA_ASYNC")||atoi(getenv("COLI_CUDA_ASYNC"));
    if(async&&(!reserve_pinned(&ctx->host_x,&ctx->host_x_cap,xb)||
               !reserve_pinned(&ctx->host_y,&ctx->host_y_cap,xb)))return 0;
    hipError_t copy_desc=async?hipMemcpyAsync(ctx->group_desc,host,(size_t)count*sizeof(GroupDesc),
                                                hipMemcpyHostToDevice,ctx->stream)
                               :hipMemcpy(ctx->group_desc,host,(size_t)count*sizeof(GroupDesc),hipMemcpyHostToDevice);
    if(!cuda_ok(copy_desc,"expert group descriptors"))return 0;
    int profile=getenv("COLI_CUDA_PROFILE")&&atoi(getenv("COLI_CUDA_PROFILE"));
    hipEvent_t ev[4]={};
    if(profile) for(int i=0;i<4;i++) if(!cuda_ok(hipEventCreate(&ev[i]),"profile event")) profile=0;
    if(profile) hipEventRecord(ev[0],ctx->stream);
    if(async)std::memcpy(ctx->host_x,x,xb);
    hipError_t copy_x=async?hipMemcpyAsync(ctx->x,ctx->host_x,xb,hipMemcpyHostToDevice,ctx->stream)
                            :hipMemcpy(ctx->x,x,xb,hipMemcpyHostToDevice);
    if(!cuda_ok(copy_x,"expert group input upload")) return 0;
    if(profile) hipEventRecord(ev[1],ctx->stream);
    GroupDesc *dev=(GroupDesc*)ctx->group_desc;
    int tc=getenv("COLI_CUDA_TC_INT4")&&atoi(getenv("COLI_CUDA_TC_INT4"));
    tc=tc&&all_s4&&D%32==0&&I%32==0&&D%8==0&&I%8==0;
    int tc_min=getenv("COLI_CUDA_TC_MIN_ROWS")?atoi(getenv("COLI_CUDA_TC_MIN_ROWS")):8;
    for(int c=0;c<count&&tc;c++)tc=rows[c]>=tc_min;
    if(tc){
        size_t qb=(size_t)(total+7)*(size_t)(D>I?D:I)/2;
        if(!reserve_bytes((void**)&ctx->qx,&ctx->qx_cap,qb)||
           !reserve(&ctx->qscale,&ctx->qscale_cap,(size_t)(total+7)*sizeof(float)))return 0;
        hipMemsetAsync(ctx->qx,0,qb,ctx->stream);
        quantize_s4_rows<<<total,256,0,ctx->stream>>>(ctx->qx,ctx->qscale,ctx->x,total,D);
        grouped_s4_wmma<<<dim3((unsigned)((I+63)/64),(unsigned)count),256,0,ctx->stream>>>(ctx->gate,ctx->qx,ctx->qscale,dev,D,I,0);
        grouped_s4_wmma<<<dim3((unsigned)((I+63)/64),(unsigned)count),256,0,ctx->stream>>>(ctx->up,ctx->qx,ctx->qscale,dev,D,I,1);
        silu_mul<<<(unsigned)(((size_t)total*I+255)/256),256,0,ctx->stream>>>(ctx->gate,ctx->up,(size_t)total*I);
        quantize_s4_rows<<<total,256,0,ctx->stream>>>(ctx->qx,ctx->qscale,ctx->gate,total,I);
        grouped_s4_wmma<<<dim3((unsigned)((D+63)/64),(unsigned)count),256,0,ctx->stream>>>(ctx->y,ctx->qx,ctx->qscale,dev,I,D,2);
    }else if(all_s4&&(!getenv("COLI_CUDA_W4_PACKED")||atoi(getenv("COLI_CUDA_W4_PACKED")))){
        dim3 hg((unsigned)I,(unsigned)max_rows,(unsigned)count),og((unsigned)D,(unsigned)max_rows,(unsigned)count);
        int dual=!getenv("COLI_CUDA_DUAL_PROJ")||atoi(getenv("COLI_CUDA_DUAL_PROJ"));
        if(dual)grouped_hidden_w4_dual<<<hg,256,0,ctx->stream>>>(ctx->gate,ctx->up,ctx->x,dev,I,D);
        else{
            grouped_hidden_w4<<<hg,256,0,ctx->stream>>>(ctx->gate,ctx->x,dev,I,D,0);
            grouped_hidden_w4<<<hg,256,0,ctx->stream>>>(ctx->up,ctx->x,dev,I,D,1);
        }
        silu_mul<<<(unsigned)(((size_t)total*I+255)/256),256,0,ctx->stream>>>(ctx->gate,ctx->up,(size_t)total*I);
        grouped_down_w4<<<og,256,0,ctx->stream>>>(ctx->y,ctx->gate,dev,D,I);
    }else{
        dim3 hg((unsigned)I,(unsigned)max_rows,(unsigned)count),og((unsigned)D,(unsigned)max_rows,(unsigned)count);
        grouped_hidden<<<hg,256,0,ctx->stream>>>(ctx->gate,ctx->x,dev,I,D,0);
        grouped_hidden<<<hg,256,0,ctx->stream>>>(ctx->up,ctx->x,dev,I,D,1);
        silu_mul<<<(unsigned)(((size_t)total*I+255)/256),256,0,ctx->stream>>>(ctx->gate,ctx->up,(size_t)total*I);
        grouped_down<<<og,256,0,ctx->stream>>>(ctx->y,ctx->gate,dev,D,I);
    }
    if(profile) hipEventRecord(ev[2],ctx->stream);
    if(!async&&!cuda_ok(hipStreamSynchronize(ctx->stream),"expert group synchronize"))return 0;
    hipError_t copy_y=async?hipMemcpyAsync(ctx->host_y,ctx->y,xb,hipMemcpyDeviceToHost,ctx->stream)
                            :hipMemcpy(y,ctx->y,xb,hipMemcpyDeviceToHost);
    if(!cuda_ok(hipGetLastError(),"expert group launch")||!cuda_ok(copy_y,"expert group output download"))return 0;
    if(async){if(!cuda_ok(hipStreamSynchronize(ctx->stream),"expert group synchronize"))return 0;
        std::memcpy(y,ctx->host_y,xb);}
    if(profile){
        hipEventRecord(ev[3],ctx->stream); hipEventSynchronize(ev[3]); float a=0,b=0,c=0;
        hipEventElapsedTime(&a,ev[0],ev[1]); hipEventElapsedTime(&b,ev[1],ev[2]);
        hipEventElapsedTime(&c,ev[2],ev[3]);
        { std::lock_guard<std::mutex> lock(g_group_stats_mu);
          g_group_h2d_ms+=a; g_group_kernel_ms+=b; g_group_d2h_ms+=c; }
        for(int i=0;i<4;i++) hipEventDestroy(ev[i]);
    }
    { std::lock_guard<std::mutex> lock(g_group_stats_mu);
      g_group_calls++; g_group_experts+=(uint64_t)count; g_group_rows+=(uint64_t)total; }
    return 1;
 }
 extern "C" int coli_cuda_attention_absorb(ColiCudaTensor *w,float *ctx,const float *q,
                                            const float *latent,const float *rope,int H,int Q,
                                            int R,int V,int K,int T,float scale){
    if(!w||!ctx||!q||!latent||!rope||H<1||Q<1||R<1||V<1||K<1||K>512||T<1||T>4096||
       w->I!=K||w->O!=H*(Q+V))return 0;
    DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0;
    size_t qb=(size_t)H*(Q+R)*sizeof(float),lb=(size_t)T*K*sizeof(float);
    size_t rb=(size_t)T*R*sizeof(float),cb=(size_t)H*V*sizeof(float);
    if(!reserve(&dc->aq,&dc->aq_cap,qb)||!reserve(&dc->al,&dc->al_cap,lb)||
       !reserve(&dc->ar,&dc->ar_cap,rb)||!reserve(&dc->ac,&dc->ac_cap,cb))return 0;
    if(!cuda_ok(hipMemcpyAsync(dc->aq,q,qb,hipMemcpyHostToDevice,dc->stream),"attention q upload")||
       !cuda_ok(hipMemcpyAsync(dc->al,latent,lb,hipMemcpyHostToDevice,dc->stream),"attention latent upload")||
       !cuda_ok(hipMemcpyAsync(dc->ar,rope,rb,hipMemcpyHostToDevice,dc->stream),"attention rope upload"))return 0;
    size_t shared=(size_t)(2*K+T)*sizeof(float);
    attention_absorb_kernel<<<H,256,shared,dc->stream>>>(dc->ac,dc->aq,dc->al,dc->ar,w->weights,w->scales,
        w->fmt,H,Q,R,V,K,T,scale);
    if(!cuda_ok(hipGetLastError(),"attention absorb launch")||
       !cuda_ok(hipMemcpyAsync(ctx,dc->ac,cb,hipMemcpyDeviceToHost,dc->stream),"attention context download")||
       !cuda_ok(hipStreamSynchronize(dc->stream),"attention synchronize"))return 0;
    return 1;
 }
 extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) {
    if (!tensor) return;
    DeviceContext *ctx = find_ctx(tensor->device);
    if (ctx) select_ctx(ctx);
    if (tensor->tracked && ctx) {
        size_t bytes = tensor->weight_bytes + (tensor->fmt ? (size_t)tensor->O * sizeof(float) : 0);
        if (ctx->tensor_count) ctx->tensor_count--;
        if (ctx->tensor_bytes >= bytes) ctx->tensor_bytes -= bytes;
    }
    if (tensor->weights) hipFree(tensor->weights);
    if (tensor->scales) hipFree(tensor->scales);
    std::free(tensor);
 }
 extern "C" size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor) {
    return tensor ? tensor->weight_bytes + (tensor->fmt ? (size_t)tensor->O * sizeof(float) : 0) : 0;
 }
 extern "C" int coli_cuda_tensor_device(const ColiCudaTensor *tensor) {
    return tensor ? tensor->device : -1;
 }
@@ -0,0 +1,219 @@
 /* backend_loader.c — Windows runtime loader for coli_cuda.dll.
 *
 * Why this exists: the engine is built with MinGW-w64 (gcc), but CUDA kernels
 * must be compiled with MSVC + nvcc. We cannot link a CUDA .o into a gcc binary
 * reliably across the MSVC/MinGW ABI, and nvcc requires cl.exe as its host
 * compiler. The clean cross-toolchain split is: build the CUDA backend into a
 * standalone coli_cuda.dll with nvcc+MSVC, then load it here at runtime via
 * LoadLibrary/GetProcAddress. The host (glm.exe) never links cudart directly.
 *
 * On Linux this file is not compiled (the Makefile links backend_cuda.o
 * directly). On Windows, when COLI_CUDA is defined, glm.c calls the
 * coli_cuda_* wrappers below, which forward through function pointers resolved
 * from the DLL at first use. If the DLL is absent, every call safely returns
 * the "not initialized" sentinel (0 / no-op) and the engine falls back to CPU.
 *
 * ABI note: ColiCudaTensor* is opaque to the host (it stores the pointer,
 * never dereferences it), so the MSVC-allocated struct is safe to pass across
 * the boundary as an opaque handle. All scalar types (int, size_t, pointers)
 * agree between MSVC and MinGW-w64 on x86-64.
 */
 #ifdef _WIN32
 #include <stdio.h>
 #include <stddef.h>
 #include <windows.h>
 #include "backend_cuda.h"
 /* Function-pointer typedefs matching each exported symbol. */
 typedef int            (*fn_init)(const int *devices, int count);
 typedef void           (*fn_shutdown)(void);
 typedef int            (*fn_device_count)(void);
 typedef int            (*fn_device_at)(int index);
 typedef int            (*fn_mem_info)(int device, size_t *free_bytes, size_t *total_bytes);
 typedef void           (*fn_stats)(int device, size_t *tensor_count, size_t *tensor_bytes);
 typedef void           (*fn_group_stats)(uint64_t *calls, uint64_t *experts, uint64_t *rows,
                                         double *h2d_ms, double *kernel_ms, double *d2h_ms);
 typedef int            (*fn_expert_mlp)(ColiCudaTensor *gate, ColiCudaTensor *up,
                                        ColiCudaTensor *down, float *y, const float *x, int S);
 typedef int            (*fn_expert_group)(ColiCudaTensor *const *gates, ColiCudaTensor *const *ups,
                                          ColiCudaTensor *const *downs, const int *rows, int count,
                                          float *y, const float *x);
 typedef int            (*fn_attention_absorb)(ColiCudaTensor *kv_b, float *ctx, const float *q,
                                              const float *latent, const float *rope, int H, int Q,
                                              int R, int V, int K, int T, float attention_scale);
 typedef int            (*fn_tensor_upload)(ColiCudaTensor **tensor, const void *weights,
                                           const float *scales, int fmt, int I, int O, int device);
 typedef int            (*fn_matmul)(ColiCudaTensor **tensor, float *y, const float *x,
                                    const void *weights, const float *scales,
                                    int fmt, int S, int I, int O, int device);
 typedef void           (*fn_tensor_free)(ColiCudaTensor *tensor);
 typedef size_t         (*fn_tensor_bytes)(const ColiCudaTensor *tensor);
 typedef int            (*fn_tensor_device)(const ColiCudaTensor *tensor);
 /* Resolved pointers, plus a flag so we attempt the load at most once. */
 static struct {
    int loaded;        /* 1 = load attempted (success or fail), 0 = not yet */
    int available;     /* 1 = DLL loaded and all symbols resolved */
    HMODULE dll;
    fn_init            init;
    fn_shutdown        shutdown;
    fn_device_count    device_count;
    fn_device_at       device_at;
    fn_mem_info        mem_info;
    fn_stats           stats;
    fn_group_stats     group_stats;
    fn_expert_mlp      expert_mlp;
    fn_expert_group    expert_group;
    fn_attention_absorb attention_absorb;
    fn_tensor_upload   tensor_upload;
    fn_matmul          matmul;
    fn_tensor_free     tensor_free;
    fn_tensor_bytes    tensor_bytes;
    fn_tensor_device   tensor_device;
 } g_cuda;
 /* Resolve the DLL and all 11 symbols. Returns 1 on success, 0 otherwise.
 * Idempotent: the first call (success or fail) sticks; later calls are no-ops
 * that return the cached result. The engine treats a 0 return as "CUDA
 * unavailable" and falls back to the CPU path without aborting. */
 static int coli_cuda_load(void){
    if(g_cuda.loaded) return g_cuda.available;
    g_cuda.loaded = 1;
    /* Search the model directory first (so a DLL shipped next to the model
     * wins), then the engine directory, then the default DLL search path. */
    g_cuda.dll = LoadLibraryA("coli_cuda.dll");
    if(!g_cuda.dll){
        fprintf(stderr, "[CUDA] coli_cuda.dll not found; GPU tier disabled "
                        "(CPU path remains active).\n");
        return 0;
    }
    #define RESOLVE(name, type) \
        /* GetProcAddress returns FARPROC (void(*)(void)); casting it to a   \
         * specific function-pointer type is the standard LoadLibrary idiom. \
         * -Wcast-function-type flags it but it is safe: the DLL exported     \
         * the symbol with extern "C" and the exact signature we expect. */   \
        _Pragma("GCC diagnostic push") \
        _Pragma("GCC diagnostic ignored \"-Wcast-function-type\"") \
        g_cuda.name = (type)GetProcAddress(g_cuda.dll, "coli_cuda_" #name); \
        _Pragma("GCC diagnostic pop") \
        if(!g_cuda.name){ \
            fprintf(stderr, "[CUDA] coli_cuda.dll missing symbol coli_cuda_" #name "\n"); \
            FreeLibrary(g_cuda.dll); g_cuda.dll=NULL; return 0; }
    RESOLVE(init,           fn_init)
    RESOLVE(shutdown,       fn_shutdown)
    RESOLVE(device_count,   fn_device_count)
    RESOLVE(device_at,      fn_device_at)
    RESOLVE(mem_info,       fn_mem_info)
    RESOLVE(stats,          fn_stats)
    RESOLVE(group_stats,    fn_group_stats)
    RESOLVE(expert_mlp,     fn_expert_mlp)
    RESOLVE(expert_group,   fn_expert_group)
    RESOLVE(attention_absorb, fn_attention_absorb)
    RESOLVE(tensor_upload,  fn_tensor_upload)
    RESOLVE(matmul,         fn_matmul)
    RESOLVE(tensor_free,    fn_tensor_free)
    RESOLVE(tensor_bytes,   fn_tensor_bytes)
    RESOLVE(tensor_device,  fn_tensor_device)
    #undef RESOLVE
    g_cuda.available = 1;
    return 1;
 }
 /* ---- Public wrappers: match backend_cuda.h signatures exactly.
 * Each forwards to the resolved pointer; if the DLL never loaded, return the
 * "not initialized" result the engine already handles (init returns 0, matmul
 * returns 0 so the caller marks the tensor cuda_failed and uses CPU). ---- */
 int coli_cuda_init(const int *devices, int count){
    if(!coli_cuda_load()) return 0;
    return g_cuda.init(devices, count);
 }
 void coli_cuda_shutdown(void){
    if(g_cuda.available && g_cuda.shutdown) g_cuda.shutdown();
 }
 int coli_cuda_device_count(void){
    if(!g_cuda.available) return 0;
    return g_cuda.device_count();
 }
 int coli_cuda_device_at(int index){
    if(!g_cuda.available) return -1;
    return g_cuda.device_at(index);
 }
 int coli_cuda_mem_info(int device, size_t *free_bytes, size_t *total_bytes){
    if(!g_cuda.available){ if(free_bytes)*free_bytes=0; if(total_bytes)*total_bytes=0; return 0; }
    return g_cuda.mem_info(device, free_bytes, total_bytes);
 }
 void coli_cuda_stats(int device, size_t *tensor_count, size_t *tensor_bytes){
    if(!g_cuda.available){ if(tensor_count)*tensor_count=0; if(tensor_bytes)*tensor_bytes=0; return; }
    g_cuda.stats(device, tensor_count, tensor_bytes);
 }
 void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64_t *rows,
                           double *h2d_ms, double *kernel_ms, double *d2h_ms){
    if(!g_cuda.available){
        if(calls)*calls=0; if(experts)*experts=0; if(rows)*rows=0;
        if(h2d_ms)*h2d_ms=0; if(kernel_ms)*kernel_ms=0; if(d2h_ms)*d2h_ms=0;
        return;
    }
    g_cuda.group_stats(calls, experts, rows, h2d_ms, kernel_ms, d2h_ms);
 }
 int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
                         ColiCudaTensor *down, float *y, const float *x, int S){
    if(!g_cuda.available) return 0;
    return g_cuda.expert_mlp(gate, up, down, y, x, S);
 }
 int coli_cuda_expert_group(ColiCudaTensor *const *gates, ColiCudaTensor *const *ups,
                           ColiCudaTensor *const *downs, const int *rows, int count,
                           float *y, const float *x){
    if(!g_cuda.available) return 0;
    return g_cuda.expert_group(gates, ups, downs, rows, count, y, x);
 }
 int coli_cuda_attention_absorb(ColiCudaTensor *kv_b, float *ctx, const float *q,
                               const float *latent, const float *rope, int H, int Q,
                               int R, int V, int K, int T, float attention_scale){
    if(!g_cuda.available) return 0;
    return g_cuda.attention_absorb(kv_b, ctx, q, latent, rope, H, Q, R, V, K, T, attention_scale);
 }
 int coli_cuda_tensor_upload(ColiCudaTensor **tensor, const void *weights,
                            const float *scales, int fmt, int I, int O, int device){
    if(!g_cuda.available) return 0;
    return g_cuda.tensor_upload(tensor, weights, scales, fmt, I, O, device);
 }
 int coli_cuda_matmul(ColiCudaTensor **tensor, float *y, const float *x,
                     const void *weights, const float *scales,
                     int fmt, int S, int I, int O, int device){
    if(!g_cuda.available) return 0;
    return g_cuda.matmul(tensor, y, x, weights, scales, fmt, S, I, O, device);
 }
 void coli_cuda_tensor_free(ColiCudaTensor *tensor){
    if(g_cuda.available && g_cuda.tensor_free) g_cuda.tensor_free(tensor);
 }
 size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor){
    if(!g_cuda.available) return 0;
    return g_cuda.tensor_bytes(tensor);
 }
 int coli_cuda_tensor_device(const ColiCudaTensor *tensor){
    if(!g_cuda.available) return -1;
    return g_cuda.tensor_device(tensor);
 }
 #endif /* _WIN32 */
@@ -110,7 +110,7 @@ def cuda_binary():
    if not os.path.exists(GLM) or sys.platform != "linux": return False
    try:
        linked=subprocess.run(["ldd",GLM],capture_output=True,text=True,timeout=3)
-        return any("libcudart" in line and "not found" not in line
+        return any("libamdhip64" in line and "not found" not in line
                   for line in linked.stdout.splitlines())
    except (OSError,subprocess.SubprocessError): return False
@@ -383,13 +383,38 @@ def cmd_chat(a):
    banner(f"chat · {os.path.basename(a.model)} · ram {a.ram or '-'}GB · topp {a.topp or 'off'}")
    errlog=tempfile.NamedTemporaryFile(mode="w+", suffix=".log", delete=False)
    e=env_for(a); e["SERVE"]="1"
    # stderr -> PIPE, NOT stderr=errlog (file). On Windows/MinGW, pointing the
    # child's stderr at a file/DEVNULL handle stalls the CRT so stdout (the byte
    # protocol coli reads one byte at a time) never flushes and chat hangs at
    # ~10 GB resident. A PIPE whose read end nobody drains still works: the
    # engine emits only ~400 bytes of status to stderr, which fits comfortably
    # in the OS pipe buffer, so it never blocks. We snapshot stderr into errlog
    # once the READY sentinel arrives, so the status-line display below works
    # exactly as before. (Do NOT add a concurrent stderr drain thread: on
    # Windows, reading two child pipes simultaneously deadlocks CPython's IO.)
    p=subprocess.Popen([GLM,str(a.cap)], env=e, stdin=subprocess.PIPE,
-                       stdout=subprocess.PIPE, stderr=errlog, bufsize=0)
+                       stdout=subprocess.PIPE, stderr=subprocess.PIPE, bufsize=0)
    sp=Spinner("waking the giant (744B)…"); sp.start()
    st=stream_turn(p, READY, lambda b: None)
    sp.stop()
    if st is None:
        try: errlog.write(p.stderr.read().decode("utf-8","replace"))
        except (OSError, ValueError): pass
        errlog.seek(0); print(errlog.read()[-1500:]); sys.exit("the engine exited while loading")
    # READY received. Drain the child's stderr into errlog without blocking:
    # the engine is still alive (blocked on stdin), so a plain read() would
    # hang forever waiting for EOF. A short bounded drain grabs the ~400 bytes
    # of load-time status ([RAM_GB], [MTP], ...) that were already emitted.
    _drain_box={"done":False}
    def _drain():
        try: errlog.write(p.stderr.read().decode("utf-8","replace"))
        except (OSError, ValueError): pass
        _drain_box["done"]=True
    threading.Thread(target=_drain, daemon=True).start()
    _drain_box["th"]=threading.current_thread()
    for _ in range(20):           # up to ~1s for the load-status lines
        if _drain_box["done"]: break
        time.sleep(0.05)
    errlog.flush()
    try:
        elog=open(errlog.name).read()
@@ -27,7 +27,9 @@
 #include <stdatomic.h>                            /* PIPE ready-flags/job queue + PILOT_REAL cross-layer handshake */
 #include <sched.h>                                /* sched_yield: PIPE spin / PILOT barrier */
 #include <unistd.h>
-#include <sys/select.h>
+#if defined(__APPLE__) || defined(__linux__)
 #include <sys/select.h>                             /* select() serve-loop polling (#68); not on native MinGW */
 #endif
 #if defined(__APPLE__) || defined(__linux__)
 #include <sys/resource.h>
 #include <sys/mman.h>                             /* mlock: inchioda le pagine in RAM / wire pages into RAM */
@@ -439,6 +441,8 @@ static void matmul_i2(float *y, const float *x, const uint8_t *q2, const float *
 * RMS per matmul (attivazione int8), IDOT=0 torna al percorso f32 esatto. */
 #if defined(__AVX512VNNI__) && defined(__AVX512BW__)
 #define IDOT_KERNEL "avx512-vnni"
 #elif defined(__AVXVNNI__) && defined(__AVX2__)
 #define IDOT_KERNEL "avx-vnni"
 #elif defined(__AVX2__)
 #define IDOT_KERNEL "avx2"
 #elif defined(__ARM_NEON)
@@ -475,6 +479,12 @@ static inline int hsum256_i32(__m256i v){
    return _mm_cvtsi128_si32(lo);
 }
 #endif
 #if defined(__AVXVNNI__) && defined(__AVX2__)
 /* hsum di un __m128i a 4 lane s32 (l'AVX-VNNI 128-bit accumula su 4 lane). */
 static inline int hsum128_i32(__m128i v){
    v=_mm_hadd_epi32(v,v); v=_mm_hadd_epi32(v,v); return _mm_cvtsi128_si32(v);
 }
 #endif
 /* dot int8·int8: trucco del segno (|w| unsigned × x·sign(w) signed). Sicuro:
 * coppie <= 128*127*2 = 32512 < 32767, accumulo s32 fino a I=16384. */
 static inline int32_t dot_i8i8(const int8_t *w, const int8_t *x, int I){
@@ -493,6 +503,18 @@ static inline int32_t dot_i8i8(const int8_t *w, const int8_t *x, int I){
        acc=_mm512_dpbusd_epi32(acc,_mm512_abs_epi8(wv),xs);
    }
    sum=_mm512_reduce_add_epi32(acc);
 #elif defined(__AVXVNNI__) && defined(__AVX2__)
    /* AVX-VNNI 128-bit: vpdpbusd u8*s8 -> s32, 16 byte/iter. Stesso trucco del
     * segno della variante 512-bit: |w| via abs, segno piegato in x con maschera
     * (w==0 -> product 0). __AVX2__ serve per _mm_sign_epi8 / abs. */
    __m128i acc=_mm_setzero_si128();
    for(;i+16<=I;i+=16){
        __m128i wv=_mm_loadu_si128((const __m128i*)(w+i));
        __m128i xv=_mm_loadu_si128((const __m128i*)(x+i));
        __m128i xs=_mm_sign_epi8(xv,wv);              /* x * sign(w); _mm_sign zona __AVX2__ */
        acc=_mm_dpbusd_epi32(acc,_mm_abs_epi8(wv),xs);
    }
    sum=hsum128_i32(acc);
 #elif defined(__AVX2__)
    __m256i acc=_mm256_setzero_si256(); const __m256i ones=_mm256_set1_epi16(1);
    for(;i+32<=I;i+=32){
@@ -563,6 +585,23 @@ static inline int32_t dot_i4i8(const uint8_t *w4, const int8_t *x, int I){
        acc=_mm512_dpbusd_epi32(acc,_mm512_abs_epi8(wv),xs);
    }
    sum=_mm512_reduce_add_epi32(acc);
 #elif defined(__AVXVNNI__) && defined(__AVX2__)
    /* AVX-VNNI 128-bit, int4: 16 byte = 32 nibble -> int8 [-8,7] in due half
     * (n0/n1), ciascuno alimentato a un vpdpbusd da 16 byte. Stesso unpack
     * 128-bit del ramo AVX2 sotto; 32 elementi/iter come li. */
    const __m128i m4=_mm_set1_epi8(0x0F); const __m128i b8=_mm_set1_epi8(8);
    __m128i acc=_mm_setzero_si128();
    for(;i+32<=I;i+=32){
        __m128i by=_mm_loadu_si128((const __m128i*)(w4+(i>>1)));   /* 16 byte = 32 nibble */
        __m128i lo=_mm_and_si128(by,m4), hi=_mm_and_si128(_mm_srli_epi16(by,4),m4);
        __m128i n0=_mm_unpacklo_epi8(lo,hi), n1=_mm_unpackhi_epi8(lo,hi);  /* nibble in ordine */
        __m128i w0=_mm_sub_epi8(n0,b8), w1=_mm_sub_epi8(n1,b8);
        __m128i x0=_mm_loadu_si128((const __m128i*)(x+i));
        __m128i x1=_mm_loadu_si128((const __m128i*)(x+i+16));
        acc=_mm_dpbusd_epi32(acc,_mm_abs_epi8(w0),_mm_sign_epi8(x0,w0));
        acc=_mm_dpbusd_epi32(acc,_mm_abs_epi8(w1),_mm_sign_epi8(x1,w1));
    }
    sum=hsum128_i32(acc);
 #elif defined(__AVX2__)
    const __m128i m4=_mm_set1_epi8(0x0F); const __m256i b8=_mm256_set1_epi8(8);
    const __m256i ones=_mm256_set1_epi16(1);
@@ -1128,7 +1167,9 @@ static pthread_mutex_t g_map_mtx = PTHREAD_MUTEX_INITIALIZER;   /* expert_load e
 static void *map_of_fd(int fd){
    pthread_mutex_lock(&g_map_mtx);
    for(int i=0;i<g_nmaps;i++) if(g_maps[i].fd==fd){ void *b=g_maps[i].base; pthread_mutex_unlock(&g_map_mtx); return b; }
-    void *base=NULL; struct stat st;
+    void *base=NULL;
 #if defined(__APPLE__) || defined(__linux__)
    struct stat st;
    if(g_nmaps<512 && fstat(fd,&st)==0){
        size_t len=((size_t)st.st_size+16383)&~(size_t)16383;
        void *p=mmap(NULL,len,PROT_READ,MAP_SHARED,fd,0);
@@ -1139,6 +1180,7 @@ static void *map_of_fd(int fd){
 #endif
        }
    }
 #endif
    pthread_mutex_unlock(&g_map_mtx);
    return base;
 }
@@ -1198,7 +1240,9 @@ static int expert_load(Model *m, int layer, int eid, ESlot *s, int fatal){
             * residency. This is pread's I/O without the copy and without the slab. */
            for(int k=0;k<3;k++){
                char *p=(char*)bw[k]+tw[k]->off; size_t n=(size_t)tw[k]->nbytes;
 #if defined(__APPLE__) || defined(__linux__)
                madvise((void*)((uintptr_t)p & ~16383UL), n+16384, MADV_WILLNEED);
 #endif
                volatile char acc=0;
                for(size_t i=0;i<n;i+=4096) acc+=p[i];
                acc+=p[n-1]; (void)acc;
@@ -3082,6 +3126,7 @@ static int mux_submit(Model *m, Tok *T, ServeCtx *ctx, ServeReq *req, int nctx,
 }
 static void run_serve_mux(Model *m, const char *snap){
 #if defined(__APPLE__) || defined(__linux__)
    char tkp[2048]; snprintf(tkp,sizeof(tkp),"%s/tokenizer.json",snap);
    Tok T; tok_load(&T,tkp); int eos=tok_id_of(&T,"<|endoftext|>"); stops_arm(&m->c,eos);
    g_draft=0; /* one scheduler owns every forward; MTP/speculation is not ragged-safe */
@@ -3123,9 +3168,33 @@ static void run_serve_mux(Model *m, const char *snap){
    usage_save(m);
    for(int i=0;i<nctx;i++) serve_ctx_free(m,&ctx[i]); free(ctx); free(req);
    m->kv=NULL; m->Lc=m->Rc=m->Ic=NULL; m->kv_start=NULL; m->max_t=0;
 #else
    /* SERVE_BATCH (continuous batching) uses select() on stdin, a Unix-ism.
     * Not yet ported to native Windows — fall back to the single-sequence
     * serve path (run_serve). Remove this stub once select()-free polling
     * (e.g. WaitForSingleObject on the stdin handle) is implemented. */
    (void)snap;
    fprintf(stderr,"[SERVE_BATCH] continuous-batching serve is not yet available on "
                   "native Windows; use the default serve path (omit SERVE_BATCH).\n");
 #endif
 }
 static void run_serve(Model *m, const char *snap){
    /* Serve mode speaks a byte protocol over BOTH stdout and stdin:
     *   stdout: \x01\x01READY\x01\x01\n, STAT lines, \x01\x01END\x01\x01\n
     *   stdin:  text lines plus \x02RESET / \x02MORE control bytes.
     * 'coli' matches the sentinels with endswith() and a "^STAT ..." regex,
     * so they must arrive byte-exact (LF, no CR). On Windows the CRT opens
     * both handles in TEXT mode: stdout translates '\n'->'\r\n' (so the READY
     * sentinel never matches and chat hangs at ~10 GB resident), and stdin
     * translates '\r\n'->'\n' and rejects writes of raw bytes with EINVAL,
     * breaking the control protocol. Put BOTH handles in BINARY mode so the
     * protocol bytes are exact in both directions. No-op on Linux/macOS. */
 #ifdef _WIN32
    _setmode(_fileno(stdin),  _O_BINARY);
    _setmode(_fileno(stdout), _O_BINARY);
    setvbuf(stdout, NULL, _IONBF, 0);
 #endif
    char tkp[2048]; snprintf(tkp,sizeof(tkp),"%s/tokenizer.json",snap);
    Tok T; tok_load(&T,tkp);
    int eos=tok_id_of(&T,"<|endoftext|>");
@@ -39,7 +39,7 @@ typedef struct { Slot *slots; int n, cap; } LCache;
 typedef struct {
    Cfg c;
    shards S;
-    int quant_bits;        /* bit di quantizzazione degli expert (2..8); 16 = f32 */
+    int quant_bits;        /* bit di quantizzazione degli expert (2..8); storage int8, niente f32 (#134) */
    float *embed, *lm_head, *final_norm;
    Layer *L;
    LCache *cache;          /* [n_layers] */
@@ -360,6 +360,10 @@ int main(int argc, char **argv) {
    if (!snap) { fprintf(stderr, "set SNAP=<snapshot directory>\n"); return 1; }
    int cap  = argc > 1 ? atoi(argv[1]) : 16;
    int bits = argc > 2 ? atoi(argv[2]) : 8;
    if (bits < 2 || bits > 8) {   /* expert storage is int8_t: bits>8 truncates in quantize_rows (#134). f32 mode is not implemented here — int8 is already token-exact vs the oracle. */
        fprintf(stderr, "quant_bits must be 2..8 (got %d); OLMoE experts are int8-backed, no f32 mode\n", bits);
        return 1;
    }
    const char *refpath = argc > 3 ? argv[3] : "ref.json";
    FILE *f = fopen(refpath, "rb"); if(!f){perror(refpath);return 1;}
@@ -7,6 +7,7 @@ import re
 import shutil
 import statistics
 import subprocess
 import sys
 from pathlib import Path
@@ -76,10 +77,45 @@ def analyze_model(model):
 def memory_available():
    # Linux (and MSYS2/Git-Bash CPython where /proc exists): MemAvailable.
    try:
        text = Path("/proc/meminfo").read_text()
        return int(re.search(r"MemAvailable:\s+(\d+)", text).group(1)) * 1024
    except (OSError, AttributeError):
        pass
    # Windows native CPython: GlobalMemoryStatusEx -> ullAvailPhys.
    # Same definition the C engine uses (compat_meminfo in compat.h):
    # standby/free/zero pages, i.e. reclaimable without swapping.
    if sys.platform == "win32":
        try:
            import ctypes
            class MEMORYSTATUSEX(ctypes.Structure):
                _fields_ = [("dwLength", ctypes.c_ulong),
                            ("dwMemoryLoad", ctypes.c_ulong),
                            ("ullTotalPhys", ctypes.c_ulonglong),
                            ("ullAvailPhys", ctypes.c_ulonglong),
                            ("ullTotalVirtual", ctypes.c_ulonglong),
                            ("ullAvailVirtual", ctypes.c_ulonglong),
                            ("ullAvailExtendedVirtual", ctypes.c_ulonglong)]
            stat = MEMORYSTATUSEX(dwLength=ctypes.sizeof(MEMORYSTATUSEX))
            kernel32 = ctypes.windll.kernel32
            kernel32.GlobalMemoryStatusEx.argtypes = [ctypes.c_void_p]
            kernel32.GlobalMemoryStatusEx.restype = ctypes.c_int
            if kernel32.GlobalMemoryStatusEx(ctypes.byref(stat)) and stat.ullAvailPhys:
                return stat.ullAvailPhys
            # Fallback (e.g. sandboxed callers where GlobalMemoryStatusEx reports
            # nothing): total installed RAM in KB. Less precise than ullAvailPhys
            # — it ignores standby/reclaimable pages — but never returns 0 on a
            # real machine, which keeps the expert cache from being mis-sized.
            total_kb = ctypes.c_ulonglong(0)
            kernel32.GetPhysicallyInstalledSystemMemory.argtypes = [ctypes.c_void_p]
            kernel32.GetPhysicallyInstalledSystemMemory.restype = ctypes.c_int
            if kernel32.GetPhysicallyInstalledSystemMemory(ctypes.byref(total_kb)):
                return total_kb.value * 1024
        except OSError:
            pass
    return 0
@@ -6,7 +6,14 @@ import tempfile
 import unittest
 from pathlib import Path
-from resource_plan import GB, analyze_model, build_plan, environment_for_plan, format_plan
+from resource_plan import (
    GB,
    analyze_model,
    build_plan,
    environment_for_plan,
    format_plan,
    memory_available,
 )
 def write_shard(path, tensors):
@@ -53,6 +60,12 @@ class ResourcePlanTest(unittest.TestCase):
        self.assertEqual(info["expert_count"], 2)
        self.assertEqual(info["per_cap_bytes"], 60)
    def test_memory_available_is_positive(self):
        # Regression: on native Windows CPython, /proc/meminfo does not exist,
        # so the Linux-only path returned 0 and the expert cache was sized to
        # 0 slots/layer. The value must be a sane positive number of bytes.
        self.assertGreater(memory_available(), 0)
    def test_builds_bounded_three_tier_plan(self):
        gpus = [{"index": 0, "name": "test-gpu", "total_bytes": 12 * GB,
                 "free_bytes": 10 * GB}]
@@ -78,14 +78,51 @@ def _quant_last_dim(x, bits, group):
    return out.reshape(*out.shape[:-2], -1) if group else out
-def quantize_param(w, bits, group):
+# --------------------------------------------------------------------------------------
 # Rotation preconditioning (QuaRot / QuIP# family, #81): multiply the input dimension by
 # an orthogonal Q = diag(signs) @ H/sqrt(n) BEFORE quantizing, and by Q^T after — the
 # round-trip Q4(W@Q)@Q.T measures exactly the weight error of a deployed scheme that
 # stores W@Q quantized and rotates activations at runtime (W'@x' = W@x since Q@Q.T = I;
 # the runtime cost is one O(D log D) transform per matmul INPUT, not per weight).
 # Spreading outliers across the block is the point: absmax scales stop being hostage to
 # one heavy coordinate, which is the failure mode #108 measured (margin erosion on MMLU).
 # --------------------------------------------------------------------------------------
 _ROT_CACHE = {}
 def rotation(dim, device, seed=417):
    key = (dim, str(device))
    if key in _ROT_CACHE:
        return _ROT_CACHE[key]
    if dim & (dim - 1):
        raise SystemExit(f"-rot needs power-of-2 input dims (got {dim}); OLMoE dims are 2048/1024")
    h = torch.ones(1, 1, device=device, dtype=torch.float32)
    while h.shape[0] < dim:                       # Sylvester recursion
        h = torch.cat([torch.cat([h, h], 1), torch.cat([h, -h], 1)], 0)
    h /= h.shape[0] ** 0.5                        # orthonormal
    g = torch.Generator().manual_seed(seed + dim)
    signs = (torch.randint(0, 2, (dim,), generator=g).float() * 2 - 1).to(device)
    q = signs[:, None] * h                        # Q = D @ H/sqrt(n), orthogonal
    _ROT_CACHE[key] = q
    return q
 def quantize_param(w, bits, group, rot=False):
    if w.ndim == 3:                        # fused experts [E, in, out] -> move input last
        x = w.transpose(1, 2).contiguous()
-        return _quant_last_dim(x, bits, group).transpose(1, 2).contiguous()
+        x = _rot_quant(x, bits, group) if rot else _quant_last_dim(x, bits, group)
        return x.transpose(1, 2).contiguous()
    if rot:
        return _rot_quant(w, bits, group)
    return _quant_last_dim(w, bits, group)  # nn.Linear [out, in] -- input already last
-SCHEME_RE = re.compile(r"^int(2|4|8)(?:-g(\d+))?(-nohead)?$")
+def _rot_quant(x, bits, group):
    """W -> Qn(W@Q) @ Q^T along the last (input) dim — see rotation() above."""
    q = rotation(x.shape[-1], x.device)
    return (_quant_last_dim(x.float() @ q, bits, group) @ q.T).contiguous()
 SCHEME_RE = re.compile(r"^int(2|3|4|8)(?:-g(\d+))?(-rot)?(-nohead)?$")
 def parse_scheme(name):
@@ -94,8 +131,8 @@ def parse_scheme(name):
        return None
    m = SCHEME_RE.match(name)
    if not m:
-        raise SystemExit(f"bad scheme '{name}' (expected fp16 | int{{2,4,8}}[-g<N>][-nohead])")
+        raise SystemExit(f"bad scheme '{name}' (expected fp16 | int{{2,3,4,8}}[-g<N>][-rot][-nohead])")
-    return int(m.group(1)), int(m.group(2) or 0), bool(m.group(3))
+    return int(m.group(1)), int(m.group(2) or 0), bool(m.group(3)), bool(m.group(4))
 def is_router(name):
@@ -115,7 +152,7 @@ def apply_scheme(model, scheme):
    spec = parse_scheme(scheme)
    if spec is None:
        return 0, 0, total
-    bits, group, skip_head = spec
+    bits, group, rot, skip_head = spec
    n = qp = 0
    with torch.no_grad():
        for name, p in model.named_parameters():
@@ -123,7 +160,7 @@ def apply_scheme(model, scheme):
                continue
            if skip_head and is_head_or_embed(name):
                continue
-            p.data.copy_(quantize_param(p.data.float(), bits, group).to(p.dtype))
+            p.data.copy_(quantize_param(p.data.float(), bits, group, rot).to(p.dtype))
            n += 1
            qp += p.numel()
    return n, qp, total
@@ -0,0 +1,146 @@
 # warmup.ps1 - overnight expert-cache warmup for colibri
 #
 # Runs `coli run` in a loop with diverse prompts so the engine records which
 # routed experts your workload actually uses into .coli_usage. At startup the
 # engine pins the hottest experts into RAM; the more history it has, the bigger
 # and more accurate that pin gets. This does NOT load random experts - it loads
 # whatever the model actually routes to for these prompts, then promotes the
 # frequent ones.
 #
 # Usage (from the c\ directory):
 #   .\warmup.ps1                          # defaults: model next to repo, 3 rounds
 #   .\warmup.ps1 -Model D:\glm52_i4 -Rounds 10 -Ngen 400
 #
 # Let it run while you sleep. Each iteration logs selections count + hit rate.
 # Ctrl-C is safe: each run saves usage atomically only on clean completion, so
 # the file is never corrupted (but a killed mid-generation run saves nothing).
 #
 # Why diverse prompts? Expert routing is content-dependent. Coding prompts
 # activate different experts than poetry or math. A spread of topics builds a
 # general-purpose pin that helps whatever YOU ask later. If you only ever warm
 # on one topic, the pin overfits to that topic.
 param(
    [string]$Model = (Resolve-Path (Join-Path $PSScriptRoot "..\glm52_i4")).Path,
    [int]$Rounds = 3,
    # Default 32 (not 500): on a cold QLC cache a 500-token run takes hours and
    # a killed mid-generation run saves nothing (usage_save runs only on clean
    # completion). 32 tokens finishes in ~5-10 min even cold, so usage saves
    # frequently and the loop accumulates selections steadily overnight. Each
    # 32-token prompt still records ~90k expert selections.
    [int]$Ngen = 32,
    [string]$Log = (Join-Path $PSScriptRoot "warmup.log")
 )
 # "Continue" (not "Stop"): the engine writes status to stderr, which "Stop"
 # treats as a fatal error and aborts the whole warmup loop on every prompt.
 $ErrorActionPreference = "Continue"
 $Coli = Join-Path $PSScriptRoot "coli"
 if (-not (Test-Path $Coli)) { Write-Error "coli not found at $Coli - run from the c\ directory"; exit 1 }
 if (-not (Test-Path $Model)) { Write-Error "model not found at $Model"; exit 1 }
 # Diverse prompts across domains - each touches a different expert distribution.
 # Kept open-ended ("explain", "write", "list") so generation runs to NGEN tokens
 # and routes through many experts rather than stopping early on a short answer.
 $Prompts = @(
    "Explain how a transformer neural network works, covering attention, feed-forward layers, and backpropagation in detail.",
    "Write a Python function that implements quicksort with in-place partitioning, including comments explaining each step.",
    "Describe the causes and major events of the French Revolution in chronological order.",
    "What is the difference between TCP and UDP? Explain handshakes, reliability, and use cases.",
    "Write a short story about a lighthouse keeper who discovers a message in a bottle.",
    "Explain the theory of general relativity, including the equivalence principle and gravitational time dilation.",
    "List and describe the major organ systems of the human body and their primary functions.",
    "How does photosynthesis work? Explain the light-dependent reactions and the Calvin cycle.",
    "Write a C program that reads a file line by line and counts word frequency using a hash table.",
    "Summarize the plot of Shakespeare's Hamlet, act by act.",
    "Explain the difference between supervised, unsupervised, and reinforcement learning with examples of each.",
    "What causes climate change? Describe the greenhouse effect, carbon cycle, and major greenhouse gases.",
    "Write a recipe for a classic French onion soup, with step-by-step instructions.",
    "Describe how the internet works, from typing a URL to rendering a webpage, including DNS, TCP, HTTP, and browsers.",
    "Explain database normalization, including first, second, and third normal forms with examples.",
    "What is quantum entanglement? Explain it as if to a curious high school student.",
    "Write a poem about the ocean and the passage of time.",
    "Describe the water cycle, including evaporation, condensation, precipitation, and transpiration.",
    "How do vaccines work? Explain the immune response, antibodies, and mRNA vaccine technology.",
    "Explain the Big Bang theory and the evidence supporting it, including cosmic microwave background and redshift.",
    "Write a Python class for a binary search tree with insert, search, and inorder traversal methods.",
    "What are the major branches of philosophy? Describe epistemology, ethics, metaphysics, and logic.",
    "Explain how a CPU executes an instruction, covering fetch, decode, execute, and writeback.",
    "Describe the life cycle of a star, from protostar to main sequence to red giant and beyond.",
    "How does public key cryptography work? Explain RSA, including key generation, encryption, and signing.",
    "Write a dialogue between two characters debating whether artificial intelligence can be conscious.",
    "Explain the economic concepts of supply and demand, elasticity, and market equilibrium.",
    "What is CRISPR gene editing and how does it work? Explain Cas9, guide RNA, and applications.",
    "Describe the major causes and consequences of World War I.",
    "How does a compiler work? Explain lexing, parsing, semantic analysis, optimization, and code generation."
 )
 function Get-Selections {
    $u = Join-Path $Model ".coli_usage"
    if (-not (Test-Path $u)) { return 0 }
    $tot = 0
    Get-Content $u | ForEach-Object {
        $p = $_ -split '\s+'
        if ($p.Count -eq 3) { $tot += [int]$p[2] }
    }
    return $tot
 }
 $start = Get-Date
 $baseline = Get-Selections
 $line = "=" * 72
 "$line"                                  | Tee-Object -FilePath $Log -Append
 "colibri warmup - started $start"        | Tee-Object -FilePath $Log -Append
 "  model:    $Model"                     | Tee-Object -FilePath $Log -Append
 "  rounds:   $Rounds x $($Prompts.Count) prompts" | Tee-Object -FilePath $Log -Append
 "  ngen:     $Ngen tokens/prompt"        | Tee-Object -FilePath $Log -Append
 "  baseline: $baseline selections"       | Tee-Object -FilePath $Log -Append
 "$line"                                  | Tee-Object -FilePath $Log -Append
 $iter = 0
 $total = $Rounds * $Prompts.Count
 for ($r = 1; $r -le $Rounds; $r++) {
    for ($i = 0; $i -lt $Prompts.Count; $i++) {
        $iter++
        $prompt = $Prompts[$i]
        $now = Get-Date -Format "HH:mm:ss"
        $sel = Get-Selections
        $header = "[$now] round $r/$Rounds prompt {0,2}/$($Prompts.Count)  (iter $iter/$total)  selections: $sel" -f ($i+1)
        $header | Tee-Object -FilePath $Log -Append
        "  prompt: $($prompt.Substring(0, [Math]::Min(70, $prompt.Length)))..." | Tee-Object -FilePath $Log -Append
        $t0 = Get-Date
        # coli run writes status to stderr (normal) and may exit non-zero on
        # EOS-early; neither is a real failure for our purpose. Relax the
        # error preference and collect ALL output streams so stderr text
        # doesn't abort the loop.
        $prev = $ErrorActionPreference
        $ErrorActionPreference = "Continue"
        try {
            $output = & python $Coli run --model $Model --ngen $Ngen $prompt 2>&1 |
                      Select-Object -Last 4
        } catch {
            $output = @("  (engine run threw: $($_.Exception.Message))")
        }
        $ErrorActionPreference = $prev
        $elapsed = ((Get-Date) - $t0).TotalSeconds
        $after = Get-Selections
        $delta = $after - $sel
        $output | ForEach-Object { "  $_" | Tee-Object -FilePath $Log -Append }
        "  -> {0:N0}s, +{1} selections (now {2})" -f $elapsed, $delta, $after | Tee-Object -FilePath $Log -Append
        "" | Tee-Object -FilePath $Log -Append
    }
 }
 $end = Get-Date
 $final = Get-Selections
 $gain = $final - $baseline
 $duration = ($end - $start).ToString("hh\:mm\:ss")
 "$line"                                           | Tee-Object -FilePath $Log -Append
 "colibri warmup - finished $end"                  | Tee-Object -FilePath $Log -Append
 "  duration:    $duration"                        | Tee-Object -FilePath $Log -Append
 "  selections:  $baseline -> $final (+$gain)"     | Tee-Object -FilePath $Log -Append
 "  next: python coli chat --model $Model"         | Tee-Object -FilePath $Log -Append
 "$line"                                           | Tee-Object -FilePath $Log -Append
@@ -0,0 +1,55 @@
 # Metal backend on M5 Max — performance report + a tuning trap in the rebase (#72)
 **TL;DR:** After the rebase, the Metal backend is **faster than the pre-rebase branch — 2.24 vs 2.06 tok/s (+8.5%)** — *once tuned*. But the **default** rebased config regresses hard (1.25 tok/s, −39%), because the base now pulls in the OMP hot-team active-spin (#77), which on Apple Silicon steals the shared SoC power budget from the GPU and throttles the Metal kernels. Disabling the spin recovers the GPU; adding `PIPE` (new since the old base) then pushes *past* the old number. Recommend Metal builds default to a passive OMP wait.
 ## Setup (held fixed across every run)
 - **Hardware:** Apple M5 Max — 18 CPU cores (12 P + 6 E), 40-core GPU, 128 GB unified memory
 - **Model:** GLM-5.2 int4 (744B MoE), experts streamed from SSD
 - **Workload:** `./coli run "Compare the myths of Lucifer and Prometheus"`, 1024 tokens generated
 - **Constant flags:** `COLI_METAL=1 DIRECT=1 MTP=0 --ram 110`
 - Every run: identical working set — ~607 experts/token, ~74–75% hit rate, RSS ~97.9 GB, `fallback CPU 0` (all eligible blocks on GPU)
 ## Results
 | config | tok/s | wall (s) | expert-disk | expert-matmul | attention | attn GPU kernel | expert GPU overhead* |
 |---|---|---|---|---|---|---|---|
 | **A** — old base (pre-rebase branch) | 2.06 | 496 | 266 | 109 | 100 | 76 | 51 |
 | **B** — rebased, **defaults** | 1.25 | 819 | 285 | 215 | 290 | 223 | 106 |
 | **C** — rebased, defaults + `PIPE=1` | 1.30 | 788 | 297 | 190 | 272 | 212 | 89 |
 | **D** — rebased + `COLI_NO_OMP_TUNE=1` | 1.90 | 539 | 266 | 143 | 109 | 85 | 84 |
 | **E** — rebased + `NO_OMP` + `PIPE=1` **(winner)** | **2.24** | **457** | **241** | **97** | **99** | **79** | **44** |
 \* *expert GPU overhead = expert gpu-wall − expert kernel time, i.e. GPU idle waiting to be fed. Expert kernel time itself was constant (~34–35 s) in every un-throttled run.*
 Winner (E) command:
 ```bash
 COLI_METAL=1 DIRECT=1 MTP=0 COLI_NO_OMP_TUNE=1 PIPE=1 PIPE_WORKERS=8 \
  ./coli run --model /path/to/glm52_i4 \
  "Compare the myths of Lucifer and Prometheus" --ram 110
 ```
 ## What happened, phase by phase
 **The default regression (A→B) is not in the kernels' work — it's GPU throttling.** Same GPU dispatch counts (≈140k blocks, 79,794 attention layer-launches, 618k experts-on-GPU) in every run, yet the **attention GPU kernel time triples, 76 s → 223 s.** Kernel *execution* time can't triple for identical work unless the GPU is clocking down. The cause is the OMP hot-team (#77): with active spin, the CPU sits at ~97% busy-waiting on the GPU, and because the M-series CPU and GPU share one power/thermal envelope, that spin robs the GPU of clock headroom. Note B→C: adding `PIPE` while the spin is still on barely helps (1.25→1.30) — the CPU has no spare cycles to run PIPE's workers.
 **Fix 1 — kill the spin (A→D).** `COLI_NO_OMP_TUNE=1` (passive waits) gives the GPU its power back: attention kernel falls 223 s → 85 s, near the old 76 s, and throughput jumps to 1.90. But a residual gap remains, and it is entirely in **expert-matmul (+33 s vs A)** — specifically the *GPU overhead*, which grew 51 s → 84 s while the expert kernel stayed at ~34 s. Passive waits stop stealing power but make the CPU slower to hand the next expert batch to the GPU, so the GPU idles between dispatches.
 **Fix 2 — hide the dispatch latency (D→E).** `PIPE=1 PIPE_WORKERS=8` keeps experts prefetched and streaming, so the GPU stops waiting: expert overhead 84 s → **44 s** (below even the old base's 51 s), and `PIPE`'s I/O overlap also trims the disk wall 266 s → **241 s**. Net **2.24 tok/s** — past the pre-rebase 2.06, because `PIPE` did not exist on the old base.
 The two levers are complementary, not additive coincidences: **`NO_OMP` restores GPU clocks; `PIPE` hides the CPU→GPU dispatch latency that `NO_OMP` introduces.** You want both.
 ## Recommendations
 1. **On Apple Silicon Metal builds, default OMP to a passive wait** (e.g. auto-set `OMP_WAIT_POLICY=passive`, or `COLI_NO_OMP_TUNE` behavior, when `COLI_METAL` is enabled). The #77 active-spin default is a −39% trap here: during Metal offload the CPU is mostly *waiting on the GPU*, and spinning actively throttles it. This is the single highest-impact change.
 2. **Document `PIPE=1 PIPE_WORKERS=8` as recommended with Metal** — it recovers the dispatch-latency cost of the passive wait and overlaps expert I/O.
 3. **Memory ceiling (128 GB M5 Max):** `--ram 110` is safe (RSS ~97 GB, compressor quiet); `--ram 120` crosses into memory compression (double penalty — compressor CPU cost *and* SoC power stolen from the GPU); `--ram 115` untested/marginal. Metal's registered buffers share unified memory, so the knee is lower than a CPU-only build would suggest.
 ## Caveats / untested levers
 - **Decode only.** These are steady-state decode numbers. The cold-prefill wall is unaffected (fused attention covers S≤4).
 - **`MTP=0` throughout.** The rebased kernel's "fused attention handles S≤4 (covers MTP verify forwards)" commit means MTP verify is now GPU-accelerated — MTP was a net loss on the old base, so re-testing `MTP` on is a live, unexplored lever.
 - **`DIRECT=1` is required, not optional.** `DIRECT=0` was A/B'd and is ~2× slower (2.16 → 1.15 tok/s at the same point in the run). It also drops RSS 98 → 84 GB — reads fall back to the OS page cache, which lowers process RSS but breaks the zero-copy GPU slab registration `DIRECT=1` enables, adding a copy to feed the GPU.
 - **Determinism (confirmed non-deterministic run-to-run):** Two identical `DIRECT=1` runs (same config, same prompt, greedy / MTP off) diverged within ~7 tokens. So the engine is not run-to-run reproducible under the parallel config, and the earlier `DIRECT=0` vs `DIRECT=1` divergence is a symptom of this, not a `DIRECT` read-path bug. Cause is expected and benign: floating-point non-associativity in parallel expert-sum reductions (PIPE worker completion order and/or GPU threadgroup reductions) occasionally flips an argmax at a token boundary. Output quality is unaffected (both completions are valid) and **throughput is identical** across runs (1.28 tok/s at the same point), so benchmark numbers are stable. Implication for the PR's **"token-exact"** claim: it holds as "GPU path matches the CPU reference under a deterministic/serial validation config," but is *not* "bit-reproducible across runs" with `PIPE`/threads enabled — worth stating so nobody diffs two runs and files a false bug.
 - Numbers are single-run per config on a warm cache; run-to-run and thermal/ambient variation not bounded (a same-session A/B of old vs rebased commit would tighten the throttling claim further).
@@ -1 +1,30 @@
-{"prompt": "The capital of France is", "prompt_ids": [510, 5347, 273, 6181, 310], "full_ids": [510, 5347, 273, 6181, 310, 7785, 15, 187, 187, 510, 5347, 273, 253, 1986, 2077, 310, 5041], "text": "The capital of France is Paris.\n\nThe capital of the United States is Washington"}
+{
  "prompt": "The capital of France is",
  "prompt_ids": [
    510,
    5347,
    273,
    6181,
    310
  ],
  "full_ids": [
    510,
    5347,
    273,
    6181,
    310,
    7785,
    15,
    187,
    187,
    510,
    3565,
    3448,
    273,
    6181,
    310,
    5112,
    15
  ],
  "text": "The capital of France is Paris.\n\nThe official language of France is French."
 }
Author	SHA1	Message	Date
cody	6989590ab1	perf: znver5 + LTO + OMP 32T + GPU DENSE/ASYNC (Strix Halo optimiert)	2026-07-14 23:56:35 +02:00
cody	53c37ec918	Initial: Colibri upstream + HIP/ROCm port (Strix Halo / Radeon 8060S)	2026-07-14 23:53:20 +02:00
JustVugg	6d3ed7e62b	README: GB10 (DGX Spark) row (#136 ) — 0.50 tok/s warm, matmul-bound, unified memory; confirms #76 fix on aarch64 (TF 32/32 + greedy 20/20 on NEON and CUDA sm_121) Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>	2026-07-13 21:13:40 +02:00
JustVugg	086b2dfb87	README: native-Windows warm-cache row (#128 , 0.03->0.5 tok/s) — honest label per #131 (pipe+RAM fixes, no GPU) Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>	2026-07-13 21:00:50 +02:00
woolcoxm	2319b942d2	Windows native port: serve-mode pipe fix + RAM detection + POSIX guards, AVX-VNNI kernel, gated CUDA DLL (#131 , fixes #123 ) Rebased onto current dev, split into 3 logical parts (all validated): 1. CPU portability (serve-mode _O_BINARY pipe fix — stock main hangs on MinGW without it; RAM detection cap 0->9/layer; POSIX guards for select/mmap/madvise; warmup script). 2. AVX-VNNI 128-bit int8/int4 dot kernel (Alder Lake+/Meteor Lake+), bit-identical to AVX2 (author-verified on Meteor Lake; compiles out to AVX2 elsewhere) + _mm256_extracti128_si256 typo fix that blocked -march=native. 3. CUDA DLL via LoadLibrary, gated behind CUDA_DLL=1 (host never links cudart; silent CPU fallback if absent; author-verified on RTX 5070 Ti). Validated here: make check 59/59, oracle 32/32 TF, Windows cross-compile clean + glm.exe loads+runs via WSL interop. Fixes the #123 Windows build failure.	2026-07-13 20:54:30 +02:00
bopof	afc259c599	Makefile: detect Windows from $(OS)=Windows_NT before uname so make works from native PowerShell/CMD (#129 ) On Windows $(OS) is Windows_NT in every shell; check it first so native PowerShell/CMD (no uname on PATH) doesn't fall through to the Linux branch. Non-Windows unchanged (else branch still uses uname). Linux build verified green.	2026-07-13 20:47:23 +02:00
RDouglas	5dd7503ee7	docs: Metal M5 Max perf report — OMP hot-team spin (#77 ) throttles Apple GPU, NO_OMP+PIPE recovers to 2.24 tok/s (#116 ) Docs-only. Documents that the OMP active-spin steals SoC power from the Metal GPU on Apple Silicon (default regresses -39%); COLI_NO_OMP_TUNE=1 + PIPE=1 recovers and beats the pre-rebase branch (2.24 vs 2.06 tok/s). Flags a follow-up: Metal builds should default to passive OMP wait.	2026-07-13 20:47:19 +02:00
ZacharyZcR	f8c0552c6d	quant_ablation: add rotation preconditioning (-rot, QuaRot/QuIP#) and int3 schemes (#132 , #81 ) Extends the ablation grammar to int{2,3,4,8}[-g<N>][-rot][-nohead]. -rot round-trips weights through an orthogonal Hadamard Q=diag(±1)·H/√n on the input dim, measuring the exact weight error of a deployed rotate-activations scheme. Engine-free tool (c/tools/quant_ablation.py only). Verified: syntax clean, scheme parser correct (int3/-rot/-nohead), no unsafe constructs. Findings feed the v2 int3-g64 direction (#81/#108).	2026-07-13 20:45:34 +02:00
JustVugg	1f0e1b7076	olmoe: reject quant_bits outside 2..8 (fixes degenerate output, #134 ) + correct ref.json to OLMoE-0125-Instruct oracle (#133 ) #134: olmoe.c stored experts as int8_t but silently accepted any bits argv; bits=16 (falsely documented as f32) truncated in quantize_rows -> wraparound garbage experts. Guard bits to 2..8 with a clear error (int8 is token-exact; f32 experts are not implemented here, unlike glm.c's fmt=0 path). Doc corrected. #133: shipped ref.json was from a different checkpoint (continued 'The capital of the United States is Washington'); the intended target is allenai/OLMoE-1B-7B-0125-Instruct (per convert_olmoe.py), whose greedy oracle continues 'The official language of France is French'. full_ids/text updated to the reporter's verified oracle (engine is already token-exact 12/12 vs it). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>	2026-07-13 20:33:47 +02:00
JustVugg	4418e16db4	README: benchmark rows for Gen5-NVMe 9950X3D2+5090 (1.23 tok/s, #120 ) and first Strix Halo datapoint (1.10 tok/s, #124 ) Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>	2026-07-13 20:26:52 +02:00