diff --git a/.gitignore b/.gitignore
index 203a322..981e740 100644
--- a/.gitignore
+++ b/.gitignore
@@ -31,6 +31,10 @@ c/tests/test_tier
c/tests/test_tier.exe
c/tests/test_grammar
c/tests/test_grammar.exe
+c/tests/test_schema_gbnf
+c/tests/test_schema_gbnf.exe
+c/tests/test_compat_direct
+c/tests/test_compat_direct.exe
# oracoli tiny generati (make_glm_oracle.py) e dati benchmark scaricati
c/glm_tiny/
@@ -60,3 +64,7 @@ c/backend_metal_test
c/tests/test_tier
c/mio_env/
c/bench/
+c/tests/test_decode_batch
+c/tests/test_i4_acc512
+c/tests/test_idot
+c/tests/test_uring
diff --git a/README.md b/README.md
index f72c4e9..2e22048 100644
--- a/README.md
+++ b/README.md
@@ -17,14 +17,40 @@ $ ./coli chat
◆ Ciao! 😊 Come posso aiutarti oggi?
```
+
+## See it running
+
+
+
+
+The web dashboard (./coli web): a 744B model answering at 4+ tok/s end-to-end on 6× RTX 5090 —
+with live token metrics, the hardware panel, and the VRAM/RAM/disk expert tiers.
+
+
+
+
+The Brain page: all 19,456 experts as a living cortex — colour is the storage tier,
+brightness is routing heat, and every expert routed in a turn flashes white. Hovering shows the expert's
+measured topic affinity.
+
+## Contents
+
+- [The idea](#the-idea)
+- [See it running](#see-it-running)
+- [What's implemented](#whats-implemented)
+- [Honest numbers](#honest-numbers-wsl2-12-cores-25-gb-ram-nvme-via-vhdx)
+- [Download the model](#download-the-model)
+- [Web dashboard](#web-dashboard)
+- [Got a better machine?](#got-a-better-machine-try-it--heres-what-to-expect)
+
## The idea
A 744B Mixture-of-Experts model activates only ~40B parameters per token — and only ~11 GB of those change from token to token (the routed experts). So:
- the **dense part** (attention, shared experts, embeddings — ~17B params) stays **resident in RAM at int4** (~9.9 GB);
-- the **21,504 routed experts** (75 MoE layers × 256 experts + the MTP head, ~19 MB each at int4) live **on disk** (~370 GB) and are **streamed on demand**, with a per-layer LRU cache, an optional pinned hot-store, and the OS page cache as a free L2.
+- the **19,456 routed experts** (75 MoE layers × 256 experts + the MTP head, ~19 MB each at int4) live **on disk** (~370 GB) and are **streamed on demand**, with a per-layer LRU cache, an optional pinned hot-store, and the OS page cache as a free L2.
-The engine is a single C file (`c/glm.c`, ~2,400 lines) plus small headers. No BLAS, no Python at runtime, no GPU required (an opt-in CUDA tier for pinned experts exists — see below).
+The engine is a single C file (`c/glm.c`) plus small headers. No BLAS, no Python at runtime, no GPU required (an opt-in CUDA tier for pinned experts exists — see below).
## What's implemented
@@ -354,6 +380,10 @@ PIN=stats.txt PIN_GB=160 SNAP=/nvme/glm52_i4 ./glm 64 4 4
COLI_CUDA=1 COLI_GPUS=0,1,2,3,4,5 CUDA_EXPERT_GB=150 \
CUDA_DENSE=1 PIN=stats.txt PIN_GB=300 RAM_GB=226 \
SNAP=/nvme/glm52_i4 ./glm 64 4 4
+# large-RAM host: fill safe VRAM, then keep every remaining expert in RAM
+COLI_CUDA=1 COLI_GPUS=0,1,2,3,4,5 CUDA_EXPERT_GB=auto \
+CUDA_DENSE=1 COLI_CUDA_ATTN=1 PIN=stats.txt PIN_GB=all RAM_GB=auto \
+SNAP=/nvme/glm52_i4 ./glm 64 4 4
```
Selected experts are uploaded during startup, so capacity failures occur before
@@ -375,6 +405,25 @@ replay from 1.87 to 2.16 tok/s (+15.7%), reduced expert disk wait from 5.144s to
3.948s, and kept the projected RAM peak below `RAM_GB=226`. The cache cap adjusts
down automatically (54 to 40 in that run) so the larger pinned tier does not exceed
the process budget. Start lower on hosts with less available RAM.
+
+`CUDA_EXPERT_GB=auto` fills each selected device only up to its measured free
+memory minus projected dense tensors and 2 GB of runtime headroom. `PIN_GB=all`
+then loads every remaining routed expert into RAM, eliminating decode-time disk
+misses when the host budget permits it. The regular `RAM_GB` guard still clamps
+the per-layer working cache and rejects unsafe projections; this mode is intended
+for dedicated high-memory inference hosts, not desktops running other workloads.
+On a dedicated 251 GiB host with six RTX 5090s, this mode selected a 176.7 GB
+VRAM expert tier and a 191.3 GB RAM tier (all 19,456 experts resident). The
+mode also adapts the VRAM tier every 16 emitted tokens by swapping hot RAM
+experts into existing GPU slots. A real 64-token greedy GLM-5.2 generation
+measured **6.00 tok/s decode**, up from
+2.20 tok/s end-to-end with the earlier 150 GB tier; expert hit rate was 100%
+and disk wait was zero. Prompt prefill is reported separately. This is a
+host-specific capacity result, not a portable default.
+
+Text-mode timing reports prefill separately from decode. The decode rate starts
+after the prompt KV is built, so it is comparable to `REPLAY` throughput without
+hiding time-to-first-token.
MTP speculation defaults off on CUDA because cold draft routes increase expert
traffic; an explicit `DRAFT=n` still overrides the default.
@@ -406,8 +455,8 @@ CUDA, CPU hot-store, and CUDA hot-expert execution with identical replay tokens.
### Web interface
-`web/` contains a community-contributed browser UI (React + TypeScript, ~390
-lines of source, a pure API client — it never touches the engine directly):
+`web/` contains a community-contributed browser UI (React + TypeScript, a pure
+API client — it never touches the engine directly):
```bash
cd web
@@ -419,7 +468,7 @@ works against the colibrì OpenAI-compatible server (in review, #21) or any othe
compatible endpoint. Nothing leaves the endpoint you configure. The terminal
`coli chat` remains the first-class interface.
-Useful knobs (env or flags): `--temp T` token sampling temperature (default 0.7 + nucleus 0.90 — tuned for int4; 0 = greedy), `--topp 0.7` adaptive expert top-p (30–40% less disk), `--ngen N` max tokens per answer (`:more` in chat continues a truncated one), `--repin N` adapt RAM/VRAM hot experts every N emitted tokens, `AUTOPIN=0` disable the learning cache's auto-pin, `THINK=1` enable GLM-5.2's reasoning block, `DRAFT=n` MTP draft depth, `GRAMMAR=g.gbnf` grammar-forced drafts for constrained JSON/NDJSON output (`GRAMMAR_DRAFT=n` caps the forced span), `TF=1` teacher-forcing validation, `PILOT=1` router-lookahead disk prefetch (experimental — see below), `CAP_RAISE=0` don't auto-grow the expert cache.
+Useful knobs (env or flags): `--temp T` token sampling temperature (default 0.7 + nucleus 0.90 — tuned for int4; 0 = greedy), `--topp 0.7` adaptive expert top-p (30–40% less disk), `--ngen N` max tokens per answer (`:more` in chat continues a truncated one), `--repin N` adapt RAM/VRAM hot experts every N emitted tokens, `AUTOPIN=0` disable the learning cache's auto-pin, `THINK=1` enable GLM-5.2's reasoning block, `DRAFT=n` MTP draft depth, `GRAMMAR=g.gbnf` grammar-forced drafts for constrained JSON/NDJSON output (`GRAMMAR_DRAFT=n` caps the forced span), `TF=1` teacher-forcing validation, `PILOT=1` router-lookahead disk prefetch (experimental — see below), `URING=1` Linux-only batched expert I/O (implies `PIPE=1`; also batches `PILOT_REAL`), `PIPE=0` disable the async expert-load pool (**default ON on Windows** — overlaps expert `pread` with the matmul so the CPU isn't idle waiting on the SSD; measured −18% disk service time), `RAM_GB=` claim more RAM for the expert cache than the conservative auto-detect (e.g. `RAM_GB=31` on a 32 GB host raises the cache cap and hit rate measurably), `CAP_RAISE=0` don't auto-grow the expert cache.
### Resource policy
@@ -444,10 +493,10 @@ Disk is an immutable recovery source, not a normal decode target. If the plan
leaves cold expert bytes on disk, speed depends on cache hit rate; output
quality does not.
-Cold expert reads use a deferred pipeline: resident RAM/VRAM experts execute
+Cold expert reads can use a deferred pipeline: resident RAM/VRAM experts execute
while missing experts are loaded in a bounded background I/O pool, then the
-cold results join before the layer completes. `IO_THREADS=n` overrides the
-default eight loader threads when foreground work exists. Profiling reports
+cold results join before the layer completes. The pool engages only under
+`PIPE=1`; `PIPE_WORKERS=n` sets its worker count (default 8). Profiling reports
both disk service time and the smaller foreground-visible wait time so overlap
is explicit rather than credited as unexplained speedup.
@@ -478,6 +527,23 @@ thrashing. Persistent `.coli_usage` remains the long-term signal and is not deca
**Conversations reopen warm** (`.coli_kv`, since 2026-07-10): `coli chat` persists the compressed MLA KV-cache to disk after every turn (~182 KB/token, appended incrementally, crash-safe). Close the chat, reopen it tomorrow — the model still remembers the whole conversation and **zero re-prefill happens**: validated byte-identical to an uninterrupted session. `:reset` clears it, `KVSAVE=0` disables it.
+## Web dashboard
+
+One command serves the OpenAI-compatible API **and** the web console on the same port, then opens your browser when the engine is ready:
+
+```bash
+cd web && npm install && npm run build # once
+./coli web --model
+```
+
+What you get:
+
+- **Chat** with live metrics: a flashing token counter while generating, then tok/s, time-to-first-token, prompt→completion counts and queue wait;
+- **Runtime panel**: your hardware (CPU, GPUs + VRAM, RAM, cores), the scheduler, and the live expert-tier bar — how many of the 19,456 experts sit in VRAM / RAM / disk right now;
+- **Brain**: the whole model as a 76×256 cortex, one cell per expert. Colour = tier, brightness = routing heat, and the experts routed in each turn flash white and decay — you watch the model think. Hover any cell for its tier, heat and [measured topic affinity](https://github.com/JustVugg/colibri/issues/175) (specialists for code, Chinese, math, law… live in layers 11–22).
+
+The dashboard talks to the engine over two tiny protocol lines (`TIERS`, `EMAP`/`HITS`) and plain JSON endpoints — nothing heavier than the engine itself.
+
## Got a better machine? Try it — here's what to expect
colibrì was built on deliberately humble hardware (12 cores, 25 GB RAM, an older DRAM-less NVMe behind a WSL2 VHDX that measured ~1 GB/s random on *this* drive — note WSL2 VHDX is not inherently slow: a community 5090 box measured 10.5 GB/s O_DIRECT through one, [#101](https://github.com/JustVugg/colibri/issues/101)). **Every one of those constraints is a knob your machine can turn up.** The engine needs: Linux (or WSL2), macOS, or **Windows 11 natively (MinGW-w64)**; gcc with OpenMP, AVX2, ≥16 GB RAM, and the ~370 GB int4 model on a local NVMe (ext4/NTFS — never a network/9p mount).
@@ -591,6 +657,7 @@ c/
├── scripts/ long-running conversion helpers
└── tests/ dependency-free C and Python tests
web/ browser UI (pure OpenAI-API client, community-maintained)
+desktop/ Tauri v2 desktop shell wrapping the web UI
```
The runtime path intentionally stays flat and readable: `glm.c` plus its small
diff --git a/c/Makefile b/c/Makefile
index c1e9319..5ae7a9c 100644
--- a/c/Makefile
+++ b/c/Makefile
@@ -14,7 +14,7 @@ TRIPLET := $(shell $(DETECT_CC) -dumpmachine 2>/dev/null)
MINGW := $(findstring mingw,$(TRIPLET))
CYGWIN := $(findstring cygwin,$(TRIPLET))
DARWIN := $(findstring darwin,$(TRIPLET))
-PPC64 := $(findstring powerpc64,$(TRIPLET))
+LINUX := $(findstring linux,$(TRIPLET))
IS_WIN := $(MINGW)$(CYGWIN)
# Fallbacks for the rare toolchain that does not answer -dumpmachine: keep the
@@ -26,10 +26,17 @@ ifeq ($(IS_WIN),)
UNAME_S := $(shell uname -s)
UNAME_M := $(shell uname -m)
DARWIN := $(findstring Darwin,$(UNAME_S))
-PPC64 := $(findstring ppc64,$(UNAME_M))
endif
endif
+TARGET_CPU := $(firstword $(subst -, ,$(TRIPLET)))
+ifeq ($(TARGET_CPU),)
+TARGET_CPU := $(UNAME_M)
+endif
+X86_64 := $(filter x86_64 amd64,$(TARGET_CPU))
+AARCH64 := $(filter aarch64 arm64,$(TARGET_CPU))
+PPC64 := $(filter powerpc64% ppc64%,$(TARGET_CPU))
+
ifneq (,$(DARWIN))
# --- macOS / Apple Silicon ---
# Apple clang non include il runtime OpenMP: se c'e' libomp di Homebrew lo usa
@@ -37,7 +44,9 @@ ifneq (,$(DARWIN))
# Niente -march: su arm64 NEON e' baseline (i kernel __ARM_NEON si attivano da soli).
CC = clang
OMPDIR := $(shell brew --prefix libomp 2>/dev/null)
-ifneq ($(OMPDIR),)
+# `brew --prefix libomp` can print the formula's prospective path even when it
+# is not installed, so verify both artifacts before adding unusable flags.
+ifneq ($(and $(OMPDIR),$(wildcard $(OMPDIR)/include/omp.h),$(wildcard $(OMPDIR)/lib/libomp.*)),)
OMPC = -Xclang -fopenmp -I$(OMPDIR)/include
OMPL = -L$(OMPDIR)/lib -lomp
else
@@ -46,6 +55,12 @@ OMPC =
OMPL =
endif
CFLAGS = -O3 $(OMPC) -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function
+# Opt-in: ARCH=native appends -mcpu=native (arm64 clang uses -mcpu, not -march),
+# which unlocks the i8mm SMMLA int8/int4 dot kernels in glm.c. ARCH unset ->
+# no -mcpu, default build byte-identical. Apple clang knows apple-m4 / native.
+ifneq ($(ARCH),)
+CFLAGS += -mcpu=$(ARCH)
+endif
LDFLAGS = -lm $(OMPL)
EXE =
else ifneq ($(IS_WIN),)
@@ -60,7 +75,11 @@ else ifneq ($(IS_WIN),)
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
-LDFLAGS = -lm -fopenmp -static
+# -lpsapi: compat.h calls GetProcessMemoryInfo (rss_gb). It's linked via
+# #pragma comment(lib,"psapi.lib") for MSVC, but MinGW gcc ignores that pragma
+# (warns), so link psapi explicitly or the build fails undefined-reference on
+# modern GCC (e.g. 16.x under UCRT). Harmless where the pragma also resolves it.
+LDFLAGS = -lm -fopenmp -static -lpsapi
EXE = .exe
else
ifneq (,$(PPC64))
@@ -70,18 +89,19 @@ ifneq (,$(PPC64))
# (validated token-exact vs the transformers oracle on a POWER8 S824).
CC = gcc
ARCH ?= native
-CFLAGS = -O3 -mcpu=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function
-LDFLAGS = -lm -fopenmp
+CFLAGS = -O3 -mcpu=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function
+LDFLAGS = -lm -fopenmp -pthread
EXE =
else
-# --- Linux x86-64 (percorso originale, invariato) ---
+# --- Linux / *BSD x86-64 (percorso originale, invariato) ---
CC = gcc
# ARCH=native -> ottimizzato per QUESTA macchina (default, piu' veloce).
# ARCH=x86-64-v3 -> binario PORTABILE su qualsiasi x86-64 moderno con AVX2 (per distribuire).
# ARCH=x86-64 -> massima compatibilita' (niente AVX2: usa il path scalare di fallback).
+# -pthread: Linux lo tira dentro via -fopenmp, i *BSD no e le pthread_* non risolvono (#219).
ARCH ?= native
-CFLAGS = -O3 -march=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function
-LDFLAGS = -lm -fopenmp
+CFLAGS = -O3 -march=$(ARCH) -fopenmp -pthread -Wall -Wextra -Wno-unused-parameter -Wno-misleading-indentation -Wno-unused-function
+LDFLAGS = -lm -fopenmp -pthread
EXE =
endif
endif
@@ -105,22 +125,51 @@ INSTALL ?= install
# clean. See backend_loader.c and README "cuda-dll" below.
CUDA ?= 0
CUDA_DLL ?= 0
+ifneq ($(IS_WIN),)
+# the CUDA installer sets CUDA_PATH system-wide (e.g. C:\Program Files\NVIDIA
+# GPU Computing Toolkit\CUDA\v13.2); fall back to it before the POSIX default.
+# NVCC defaults to plain `nvcc` from PATH: CUDA_PATH contains spaces, which the
+# unquoted recipe checks cannot survive — and the cuda-dll recipe already
+# requires an MSVC environment (vcvars64) on PATH, so requiring the CUDA bin
+# directory too is symmetric. The installer adds it by default.
+CUDA_HOME ?= $(subst \,/,$(CUDA_PATH))
+NVCC ?= nvcc
+# Host compiler override for nvcc (e.g. Fedora ships g++16 but CUDA needs 15):
+# make glm CUDA=1 NVCC_CCBIN=g++-15 (opt-in: unset = nvcc's own default)
+ifneq ($(NVCC_CCBIN),)
+NVCCFLAGS += -ccbin $(NVCC_CCBIN)
+endif
+else
CUDA_HOME ?= /usr/local/cuda
NVCC ?= $(CUDA_HOME)/bin/nvcc
+endif
CUDA_ARCH ?= native
-# -Xcompiler passes flags to nvcc's HOST compiler. On Windows that host is MSVC
-# cl.exe, which rejects the gcc/clang-style -Wall,-Wextra (D8021). Gate them
-# behind a non-MSVC host so the Linux/macOS (gcc/clang) build keeps the warnings
-# and the Windows cuda-dll build doesn't fail. (#306)
-NVCC_HOST_WARN ?= $(if $(IS_WIN),,-Xcompiler=-Wall,-Wextra)
-NVCCFLAGS ?= -O3 -std=c++17 -arch=$(CUDA_ARCH) $(NVCC_HOST_WARN)
-ifeq ($(IS_WIN),)
-PYTHON ?= python3
+ifneq ($(IS_WIN),)
+# nvcc's host compiler on Windows is MSVC cl.exe: the GCC-style
+# -Xcompiler=-Wall,-Wextra is rejected ("D8021 invalid numeric argument
+# '/Wextra'"), which made `make cuda-dll` unbuildable as shipped. -W3 is
+# the MSVC warning level (dash form, NOT /W3: MSYS make would mangle the
+# slash-form into a filesystem path).
+NVCCFLAGS ?= -O3 -std=c++17 -arch=$(CUDA_ARCH) -Xcompiler=-W3
else
+NVCCFLAGS ?= -O3 -std=c++17 -arch=$(CUDA_ARCH) -Xcompiler=-Wall,-Wextra
+endif
+# PYTHON is a HOST tool (clean/test-c/test-python run it on the build machine),
+# so key it off the host, NOT $(IS_WIN) — which is derived from the *target*
+# triple ($(CC) -dumpmachine). Otherwise a cross build (make CC=x86_64-w64-
+# mingw32-gcc ... on Linux) sets IS_WIN and picks `python`, breaking clean/test
+# on hosts where only `python3` exists. $(OS)=Windows_NT in every Windows shell
+# (the #129 signal) and is empty on Linux/macOS.
+ifeq ($(OS),Windows_NT)
PYTHON ?= python
+else
+PYTHON ?= python3
endif
CUDA_OBJ =
-TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE)
+TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE)
+ifneq (,$(LINUX))
+TEST_BINS += tests/test_uring$(EXE)
+endif
# Windows CUDA DLL path: host links the loader, NOT cudart.
ifneq ($(IS_WIN),)
@@ -148,7 +197,7 @@ endif
# runtime, so no Xcode / offline metal compiler is required. Default build unchanged.
METAL ?= 0
METAL_OBJ =
-METALXX = clang++ -x objective-c++ -fobjc-arc -O3
+METALXX = clang++ -x objective-c++ -std=gnu++17 -fobjc-arc -O3
ifeq ($(METAL),1)
ifeq (,$(DARWIN))
$(error METAL=1 is supported only on macOS)
@@ -177,7 +226,7 @@ $(shell printf '%s' '$(BUILD_CONFIG)' > .build-config)
endif
.build-config: ;
-glm$(EXE): glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h $(CUDA_OBJ) $(METAL_OBJ) .build-config
+glm$(EXE): glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h $(CUDA_OBJ) $(METAL_OBJ) .build-config
$(CC) $(CFLAGS) glm.c $(CUDA_OBJ) $(METAL_OBJ) -o glm$(EXE) $(LDFLAGS)
# Windows runtime loader object: resolves coli_cuda_* from coli_cuda.dll.
@@ -188,7 +237,7 @@ backend_loader.o: backend_loader.c backend_cuda.h compat.h .build-config
# 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.
+# __declspec(dllexport) so the 15 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; }
@@ -220,9 +269,23 @@ cuda-bench: backend_cuda.cu backend_cuda.h tests/bench_tensor_core.cu
olmoe$(EXE): olmoe.c st.h json.h compat.h
$(CC) $(CFLAGS) olmoe.c -o olmoe$(EXE) $(LDFLAGS)
-# binario portabile da distribuire su altre macchine x86-64
+# Use a baseline that matches the compiler target. macOS already targets a
+# portable baseline when ARCH is empty; forcing the x86 value there breaks
+# Apple Silicon. Unknown targets use native rather than an invalid x86 flag.
+ifneq (,$(DARWIN))
+PORTABLE_ARCH =
+else ifneq (,$(AARCH64))
+PORTABLE_ARCH = armv8-a
+else ifneq (,$(PPC64))
+PORTABLE_ARCH = power8
+else ifneq (,$(X86_64))
+PORTABLE_ARCH = x86-64-v3
+else
+PORTABLE_ARCH = native
+endif
+
portable:
- $(MAKE) glm$(EXE) ARCH=x86-64-v3
+ $(MAKE) glm$(EXE) ARCH=$(PORTABLE_ARCH)
iobench$(EXE): iobench.c compat.h
$(CC) $(CFLAGS) iobench.c -o iobench$(EXE) $(LDFLAGS)
@@ -239,10 +302,16 @@ tests/test_tier$(EXE): tests/test_tier.c tier.h
tests/test_grammar$(EXE): tests/test_grammar.c grammar.h
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
+tests/test_schema_gbnf$(EXE): tests/test_schema_gbnf.c schema_gbnf.h grammar.h json.h
+ $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
+
tests/test_decode_batch$(EXE): tests/test_decode_batch.c decode_batch.h
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
-tests/test_idot$(EXE): tests/test_idot.c glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h
+tests/test_idot$(EXE): tests/test_idot.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h
+ $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
+
+tests/test_kv_alloc$(EXE): tests/test_kv_alloc.c glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
tests/test_i4_acc512$(EXE): tests/test_i4_acc512.c
@@ -251,6 +320,9 @@ tests/test_i4_acc512$(EXE): tests/test_i4_acc512.c
tests/test_compat_direct$(EXE): tests/test_compat_direct.c compat.h
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
+tests/test_uring$(EXE): tests/test_uring.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h
+ $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
+
test-c: $(TEST_BINS)
$(PYTHON) tools/run_tests.py $(TEST_BINS)
@@ -284,4 +356,4 @@ clean:
bench: iobench$(EXE)
@if [ -n "$(ARGS)" ]; then ./iobench$(EXE) $(ARGS); else echo "built iobench$(EXE) — run: ./iobench$(EXE) "; fi
-.PHONY: all glm cuda-test cuda-bench cuda-dll portable test-c test-python test check clean install uninstall bench
\ No newline at end of file
+.PHONY: all glm cuda-test cuda-bench cuda-dll portable test-c test-python test check clean install uninstall bench
diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu
index d09a9fa..9ce142f 100644
--- a/c/backend_cuda.cu
+++ b/c/backend_cuda.cu
@@ -18,12 +18,14 @@ struct ColiCudaTensor {
typedef struct {
int device;
+ int compute_major,compute_minor;
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;
+ float *pipe_buf[24]; size_t pipe_cap[24]; /* scratch persistenti del resident pipeline */
cudaStream_t stream;
void *group_desc; size_t group_desc_cap;
size_t tensor_count, tensor_bytes;
@@ -120,6 +122,68 @@ __global__ static void silu_mul(float *gate, const float *up, size_t n) {
}
}
+/* Four warps share one A tile and compute 16x64 outputs. This matters for
+ * prefill: the first prototype reloaded/converter A once per 16 output cols. */
+__global__ static void w4a16_matmul(float *y,const float *x,const uint8_t *w,
+ const float *scale,int M,int K,int N){
+#if __CUDA_ARCH__ >= 700
+ using namespace nvcuda;int warp=threadIdx.x>>5,lane=threadIdx.x&31;
+ int m0=blockIdx.y*16,n0=blockIdx.x*64+warp*16;
+ __shared__ __half ah[256],bh[4][256];
+ wmma::fragment acc;wmma::fill_fragment(acc,0.f);
+ size_t rb=(size_t)(K+1)/2;
+ for(int k0=0;k0>1)];int a=(gk&1)?q>>4:q&15;
+ v=(float)(a&8?a-16:a)*scale[gn];}
+ bh[warp][z]=__float2half(v); /* [Ntile,Ktile] == B col-major */
+ }
+ __syncthreads();
+ wmma::fragment af;
+ wmma::fragment bf;
+ wmma::load_matrix_sync(af,ah,16);wmma::load_matrix_sync(bf,bh[warp],16);
+ wmma::mma_sync(acc,af,bf,acc);__syncthreads();
+ }
+ __shared__ float out[4][256];wmma::store_matrix_sync(out[warp],acc,16,wmma::mem_row_major);__syncwarp();
+ for(int z=lane;z<256;z+=32){int m=z/16,n=z%16;
+ if(m0+mFP16 conversion of A. */
+__global__ static void w4a16_gate_up(float *gate,float *up,const float *x,
+ const uint8_t *gw,const uint8_t *uw,const float *gs,const float *us,
+ int M,int K,int N){
+#if __CUDA_ARCH__ >= 700
+ using namespace nvcuda;int warp=threadIdx.x>>5,lane=threadIdx.x&31,which=warp&1,tile=warp>>1;
+ int m0=blockIdx.y*16,n0=blockIdx.x*64+tile*16;const uint8_t *w=which?uw:gw;
+ const float *scale=which?us:gs;float *y=which?up:gate;size_t rb=(size_t)(K+1)/2;
+ __shared__ __half ah[256],bh[8][256];
+ wmma::fragment acc;wmma::fill_fragment(acc,0.f);
+ for(int k0=0;k0>1)];int a=(gk&1)?q>>4:q&15;
+ v=(float)(a&8?a-16:a)*scale[gn];}bh[warp][z]=__float2half(v);}
+ __syncthreads();
+ wmma::fragment af;
+ wmma::fragment bf;
+ wmma::load_matrix_sync(af,ah,16);wmma::load_matrix_sync(bf,bh[warp],16);
+ wmma::mma_sync(acc,af,bf,acc);__syncthreads();
+ }
+ __shared__ float out[8][256];wmma::store_matrix_sync(out[warp],acc,16,wmma::mem_row_major);__syncwarp();
+ for(int z=lane;z<256;z+=32){int m=z/16,n=z%16;
+ if(m0+m=S)return; const float *xs=x+(size_t)s*K;
float v=0; for(int i=threadIdx.x;i=S||nt<1)return;
+ extern __shared__ float sm[];float *qa=sm,*cl=qa+K,*scores=cl+K,*red=scores+T;
+ const float *qs=q+((size_t)s*H+h)*(Q+R);
+ for(int k=tid;k>1;n;n>>=1){if(tid>1;n;n>>=1){if(tid= bytes) return 1;
if (*ptr) cudaFree(*ptr);
@@ -286,6 +381,7 @@ extern "C" int coli_cuda_init(const int *devices, int count) {
if (!select_ctx(ctx)) { g_nctx = 0; return 0; }
cudaDeviceProp prop{};
if (!cuda_ok(cudaGetDeviceProperties(&prop, device), "device properties")) { g_nctx = 0; return 0; }
+ ctx->compute_major=prop.major;ctx->compute_minor=prop.minor;
if(!cuda_ok(cudaStreamCreateWithFlags(&ctx->stream,cudaStreamNonBlocking),"stream creation")){
g_nctx=0;return 0;
}
@@ -307,6 +403,7 @@ extern "C" void coli_cuda_shutdown(void) {
if (ctx->qx) cudaFree(ctx->qx);
if (ctx->qscale) cudaFree(ctx->qscale);
if(ctx->aq)cudaFree(ctx->aq);if(ctx->al)cudaFree(ctx->al);if(ctx->ar)cudaFree(ctx->ar);if(ctx->ac)cudaFree(ctx->ac);
+ for(int b=0;b<24;b++) if(ctx->pipe_buf[b]) cudaFree(ctx->pipe_buf[b]);
if (ctx->host_x) cudaFreeHost(ctx->host_x);
if (ctx->host_y) cudaFreeHost(ctx->host_y);
if (ctx->stream) cudaStreamDestroy(ctx->stream);
@@ -388,6 +485,23 @@ extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
return 1;
}
+extern "C" int coli_cuda_tensor_update(ColiCudaTensor *tensor,
+ const void *weights,
+ const float *scales) {
+ if (!tensor || !weights || (tensor->fmt && !scales)) return 0;
+ DeviceContext *ctx=find_ctx(tensor->device);
+ if (!select_ctx(ctx)) return 0;
+ if (!cuda_ok(cudaMemcpy(tensor->weights,weights,tensor->weight_bytes,
+ cudaMemcpyHostToDevice),"tensor refresh")) return 0;
+ if(tensor->fmt==2){
+ offset_to_signed_s4<<<(unsigned)((tensor->weight_bytes+255)/256),256>>>(
+ (uint8_t*)tensor->weights,tensor->weight_bytes);
+ if(!cuda_ok(cudaGetLastError(),"int4 weight refresh")) return 0;
+ }
+ return !tensor->fmt || cuda_ok(cudaMemcpy(tensor->scales,scales,
+ (size_t)tensor->O*sizeof(float),cudaMemcpyHostToDevice),"scale refresh");
+}
+
extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor,
float *y, const float *x,
const void *weights, const float *scales,
@@ -436,6 +550,34 @@ extern "C" int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
return 1;
}
+extern "C" int coli_cuda_shared_mlp_w4a16(ColiCudaTensor *gate,ColiCudaTensor *up,
+ ColiCudaTensor *down,float *y,const float *x,int S){
+ if(!gate||!up||!down||!x||!y||S<1||gate->fmt!=2||up->fmt!=2||down->fmt!=2||
+ 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)||ctx->compute_major<7)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);
+ if(!reserve(&ctx->x,&ctx->x_cap,xb)||!reserve(&ctx->gate,&ctx->gate_cap,ib)||
+ !reserve(&ctx->up,&ctx->up_cap,ib)||!reserve(&ctx->y,&ctx->y_cap,xb)||
+ !reserve_pinned(&ctx->host_x,&ctx->host_x_cap,xb)||
+ !reserve_pinned(&ctx->host_y,&ctx->host_y_cap,xb))return 0;
+ std::memcpy(ctx->host_x,x,xb);
+ if(!cuda_ok(cudaMemcpyAsync(ctx->x,ctx->host_x,xb,cudaMemcpyHostToDevice,ctx->stream),
+ "shared w4a16 input upload"))return 0;
+ dim3 hidden((unsigned)((I+63)/64),(unsigned)((S+15)/16));
+ dim3 output((unsigned)((D+63)/64),(unsigned)((S+15)/16));
+ w4a16_gate_up<<stream>>>(ctx->gate,ctx->up,ctx->x,
+ (const uint8_t*)gate->weights,(const uint8_t*)up->weights,gate->scales,up->scales,S,D,I);
+ silu_mul<<<(unsigned)(((size_t)S*I+255)/256),256,0,ctx->stream>>>(ctx->gate,ctx->up,(size_t)S*I);
+ w4a16_matmul<<