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.
coli_cuda_load did LoadLibraryA("coli_cuda.dll") with a bare name. Windows'
default search order includes the current working directory (and, without
SafeDllSearchMode, other writable locations), so an attacker who plants a
coli_cuda.dll where the user launches glm.exe — or inside a downloaded model
directory the user cd's into — gets their DllMain executed at load time:
DLL hijacking -> arbitrary code execution.
Now the loader resolves the path next to glm.exe via GetModuleFileNameA and
loads that exact file with LOAD_WITH_ALTERED_SEARCH_PATH, so both the DLL and
its dependency search are anchored to the trusted install directory. Fallback
(if GetModuleFileNameA ever fails) uses LOAD_LIBRARY_SEARCH_APPLICATION_DIR |
LOAD_LIBRARY_SEARCH_SYSTEM32 — which also excludes the CWD. Cross-compiles clean
under mingw-w64; the CPU path is unaffected (this file is _WIN32-only).
Note: grammar.h gr__rule memcpy was reviewed in the same pass and is safe (len
is clamped to 63 into a name[64] buffer).
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
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>
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.
Vincenzo Fornaro