The new g64 model quantizes experts with group-size 64 (fmt=4): one f32
scale per 64 elements instead of per-row. This required changes across
the entire compute stack — CUDA kernel, engine dispatch, and dequant
helpers — plus a new fused gate+up kernel to recover the ~40% throughput
that the generic grouped path costs.
CUDA backend (backend_cuda.cu):
- ColiCudaTensor: add gs/ng fields for group-size metadata
- row_bytes/weight_at: fmt=4 uses same packed int4 layout as fmt=2
- quant_matmul: apply per-group scales (scales[o*ng+g]) for fmt=4,
matching the CPU matmul_i4_grouped accumulation exactly
- coli_cuda_tensor_upload: allocate O*ng scales (not O) for fmt=4,
store gs/ng on the tensor
- coli_cuda_tensor_update: handle grouped scales on refresh
- coli_cuda_tensor_free/tensor_bytes: VRAM accounting uses O*ng
- coli_cuda_expert_group: return 0 for fmt=4 (fall back to correct
per-expert path, since grouped_hidden/grouped_down kernels only
handle per-row scales)
- All 11 quant_matmul call sites updated to pass gs,ng
Engine (glm.c):
- qt_cuda_upload: pass t->gs to tensor_upload
- matmul dispatch (line 816): pass w->gs to coli_cuda_matmul
- kv_b shard upload: pass l->kv_b.gs
- kv_b shard rb computation: add fmt=4 case ((I+1)/2)
- embed_row: add fmt=4 branch with per-group scale dequant
- qt_addrow/qt_matvec_rows: add fmt=4 branches (CPU MLA absorption path)
- expert_gate_up: add matmul_i4_grouped_pair — fused gate+up for fmt=4
that reads x once instead of twice (+40% tok/s, 0.77 -> 1.08)
- run_text: prepend [gMASK]<sop> prefix for GLM models (#108 fix —
without it, PROMPT mode generates garbage on all GLM snapshots)
API (backend_cuda.h, backend_loader.c):
- coli_cuda_tensor_upload and coli_cuda_matmul signatures: add gs param
- Loader typedefs and wrappers thread gs through the DLL boundary
Tests:
- bench_tensor_core.cu, test_backend_cuda.cu, test_pipe_cuda.cu:
update all calls to new API signature (append gs=0 for non-grouped)
New tool: c/tools/diag_harness.py
- Comprehensive model diagnostic harness (system probe, correctness
smoke, deep PROFILE diagnostic, quality benchmarks via eval_glm.py,
throughput with/without MTP). Outputs JSON + Markdown reports.
Performance (GLM-5.2 744B g64 / RTX 5070 Ti / 32GB RAM):
broken CUDA: 0.05 tok/s (every tensor fell back to CPU)
fixed CUDA: 0.30 tok/s (6x — CUDA working)
+ full opt stack: 0.77 tok/s (CACHE_ROUTE + EXPERT_BUDGET)
+ fused grouped pair: 1.08 tok/s (+40% from gate+up fusion)
The grouped MoE kernels were per-row-only: GroupDesc had no group-size
fields, row_bytes() returned 0 for fmt=4, the scale buffer was hardcoded
to O floats, and a fmt=4 group that reached the generic path would have
been silently decoded as int2. This closed the GPU expert tier to every
grouped container — including the g64 quality line (#225) and the E8
lattice route (#347) whose whole point is fitting more experts in VRAM.
- ColiCudaTensor gains gs/scale_count; upload allocates O*ceil(I/gs)
scales for fmt=4 and applies the same offset->signed nibble conversion
as fmt=2 (identical packing). New ABI entry coli_cuda_tensor_upload_g
carries gs without touching the existing symbol — an old Windows DLL
missing it returns 0 and the tensor simply stays CPU-side.
- GroupDesc gains per-tensor group sizes; new grouped_hidden_g4_dual /
grouped_down_g4 apply the per-group scale inside the accumulation
(gs is required even, so a packed byte never straddles groups; gs=0
degrades to per-row, letting fmt=2 members ride the same launch).
- coli_cuda_expert_group routes any group containing fmt=4 through the
g4 kernels; pure-fmt=2 groups keep the existing paths byte-identical.
The generic fallback now explicitly rejects fmt=4 instead of decoding
garbage (#334's prevention note, made real).
tests/test_grouped_g4_cuda.cu: kernel-vs-CPU oracle over 50 trials x 3
experts — gs=64, a non-divisible tail group (200 % 64), and a per-row
member in the same launch: zero mismatches on a 5090.
make check 77/77; CPU, CUDA and MinGW builds clean.
First increment of the #431 plan (device router -> indirect kernels ->
one-graph decode). At decode (S=1) on the pipe2 path, the router runs on
the layer's home device: a tiny E x D logits GEMV + sigmoid, then a
single-thread selection kernel that clones moe()'s plain routing path
verbatim — bias-augmented top-K by choice with strict-> tie-breaking,
weights from the raw logit, route-level TOPP truncation, norm_topk,
routed_scale. Results pack into one scratch buffer and come back in a
single ~68-byte D2H; moe() consumes them through the same pre-routed
shortcut the Metal layer-CB uses (g_pre_idx, #417 bookkeeping included),
so usage/heat/recency accounting is identical to the CPU router.
Structural value: routing becomes available ON the device timeline,
which is what PR-B (indirect expert kernels, static topology) and PR-C
(whole-decode CUDA Graph) build on.
Opt-in, default off. Gated to the plain routing path — CACHE_ROUTE,
ROUTE_P and ROUTE_TRACE keep the CPU ranking they need; any upload or
launch failure falls back to the CPU router silently. Router weights
(E x D f32, ~6.3 MB/layer) upload lazily to the layer's home device.
tests/test_router_cuda.cu: kernel-vs-CPU-reference oracle over 200
random trials (mixed TOPP/norm_topk/scale): 200/200 exact selections,
zero near-tie flips, zero weight mismatches on a 5090.
PR #111 added 24 functions (pipe_*, attention_*_batch*, shared_mlp_w4a16,
tensor_update) to backend_cuda.h without COLI_CUDA_DLLEXPORT and without
backend_loader.c wrappers: on Windows the host failed to link (undefined
references from glm.c) and the DLL exported none of the new entry points.
Decorate the declarations and add matching typedefs/RESOLVEs/wrappers
following the loader's existing pattern. DLL now exports 39 symbols;
engine links; TF oracle 32/32 CPU + dual-GPU (sm_120/sm_89).
Co-authored-by: olorin <io@zyphyr.co>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
* Fuse CUDA expert MLP execution
* Group CUDA expert transfers by device
* Instrument grouped CUDA expert execution
* Bound grouped CUDA decode scratch
* Execute expert groups across GPUs in parallel
* Release host backing for multi-GPU experts
* Define quality-preserving memory policies
* Overlap cold expert loading with resident compute
* Adapt expert placement with session LFRU
* Fuse q4 expert gate and up dispatch
* Plan CPU work on physical cores
* Batch grouped expert CUDA kernels
* Separate VRAM and RAM expert placement
* Add ragged multi-sequence decode forward
* feat(runtime): add continuous decode scheduler
* Route concurrent API requests through batch scheduler
* Harden multiplex request lifecycle and framing
* Cancel disconnected multiplex requests
* Bind API port before starting the engine
* fix automatic KV slot allocation
* add native int4 Tensor Core grouped GEMM
* add Tensor Core throughput benchmark
* optimize packed int4 low-row kernels
* add asynchronous CUDA staging streams
* document validated six-GPU dense acceleration
* tune six-GPU expert hot set
* raise validated expert hot-set target
* add CUDA MLA absorption core
* fuse grouped expert gate and up projections
* Warn for explicit lossy routing flags
* Add full-resident expert placement mode
* Adapt VRAM expert slots to live routes
* Accelerate int4 matvec on AVX-512
* Reduce AVX-512 and RoPE decode overhead
* Seed every GPU expert layer after prefill
* Limit live GPU swaps during decode
* CUDA batch MLA attention, kv_b head-sharding, fused o_proj, expert-group dispatch, W4A16 kernels
Lab-qualified on the 6x RTX 5090 machine (914-token request benchmark):
- batch MLA absorption kernel (COLI_CUDA_ATTN=1): whole-batch attention on
device, 154.8s -> 102.4s
- attention -> o_proj fusion on the layer device: -> 97.4s
- kv_b head-sharding across cards (COLI_CUDA_ATTN_SHARD=1), no weight
duplication: -> 94.05s
- per-device expert-group dispatch with pinned-buffer async transfers,
W4A16 tensor-core kernels for the shared expert, OMP hot-thread tuning
Negative results (reverted, kept out): GPU-side weighted scatter-add
(atomics + per-layer D2H lose 43.8%), shared-expert fused small-batch
kernel (-38.8%), W4A4 grouped tensor cores (int4 activations corrupt
output). Details in the lab research log.
* GPU resident pipeline: device-resident prefill attention chain, GPU expert groups in prefill, batched router, W4A16 mixed dispatch
COLI_CUDA_PIPE=1 keeps the prefill data plane on the layer home device;
control flow (routing, cache/pin management) stays on CPU. Any CUDA
failure falls back to the unchanged CPU path.
- Device primitives + unit tests (tests/test_pipe_cuda.cu): rmsnorm
(strided), interleaved RoPE, silu-mul, residual add, fixed-order row
merge (no atomics), device-input GEMM, persistent per-device scratch.
All verified against the engine's CPU math on SM120 (worst 1.2e-5).
- attn_pipe_prefill: q_a -> norm -> q_b -> rope -> kv_a -> norm -> rope ->
batch attention -> o_proj in one device chain (q_a/q_b/kv_a colocated
with kv_b); only the final [S,D] and the new KV rows return to host.
Attention 41.2s -> 30.8s on the 1571-token benchmark.
- Prefill batch-union now uses the GPU expert groups (previously gated to
S<=64, leaving all VRAM-resident experts idle during prefill - measured
21ms of GPU expert time in a 148s prefill). Expert phase 78.9s -> 69.0s.
- Router computed as one batched matmul instead of S sequential rows
(bit-identical math).
- W4A16 tensor-core path for expert groups (COLI_CUDA_TC_W4A16=1) with
row-count mixed dispatch: >=16 rows per expert use tensor cores, smaller
batches keep the naive kernel (tensor cores measured negative below
~16 rows). Expert phase 69.0s -> 64.3s, decode unaffected.
Net on the 1571-token prefill benchmark: 148.8s -> 114.3-126.8s
(component timings stable across runs; wall drifts +-3-5s because
.coli_usage placement learning shifts the expert tiers between runs).
PROFILO now also prints the prefill-phase breakdown.
* Skip OMP hot-thread tuning when CUDA is enabled
The active-spin worker team measured 66.9s->20.9s on the CPU-only Zen5
build, but on the six-GPU full-residency workload the spinning workers
contend with the CUDA dispatch threads: ~4x slower prefill with the
process stuck near 1.8 cores. Gate the tuning on COLI_CUDA so each
configuration keeps the behavior it was measured to prefer.
* Inc.2a: sparse layers fully resident on the layer device, residual hops cards at layer boundaries
COLI_CUDA_PIPE=2 keeps the residual stream on the layer home device for
consecutive sparse layers (cudaMemcpyPeer at boundaries): in/post norms,
attention chain, both residual adds and the shared-expert MLP run on
device. Per layer only the post-norm activations (router + CPU-tier
experts + group gather), the new KV rows and, on DSA indexer layers, the
pre-attention norm leave the card. Per-layer transfers drop from ~130MB
to ~70MB. A device-side snapshot at layer entry makes any mid-layer CUDA
failure fall back to the unchanged CPU path idempotently.
1571-token prefill: 127.1s (PIPE=1 control) -> 117.6/118.9s, components
attention 30.8->26.1, other 31.8->22.5-24.5; output verified coherent
against the control.
* Head-sharded attention inside the pipe: negative on PCIe star topology, gated opt-in
Slicing q per card from the home device and collecting ctx back
serializes ~95MB/layer through the home card's PCIe link: attention
26.1s -> 41.4/44.4s on the 1571-token benchmark (two repeats), wall
117.6 -> 135-138s. The standalone host-path sharding won because six
cards uploaded from host RAM in parallel; a home-device star has no
such parallelism without NVLink. Kept behind COLI_CUDA_PIPE_SHARD=1
for interconnects where peer bandwidth does not share one root port.
* Inc.3: device-resident KV shadow for decode attention
Decode re-uploaded the whole latent+rope window per layer per token
(~300MB/token at 1571 context). Each layer now keeps a device shadow of
the compressed KV on its kv_b card, bulk-synced when behind and appended
incrementally; the host cache stays canonical. Invalidation on kv_bind
(slot switch), kv_alloc (resize) and on any overwrite of mirrored rows,
with the legacy full-upload path as fallback.
Measured (COLI_CUDA_PIPE gate): short-context decode 5.48 -> 5.59/5.87
tok/s, 1571-context decode 4.14 -> 4.22 tok/s. Decode remains CPU-expert
bound; the shadow removes the transfer tax, not the compute.
* tools: unified user-experience benchmark (bench_ux.sh)
Two fixed scenarios (short chat, long-document QA), TTFT + decode tok/s
+ first-line drift check, TEMP=0 DRAFT=0 enforced, medians over REPS
runs. Encodes the measurement discipline from the lab record: same
binary per comparison, judge medians because .coli_usage placement
learning drifts wall times between runs.
* tools: bench_ux.sh executable bit
* gitignore compiled test binaries
* tools: expert_atlas.py — measure per-expert topic affinity (#175)
Diffs .coli_usage across 10 themed probe batches (code/math/chinese/
prose/science/law/poetry/structured/translation/casual, 3 prompts each)
driven through a running API server — one engine load total. Every
touched expert gets a topic-affinity vector, entropy, and a specialist/
generalist label; output experts.json feeds the Brain page hover.
* serve: persist .coli_usage after every turn in mux mode, not only at exit
run_serve_mux saved the learning cache once at shutdown; a crash lost
the whole session's routing history, and live consumers of the file
(expert_atlas.py diffs it between probe batches) saw a frozen snapshot.
Now saved per turn like the interactive path (165KB write, negligible).
* web: Brain hover shows measured expert atlas when published
If /experts.json (from tools/expert_atlas.py, #175) is served next to
the app, the tooltip upgrades from the depth heuristic to measured
data: specialist/generalist label, entropy, and the top-3 topic
affinities. Row index maps to real layer (row+3, last row = MTP 78).
Falls back to the heuristic when no atlas is published.
---------
Co-authored-by: JustVugg <JustVugg@users.noreply.github.com>
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.
woolcoxm