Tiered CUDA acceleration for routed experts (opt-in, CPU default untouched) + REPLAY fixture harness (#16)

* feat: add experimental CUDA backend for resident tensors

* feat: promote pinned experts to a bounded VRAM tier

* feat: preload the GPU expert tier at startup

* fix: harden CUDA backend failure handling

* feat: add deterministic multi-GPU tensor placement

* test: add deterministic CUDA benchmark fixture

* perf: make routed experts the default CUDA path

.gitignore

@@ -9,12 +9,15 @@ c/glm
 c/olmoe
 c/iobench
 c/tok_test
+c/backend_cuda.o
+c/backend_cuda_test
 
 # 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

README.md

@@ -89,6 +89,70 @@ COLI_MODEL=/nvme/glm52_i4 ./coli chat
 
 The engine at runtime is pure C — python is only used by the one-time converter.
 
+### Experimental resident CUDA backend
+
+This fork includes an opt-in CUDA backend for model-resident tensors. Streaming
+experts deliberately remain on the original CPU path for now: copying an expert
+from NVMe to the GPU on every use would only replace the disk bottleneck with a
+PCIe bottleneck. Resident quantized tensors are uploaded lazily once and reused.
+
+```bash
+cd c
+make cuda-test CUDA=1                  # q8/q4/q2/f32 kernel correctness
+make CUDA=1
+# optional dense-path experiment (hot experts are configured below)
+COLI_CUDA=1 COLI_GPU=0 CUDA_DENSE=1 SNAP=/nvme/glm52_i4 ./glm 64 4 4
+```
+
+Requirements: Linux, an NVIDIA driver, and a CUDA Toolkit under
+`/usr/local/cuda` (override with `CUDA_HOME=/path/to/cuda`). `CUDA_ARCH=native`
+builds for the GPU in the current machine; set an explicit architecture when
+cross-compiling. Requesting CUDA with a CPU-only binary, an invalid device, or
+an unavailable runtime fails at startup instead of silently falling back.
+
+The normal `make` build and runtime behavior are unchanged. CUDA defaults to an
+expert-only accelerator: resident dense/attention tensors stay on CPU because
+fixture measurements show that moving them does not help while expert I/O is
+the bottleneck. `CUDA_DENSE=1` keeps the earlier all-resident experimental path.
+A measured `PIN` profile can promote its hottest experts into the persistent
+VRAM tier while keeping the rest in RAM:
+
+```bash
+STATS=stats.txt SNAP=/nvme/glm52_i4 ./glm 64 4 4   # collect routing frequencies first
+COLI_CUDA=1 COLI_GPU=0 CUDA_EXPERT_GB=16 \
+PIN=stats.txt PIN_GB=160 SNAP=/nvme/glm52_i4 ./glm 64 4 4
+# multi-GPU expert tier, 96 GB total budget across six devices
+COLI_CUDA=1 COLI_GPUS=0,1,2,3,4,5 CUDA_EXPERT_GB=96 \
+PIN=stats.txt PIN_GB=160 SNAP=/nvme/glm52_i4 ./glm 64 4 4
+```
+
+Selected experts are uploaded during startup, so capacity failures occur before
+inference and the log reports their exact tensor footprint. The budget is clamped
+against free VRAM after reserving the projected dense resident set and 2 GB of
+runtime headroom per selected device. With `COLI_GPUS`, `CUDA_EXPERT_GB` is a
+total budget across the device set; experts are assigned whole to the
+least-loaded device that can hold them. A NUMA-local RAM backing store is not
+implemented yet.
+
+Current limitations: devices use independent contexts and synchronous
+host-staged activation copies—there is no P2P/NCCL dependency yet. The kernels
+are correctness-first custom kernels rather than cuBLAS/Tensor Core kernels.
+This draft intentionally makes no end-to-end speedup claim before the full model
+is benchmarked.
+
+For a reproducible backend A/B without the full checkpoint, generate the
+deterministic 313M-parameter `glm_moe_dsa` fixture and run fixed-token replay:
+
+```bash
+cd c
+python make_glm_bench_model.py --output /nvme/colibri-bench-medium --device cuda
+python benchmark_cuda_fixture.py --model /nvme/colibri-bench-medium --gpu 0
+```
+
+The fixture has random weights and is not a language model. It exists only to
+preserve the real MLA/MoE/streaming shapes and compare CPU streaming, dense-only
+CUDA, CPU hot-store, and CUDA hot-expert execution with identical replay tokens.
+
 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 (`:piu` in chat continues a truncated one), `AUTOPIN=0` disable the learning cache's auto-pin, `THINK=1` enable GLM-5.2's reasoning block, `DRAFT=n` MTP draft depth, `TF=1` teacher-forcing validation.
 
 **The learning cache**: the engine records which experts your usage actually routes to (`.coli_usage` next to the model, updated every turn) and at startup automatically pins the hottest ones in spare RAM. colibrì literally gets faster the more you use it.

c/Makefile

@@ -28,10 +28,36 @@ CFLAGS  = -O3 -march=$(ARCH) -fopenmp -Wall -Wextra -Wno-unused-parameter -Wno-m
 LDFLAGS = -lm -fopenmp
 endif
 
+# CUDA=1 adds an opt-in backend for resident tensors. The default build remains
+# pure C and keeps the original zero-dependency runtime.
+CUDA      ?= 0
+CUDA_HOME ?= /usr/local/cuda
+NVCC      ?= $(CUDA_HOME)/bin/nvcc
+CUDA_ARCH ?= native
+NVCCFLAGS ?= -O3 -std=c++17 -arch=$(CUDA_ARCH) -Xcompiler=-Wall,-Wextra
+CUDA_OBJ  =
+ifeq ($(CUDA),1)
+ifeq ($(UNAME_S),Darwin)
+$(error CUDA=1 is supported only on Linux)
+endif
+CFLAGS   += -DCOLI_CUDA
+LDFLAGS  += -L$(CUDA_HOME)/lib64 -Wl,-rpath,$(CUDA_HOME)/lib64 -lcudart -lstdc++
+CUDA_OBJ  = backend_cuda.o
+endif
+
 all: glm
 
-glm: glm.c st.h json.h tok.h tok_unicode.h compat.h
+glm: glm.c st.h json.h tok.h tok_unicode.h compat.h $(CUDA_OBJ)
-	$(CC) $(CFLAGS) glm.c -o glm $(LDFLAGS)
+	$(CC) $(CFLAGS) glm.c $(CUDA_OBJ) -o glm $(LDFLAGS)
+
+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 $@
+
+cuda-test: backend_cuda.cu backend_cuda.h backend_cuda_test.cu
+	@command -v $(NVCC) >/dev/null 2>&1 || { echo "nvcc not found: set CUDA_HOME or NVCC" >&2; exit 1; }
+	$(NVCC) $(NVCCFLAGS) backend_cuda.cu backend_cuda_test.cu -o backend_cuda_test
+	./backend_cuda_test
 
 olmoe: olmoe.c st.h json.h
 	$(CC) $(CFLAGS) olmoe.c -o olmoe $(LDFLAGS)
@@ -41,4 +67,4 @@ portable:
 	$(MAKE) glm ARCH=x86-64-v3
 
 clean:
-	rm -f olmoe glm
+	rm -f olmoe glm backend_cuda.o backend_cuda_test

c/backend_cuda.cu

@@ -0,0 +1,230 @@
+#include "backend_cuda.h"
+
+#include <cuda_runtime.h>
+
+#include <cstdio>
+#include <cstdlib>
+
+struct ColiCudaTensor {
+    void *weights;
+    float *scales;
+    size_t weight_bytes;
+    int fmt, I, O, device;
+    int tracked;
+};
+
+typedef struct {
+    int device;
+    float *x, *y;
+    size_t x_cap, y_cap;
+    size_t tensor_count, tensor_bytes;
+} DeviceContext;
+
+static DeviceContext g_ctx[COLI_CUDA_MAX_DEVICES];
+static int g_nctx;
+
+static int cuda_ok(cudaError_t err, const char *what) {
+    if (err == cudaSuccess) return 1;
+    std::fprintf(stderr, "[CUDA] %s: %s\n", what, cudaGetErrorString(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(cudaSetDevice(ctx->device), "select 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];
+        return static_cast<float>(((i & 1) ? (v >> 4) : (v & 15)) - 8);
+    }
+    uint8_t v = q[i >> 2];
+    return static_cast<float>(((v >> ((i & 3) * 2)) & 3) - 2);
+}
+
+__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);
+}
+
+static int reserve(float **ptr, size_t *cap, size_t bytes) {
+    if (*cap >= bytes) return 1;
+    if (*ptr) cudaFree(*ptr);
+    *ptr = nullptr;
+    *cap = 0;
+    if (!cuda_ok(cudaMalloc(ptr, bytes), "scratch 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(cudaGetDeviceCount(&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; }
+        cudaDeviceProp prop{};
+        if (!cuda_ok(cudaGetDeviceProperties(&prop, device), "device properties")) { 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) cudaFree(ctx->x);
+        if (ctx->y) cudaFree(ctx->y);
+        ctx->x = ctx->y = nullptr;
+        ctx->x_cap = ctx->y_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(cudaMemGetInfo(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" 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(cudaMalloc(&t->weights, t->weight_bytes), "tensor allocation") ||
+        !cuda_ok(cudaMemcpy(t->weights, weights, t->weight_bytes, cudaMemcpyHostToDevice), "tensor upload")) {
+        coli_cuda_tensor_free(t);
+        return 0;
+    }
+    if (fmt) {
+        if (!cuda_ok(cudaMalloc(&t->scales, (size_t)O * sizeof(float)), "scale allocation") ||
+            !cuda_ok(cudaMemcpy(t->scales, scales, (size_t)O * sizeof(float), cudaMemcpyHostToDevice), "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(cudaMemcpy(ctx->x, x, xb, cudaMemcpyHostToDevice), "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(cudaGetLastError(), "matmul launch") ||
+        !cuda_ok(cudaMemcpy(y, ctx->y, yb, cudaMemcpyDeviceToHost), "output download")) 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) cudaFree(tensor->weights);
+    if (tensor->scales) cudaFree(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;
+}

c/backend_cuda.h

@@ -0,0 +1,49 @@
+#ifndef COLIBRI_BACKEND_CUDA_H
+#define COLIBRI_BACKEND_CUDA_H
+
+#include <stddef.h>
+#include <stdint.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#define COLI_CUDA_MAX_DEVICES 16
+
+/* Opaque, persistent device copy of one resident quantized tensor. */
+typedef struct ColiCudaTensor ColiCudaTensor;
+
+/* Devices are CUDA ordinals, not positions in the input list. */
+int coli_cuda_init(const int *devices, int count);
+void coli_cuda_shutdown(void);
+int coli_cuda_device_count(void);
+int coli_cuda_device_at(int index);
+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);
+
+/* Upload without executing, so capacity failures happen during model startup. */
+int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
+                            const void *weights, const float *scales,
+                            int fmt, int I, int O, int device);
+
+/*
+ * y[S,O] = x[S,I] @ W[O,I]^T.
+ * fmt matches QT in glm.c: 0=f32, 1=int8, 2=int4, 3=int2.
+ * 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,
+                     float *y, const float *x,
+                     const void *weights, const float *scales,
+                     int fmt, int S, int I, int O, int device);
+
+void coli_cuda_tensor_free(ColiCudaTensor *tensor);
+size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor);
+int coli_cuda_tensor_device(const ColiCudaTensor *tensor);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif

c/backend_cuda_test.cu

@@ -0,0 +1,81 @@
+#include "backend_cuda.h"
+
+#include <cmath>
+#include <cstdio>
+#include <cstdint>
+#include <cstdlib>
+
+static int close_enough(const float *got, const float *want, int n) {
+    for (int i = 0; i < n; i++) {
+        if (std::fabs(got[i] - want[i]) > 1e-4f) {
+            std::fprintf(stderr, "mismatch %d: got %.6f want %.6f\n", i, got[i], want[i]);
+            return 0;
+        }
+    }
+    return 1;
+}
+
+int main(int argc, char **argv) {
+    int devices[COLI_CUDA_MAX_DEVICES], ndev = argc > 1 ? argc - 1 : 1;
+    if (ndev > COLI_CUDA_MAX_DEVICES) return 2;
+    for (int i = 0; i < ndev; i++) devices[i] = argc > 1 ? std::atoi(argv[i + 1]) : 0;
+    if (!coli_cuda_init(devices, ndev)) return 77;
+    if (coli_cuda_device_count() != ndev) return 1;
+    int d0 = devices[0], d1 = devices[ndev > 1 ? 1 : 0];
+    size_t count = 99, bytes = 99;
+    coli_cuda_stats(-1, &count, &bytes);
+    if (count || bytes) return 1;
+    const float x[8] = {1, -2, 3, -4, 2, 1, -1, 0.5f};
+    float got[4];
+
+    const int8_t q8[8] = {1, 2, 3, 4, -1, 2, -3, 4};
+    const float s8[2] = {0.5f, 2.0f};
+    const float want8[4] = {-5.0f, -60.0f, 1.5f, 10.0f};
+    ColiCudaTensor *t8 = nullptr;
+    if (!coli_cuda_tensor_upload(&t8, q8, s8, 1, 4, 2, d0)) return 1;
+    if (coli_cuda_tensor_upload(&t8, q8, s8, 1, 5, 2, d0)) return 1;
+    if (ndev > 1 && coli_cuda_tensor_upload(&t8, q8, s8, 1, 4, 2, d1)) return 1;
+    if (!coli_cuda_matmul(&t8, got, x, q8, s8, 1, 2, 4, 2, d0) || !close_enough(got, want8, 4)) return 1;
+
+    /* Rows [-8,-1,0,7] and [1,2,3,4], packed low nibble first. */
+    const uint8_t q4[4] = {0x70, 0xf8, 0xa9, 0xcb};
+    const float s4[2] = {1.0f, 0.25f};
+    const float want4[2] = {-34.0f, -2.5f};
+    ColiCudaTensor *t4 = nullptr;
+    if (!coli_cuda_matmul(&t4, got, x, q4, s4, 2, 1, 4, 2, d1) || !close_enough(got, want4, 2)) return 1;
+
+    const uint8_t q2[2] = {0xe4, 0x1b};
+    const float s2[2] = {0.5f, 2.0f};
+    const float want2[2] = {-2.0f, 12.0f};
+    ColiCudaTensor *t2 = nullptr;
+    if (!coli_cuda_matmul(&t2, got, x, q2, s2, 3, 1, 4, 2, d1) || !close_enough(got, want2, 2)) return 1;
+
+    const float wf[8] = {1, 0, -1, 2, 0.5f, 0.5f, 0.5f, 0.5f};
+    const float wantf[2] = {-10.0f, -1.0f};
+    ColiCudaTensor *tf = nullptr;
+    if (!coli_cuda_matmul(&tf, got, x, wf, nullptr, 0, 1, 4, 2, d0) || !close_enough(got, wantf, 2)) return 1;
+
+    coli_cuda_stats(-1, &count, &bytes);
+    if (count != 4 || bytes != 70) {
+        std::fprintf(stderr, "unexpected CUDA stats: %zu tensors, %zu bytes\n", count, bytes);
+        return 1;
+    }
+    if (coli_cuda_tensor_device(t8) != d0 || coli_cuda_tensor_device(tf) != d0 ||
+        coli_cuda_tensor_device(t4) != d1 || coli_cuda_tensor_device(t2) != d1) return 1;
+    coli_cuda_stats(d0, &count, &bytes);
+    if (ndev > 1) {
+        if (count != 2 || bytes != 48) return 1;
+        coli_cuda_stats(d1, &count, &bytes);
+        if (count != 2 || bytes != 22) return 1;
+    } else if (count != 4 || bytes != 70) return 1;
+
+    coli_cuda_tensor_free(t8);
+    coli_cuda_tensor_free(t4);
+    coli_cuda_tensor_free(t2);
+    coli_cuda_tensor_free(tf);
+    coli_cuda_stats(-1, &count, &bytes);
+    if (count || bytes) return 1;
+    coli_cuda_shutdown();
+    std::printf("cuda backend: q8/q4/q2/f32 correctness ok on %d device(s)\n", ndev);
+    return 0;
+}

c/benchmark_cuda_fixture.py

@@ -0,0 +1,102 @@
+"""Reproducible CPU/CUDA A/B benchmark for make_glm_bench_model.py output."""
+
+import argparse
+import json
+import os
+import re
+import statistics
+import subprocess
+from pathlib import Path
+
+
+SPEED_RE = re.compile(r"REPLAY decode:.*\| ([0-9.]+) tok/s")
+PROFILE_RE = re.compile(
+    r"PROFILO: expert-disk ([0-9.]+)s \| expert-matmul ([0-9.]+)s "
+    r"\| attention ([0-9.]+)s .* lm_head ([0-9.]+)s \| altro ([0-9.-]+)s"
+)
+PROFILE_KEYS = ("disk", "expert_matmul", "attention", "lm_head", "other")
+
+
+def execute(engine: str, env: dict[str, str]) -> tuple[float, list[float]]:
+    run = subprocess.run(
+        [engine, "4", "4", "4"], env=env, text=True, capture_output=True, check=True
+    )
+    speed = SPEED_RE.search(run.stdout)
+    profile = PROFILE_RE.search(run.stdout)
+    if not speed or not profile:
+        raise RuntimeError(f"benchmark output missing\nstdout:\n{run.stdout}\nstderr:\n{run.stderr}")
+    return float(speed.group(1)), [float(value) for value in profile.groups()]
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--model", required=True)
+    parser.add_argument("--engine", default="./glm")
+    parser.add_argument("--gpu", default="0")
+    parser.add_argument("--runs", type=int, default=7)
+    parser.add_argument("--threads", type=int, default=os.cpu_count() or 1)
+    parser.add_argument("--pin-gb", default="1")
+    parser.add_argument("--cuda-expert-gb", default="2")
+    args = parser.parse_args()
+
+    model = Path(args.model).resolve()
+    stats = model / "bench_stats.txt"
+    base = os.environ.copy()
+    for key in (
+        "COLI_CUDA", "COLI_GPU", "COLI_GPUS", "CUDA_EXPERT_GB",
+        "PIN", "PIN_GB", "STATS", "TF", "REPLAY", "CUDA_DENSE",
+    ):
+        base.pop(key, None)
+    base.update(
+        SNAP=str(model),
+        REF=str(model / "ref_glm.json"),
+        REPLAY="1",
+        OMP_NUM_THREADS=str(args.threads),
+        OMP_PROC_BIND="spread",
+        OMP_PLACES="cores",
+    )
+
+    execute(args.engine, base | {"STATS": str(stats)})
+    modes = {
+        "cpu_stream": {},
+        "dense_cuda": {"COLI_CUDA": "1", "COLI_GPU": args.gpu, "CUDA_DENSE": "1"},
+        "cpu_pin": {"PIN": str(stats), "PIN_GB": args.pin_gb},
+        "cuda_pin": {
+            "COLI_CUDA": "1", "COLI_GPU": args.gpu,
+            "PIN": str(stats), "PIN_GB": args.pin_gb,
+            "CUDA_EXPERT_GB": args.cuda_expert_gb,
+        },
+        "cuda_pin_dense": {
+            "COLI_CUDA": "1", "COLI_GPU": args.gpu, "CUDA_DENSE": "1",
+            "PIN": str(stats), "PIN_GB": args.pin_gb,
+            "CUDA_EXPERT_GB": args.cuda_expert_gb,
+        },
+    }
+
+    for extra in modes.values():
+        execute(args.engine, base | extra)  # warm-up
+    speeds = {name: [] for name in modes}
+    profiles = {name: [] for name in modes}
+    names = list(modes)
+    for run_index in range(args.runs):
+        order = names[run_index % len(names):] + names[:run_index % len(names)]
+        for name in order:
+            speed, profile = execute(args.engine, base | modes[name])
+            speeds[name].append(speed)
+            profiles[name].append(profile)
+
+    result = {}
+    for name in names:
+        result[name] = {
+            "runs_tok_s": speeds[name],
+            "median_tok_s": statistics.median(speeds[name]),
+            "median_profile_s": {
+                key: statistics.median(row[index] for row in profiles[name])
+                for index, key in enumerate(PROFILE_KEYS)
+            },
+        }
+    print(json.dumps(result, indent=2))
+
+
+if __name__ == "__main__":
+    main()

c/glm.c

@@ -22,12 +22,17 @@
 #include <string.h>
 #include <math.h>
 #include <time.h>
+#include <limits.h>
 #include <sys/resource.h>
 #if defined(__APPLE__) || defined(__linux__)
 #include <sys/mman.h>                             /* mlock: inchioda le pagine in RAM / wire pages into RAM */
 #endif
 #include "st.h"
 #include "tok.h"
+#ifdef COLI_CUDA
+#include <omp.h>
+#include "backend_cuda.h"
+#endif
 #ifdef __AVX2__
 #include <immintrin.h>
 static inline float hsum256(__m256 v){            /* somma orizzontale di 8 float */
@@ -58,7 +63,13 @@ typedef struct {
  *   fmt=2 INT4  -> q4 (2 valori per byte, impacchettati) + scala per riga
  * INT4 e' cio' che fa stare la densa residente nei 15 GB (0.5 byte/param). */
 /* fmt: 0 F32, 1 INT8, 2 INT4 (2/byte), 3 INT2 (4/byte). q4 ospita sia int4 che int2 packed. */
-typedef struct { int fmt; float *qf; int8_t *q8; uint8_t *q4; float *s; int O, I; } QT;
+typedef struct {
+    int fmt; float *qf; int8_t *q8; uint8_t *q4; float *s; int O, I;
+#ifdef COLI_CUDA
+    ColiCudaTensor *cuda;
+#endif
+    int cuda_eligible, cuda_failed, cuda_device;  /* resident tensor, never a reused expert slot */
+} QT;
 static int64_t qt_bytes(const QT *t){    /* byte residenti del tensore */
     int64_t n=(int64_t)t->O*t->I;
     if(t->fmt==0) return n*4;
@@ -113,12 +124,53 @@ typedef struct {
     uint64_t mtp_prop, mtp_acc;                  /* statistica acceptance */
     int **eroute; int *enr;                      /* metodo C: routing dell'ULTIMO token per layer */
     uint64_t eclock, hits, miss, ereq;
+    uint64_t gpu_expert_calls; int gpu_expert_count; int64_t gpu_expert_bytes;
     uint64_t n_fw, n_emit;                       /* metodo E: forward di decode / token emessi */
     double t_edisk, t_emm, t_attn, t_kvb, t_head;/* profiling: dove va il tempo (sempre attivo) */
     int64_t resident_bytes;
 } Model;
 
 static void usage_save(Model *m);        /* cache che impara: definita accanto a stats_dump */
+#ifdef COLI_CUDA
+static int g_cuda_enabled;
+static double g_cuda_expert_gb;
+static int g_cuda_dense;
+static int g_cuda_devices[COLI_CUDA_MAX_DEVICES], g_cuda_ndev, g_cuda_rr;
+static int64_t g_cuda_dense_projected[COLI_CUDA_MAX_DEVICES];
+static void qt_cuda_reset(QT *t){
+    if(t->cuda){ coli_cuda_tensor_free(t->cuda); t->cuda=NULL; }
+    t->cuda_failed=0;
+}
+static int qt_cuda_upload(QT *t){
+    const void *weights = t->fmt==0 ? (const void*)t->qf
+                        : t->fmt==1 ? (const void*)t->q8 : (const void*)t->q4;
+    return coli_cuda_tensor_upload(&t->cuda,weights,t->s,t->fmt,t->I,t->O,t->cuda_device);
+}
+static void cuda_stats_print(void){
+    size_t n=0,b=0; coli_cuda_stats(-1,&n,&b);
+    fprintf(stderr,"[CUDA] resident set: %zu tensor, %.2f GB VRAM\n",n,b/1e9);
+    if(g_cuda_ndev>1) for(int i=0;i<g_cuda_ndev;i++){
+        coli_cuda_stats(g_cuda_devices[i],&n,&b);
+        fprintf(stderr,"[CUDA]   device %d: %zu tensor, %.2f GB\n",g_cuda_devices[i],n,b/1e9);
+    }
+}
+static int parse_cuda_devices(const char *list, int *out){
+    if(!list||!*list) return 0;
+    int n=0; const char *p=list;
+    while(*p){
+        char *end=NULL; long v=strtol(p,&end,10);
+        if(end==p||v<0||v>INT_MAX||n>=COLI_CUDA_MAX_DEVICES) return 0;
+        for(int i=0;i<n;i++) if(out[i]==(int)v) return 0;
+        out[n++]=(int)v; p=end;
+        while(*p==' '||*p=='\t') p++;
+        if(!*p) break;
+        if(*p++!=',') return 0;
+        while(*p==' '||*p=='\t') p++;
+        if(!*p) return 0;
+    }
+    return n;
+}
+#endif
 static double now_s(void){ struct timespec t; clock_gettime(CLOCK_MONOTONIC,&t); return t.tv_sec+t.tv_nsec*1e-9; }
 static double rss_gb(void){ struct rusage r; getrusage(RUSAGE_SELF,&r);
 #ifdef __APPLE__
@@ -385,7 +437,21 @@ static void matmul_i4_idot(float *y, const int8_t *xq, const float *sx, const ui
         for(int s=0;s<S;s++) y[(int64_t)s*O+o]=(float)dot_i4i8(w,xq+(int64_t)s*I,I)*sc*sx[s]; }
 }
 
-static void matmul_qt(float *y, const float *x, const QT *w, int S){
+static void matmul_qt(float *y, const float *x, QT *w, int S){
+#ifdef COLI_CUDA
+    /* The CUDA backend owns persistent copies only for model-resident tensors.
+     * Streaming expert slots are reused for different IDs and must never enter
+     * this cache. Nested OpenMP calls stay on CPU because each device context
+     * intentionally owns one synchronous scratch stream in this stage. */
+    if(g_cuda_enabled && w->cuda_eligible && !w->cuda_failed && !omp_in_parallel()){
+        const void *weights = w->fmt==0 ? (const void*)w->qf
+                            : w->fmt==1 ? (const void*)w->q8 : (const void*)w->q4;
+        if(coli_cuda_matmul(&w->cuda,y,x,weights,w->s,w->fmt,S,w->I,w->O,w->cuda_device)) return;
+        w->cuda_failed=1;
+        fprintf(stderr,"[CUDA] tensor [%d,%d] su device %d disabilitato dopo errore; fallback CPU\n",
+            w->O,w->I,w->cuda_device);
+    }
+#endif
     if(w->fmt==0){ matmul(y,x,w->qf,S,w->I,w->O); return; }
     /* int8 IDOT vince sempre (1.4-2.5x). int4 IDOT: l'autore su AVX2 trovo' che a S=1
      * non ripaga (soglia S>=2); ma su ARM/SDOT il singolo token CONVIENE (vedi g_i4s /
@@ -580,7 +646,15 @@ static void qt_from_disk(Model *m, const char *name, int O, int I, int bits, int
     }
 }
 static QT qt_load(Model *m, const char *name, int O, int I, int bits){
-    QT t; memset(&t,0,sizeof(t)); qt_from_disk(m,name,O,I,bits,0,&t); return t;
+    QT t; memset(&t,0,sizeof(t)); qt_from_disk(m,name,O,I,bits,0,&t);
+#ifdef COLI_CUDA
+    if(g_cuda_enabled&&g_cuda_dense){
+        t.cuda_eligible=1;
+        int slot=g_cuda_rr++%g_cuda_ndev; t.cuda_device=g_cuda_devices[slot];
+        g_cuda_dense_projected[slot]+=qt_bytes(&t);
+    }
+#endif
+    return t;
 }
 static float *ld(Model *m, const char *name){   /* tensore 1D f32 residente (norme/bias) */
     int64_t n=st_numel(&m->S,name); if(n<0){fprintf(stderr,"manca %s\n",name);exit(1);}
@@ -744,6 +818,11 @@ static void embed_row(Model *m, int tok, float *x){
  * viste dentro lo slab (zero copie). Fallback per modelli non quantizzati (oracolo tiny).
  * THREAD-SAFE su slot distinti (pread posizionale, st_find read-only). */
 static void expert_load(Model *m, int layer, int eid, ESlot *s){
+#ifdef COLI_CUDA
+    /* A live REPIN may reuse a GPU-enabled pinned slot for a different expert.
+     * Keep its tier assignment, but invalidate the old device weights. */
+    if(s->eid!=eid){ qt_cuda_reset(&s->g); qt_cuda_reset(&s->u); qt_cuda_reset(&s->d); }
+#endif
     Cfg *c=&m->c; int I=c->moe_inter, D=c->hidden, b=m->ebits;
     char nm[3][288]; const char *suf[3]={"gate_proj","up_proj","down_proj"};
     for(int k=0;k<3;k++) snprintf(nm[k],sizeof(nm[k]),"model.layers.%d.mlp.experts.%d.%s.weight",layer,eid,suf[k]);
@@ -1092,6 +1171,9 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out){
             for(int s=0;s<S;s++) for(int kk=0;kk<keff[s];kk++)
                 if(idxs[(int64_t)s*K+kk]==eid){ rows[nr]=s; rw[nr]=ws[(int64_t)s*K+kk]; nr++; break; }
             if(!nr) continue;
+#ifdef COLI_CUDA
+            if(g_cuda_enabled && e->g.cuda_eligible) m->gpu_expert_calls++;
+#endif
             for(int r=0;r<nr;r++) memcpy(xg+(int64_t)r*D, x+(int64_t)rows[r]*D, D*sizeof(float));
             double t0=now_s();
             matmul_qt(gg, xg, &e->g, nr);
@@ -1479,6 +1561,37 @@ static void generate(Model *m, const int *prompt, int np, int n_new, int *out){
     spec_decode(m,out,np,n_new,-1,logit,emit_store,&es,NULL);
 }
 
+static void profile_print(Model *m, double elapsed){
+    double accounted=m->t_edisk+m->t_emm+m->t_attn+m->t_head;
+    printf("PROFILO: expert-disk %.3fs | expert-matmul %.3fs | attention %.3fs "
+           "(di cui kvb %.3fs) | lm_head %.3fs | altro %.3fs\n",
+        m->t_edisk,m->t_emm,m->t_attn,m->t_kvb,m->t_head,elapsed-accounted);
+}
+
+/* Fixed-token decode benchmark: prefill all but the prompt's last token, then
+ * replay the oracle sequence one token at a time. CPU and CUDA therefore see
+ * identical hidden-state inputs even if their argmax predictions differ. */
+static void run_replay(Model *m, const int *full, int nfull, int np){
+    if(np<2||nfull<=np){ fprintf(stderr,"REPLAY richiede prompt e continuation non vuoti\n"); return; }
+    kv_alloc(m,nfull+2);
+    float *logit=step(m,full,np-1,0); free(logit);
+    m->hits=m->miss=m->ereq=m->gpu_expert_calls=0;
+    m->t_edisk=m->t_emm=m->t_attn=m->t_kvb=m->t_head=0;
+    double t0=now_s(); int steps=0;
+    for(int i=np-1;i<nfull-1;i++){
+        logit=step(m,full+i,1,i); free(logit); steps++;
+    }
+    double dt=now_s()-t0, tot=m->hits+m->miss;
+    printf("REPLAY decode: %d token in %.3fs | %.2f tok/s | expert hit %.1f%%\n",
+        steps,dt,steps/dt,tot?100.0*m->hits/tot:0.0);
+    profile_print(m,dt);
+#ifdef COLI_CUDA
+    if(m->gpu_expert_count) printf("CUDA expert tier: %d residenti (%.2f GB) | %llu chiamate servite da VRAM\n",
+        m->gpu_expert_count,m->gpu_expert_bytes/1e9,(unsigned long long)m->gpu_expert_calls);
+    if(g_cuda_enabled) cuda_stats_print();
+#endif
+}
+
 /* generazione reale: tokenizza PROMPT, prefill + decode greedy con stop su EOS,
  * detokenizza e stampa il testo in streaming. */
 static void run_text(Model *m, const char *snap, const char *prompt, int ngen){
@@ -1509,9 +1622,12 @@ static void run_text(Model *m, const char *snap, const char *prompt, int ngen){
     printf("speculazione: %.2f token/forward (%llu fw per %llu tok) | MTP acceptance %.0f%% (%llu/%llu)\n",
         m->n_fw?(double)m->n_emit/m->n_fw:1.0, (unsigned long long)m->n_fw, (unsigned long long)m->n_emit,
         m->mtp_prop?100.0*m->mtp_acc/m->mtp_prop:0.0, (unsigned long long)m->mtp_acc, (unsigned long long)m->mtp_prop);
-    double acc=m->t_edisk+m->t_emm+m->t_attn+m->t_head;
+#ifdef COLI_CUDA
-    printf("PROFILO: expert-disk %.1fs | expert-matmul %.1fs | attention %.1fs (di cui kvb %.1fs) | lm_head %.1fs | altro %.1fs\n",
+    if(m->gpu_expert_count) printf("CUDA expert tier: %d residenti (%.2f GB) | %llu chiamate servite da VRAM\n",
-        m->t_edisk, m->t_emm, m->t_attn, m->t_kvb, m->t_head, dt-acc);
+        m->gpu_expert_count,m->gpu_expert_bytes/1e9,(unsigned long long)m->gpu_expert_calls);
+    if(g_cuda_enabled) cuda_stats_print();
+#endif
+    profile_print(m,dt);
     free(pids); free(all);
     usage_save(m);
 }
@@ -1775,6 +1891,59 @@ static void pin_load(Model *m, const char *statspath, double gb){
     m->resident_bytes += (int64_t)npin*eb;
     fprintf(stderr,"[PIN] hot-store: %d expert in RAM (%.1f GB) in %.0fs da %s\n",
         npin, npin*eb/1e9, now_s()-t0, statspath);
+#ifdef COLI_CUDA
+    if(g_cuda_enabled && g_cuda_expert_gb>0){
+        double remaining[COLI_CUDA_MAX_DEVICES]={0}, placed_b[COLI_CUDA_MAX_DEVICES]={0};
+        int placed_n[COLI_CUDA_MAX_DEVICES]={0};
+        double budget=g_cuda_expert_gb*1e9, safe_total=0;
+        for(int i=0;i<g_cuda_ndev;i++){
+            size_t free_b=0,total_b=0;
+            if(coli_cuda_mem_info(g_cuda_devices[i],&free_b,&total_b)){
+                /* Dense tensors are assigned round-robin and upload lazily.
+                 * Reserve their projected footprint plus 2 GB per device. */
+                remaining[i]=(double)free_b-(double)g_cuda_dense_projected[i]-2e9;
+                if(remaining[i]<0) remaining[i]=0;
+                safe_total+=remaining[i];
+            }
+        }
+        if(budget>safe_total) budget=safe_total;
+        for(int a=0;a<npin && m->gpu_expert_bytes<budget;a++){
+            int li=r[a].l;
+            for(int z=0;z<m->npin[li];z++) if(m->pin[li][z].eid==r[a].e){
+                ESlot *s=&m->pin[li][z];
+                int64_t need=qt_bytes(&s->g)+qt_bytes(&s->u)+qt_bytes(&s->d);
+                if(m->gpu_expert_bytes+need>budget) break;
+                int tried[COLI_CUDA_MAX_DEVICES]={0}, placed=0;
+                for(int attempt=0;attempt<g_cuda_ndev && !placed;attempt++){
+                    int best=-1;
+                    for(int i=0;i<g_cuda_ndev;i++) if(!tried[i] && remaining[i]>=need &&
+                        (best<0||placed_b[i]<placed_b[best])) best=i;
+                    if(best<0) break;
+                    tried[best]=1;
+                    s->g.cuda_device=s->u.cuda_device=s->d.cuda_device=g_cuda_devices[best];
+                    s->g.cuda_eligible=s->u.cuda_eligible=s->d.cuda_eligible=1;
+                    if(qt_cuda_upload(&s->g) && qt_cuda_upload(&s->u) && qt_cuda_upload(&s->d)){
+                        int64_t actual=(int64_t)coli_cuda_tensor_bytes(s->g.cuda)
+                                      +(int64_t)coli_cuda_tensor_bytes(s->u.cuda)
+                                      +(int64_t)coli_cuda_tensor_bytes(s->d.cuda);
+                        m->gpu_expert_count++; m->gpu_expert_bytes+=actual;
+                        remaining[best]-=actual; placed_b[best]+=actual; placed_n[best]++;
+                        placed=1;
+                    } else {
+                        qt_cuda_reset(&s->g); qt_cuda_reset(&s->u); qt_cuda_reset(&s->d);
+                        s->g.cuda_eligible=s->u.cuda_eligible=s->d.cuda_eligible=0;
+                        remaining[best]=0;             /* device rejected its projected capacity */
+                    }
+                }
+                break;
+            }
+        }
+        fprintf(stderr,"[CUDA] hot expert tier: %d/%d expert, VRAM %.2f GB (budget totale %.1f GB)\n",
+            m->gpu_expert_count,npin,m->gpu_expert_bytes/1e9,g_cuda_expert_gb);
+        for(int i=0;i<g_cuda_ndev;i++) fprintf(stderr,"[CUDA]   device %d: %d expert, %.2f GB\n",
+            g_cuda_devices[i],placed_n[i],placed_b[i]/1e9);
+    }
+#endif
     pin_wire(m);                                   /* inchioda in RAM (no compressione) / wire in RAM (no compression) */
     free(r); free(cnt_l);
 }
@@ -1873,6 +2042,32 @@ int main(int argc, char **argv){
     int cap  = argc>1?atoi(argv[1]):64;
     int ebits= argc>2?atoi(argv[2]):8;
     int dbits= argc>3?atoi(argv[3]):ebits;
+#ifdef COLI_CUDA
+    if(getenv("COLI_CUDA") && atoi(getenv("COLI_CUDA"))){
+        const char *one=getenv("COLI_GPU"), *many=getenv("COLI_GPUS");
+        if(one&&many){ fprintf(stderr,"usa COLI_GPU oppure COLI_GPUS, non entrambi\n"); return 2; }
+        if(many) g_cuda_ndev=parse_cuda_devices(many,g_cuda_devices);
+        else if(one) g_cuda_ndev=parse_cuda_devices(one,g_cuda_devices);
+        else { g_cuda_ndev=1; g_cuda_devices[0]=0; }
+        if(g_cuda_ndev<1){ fprintf(stderr,"COLI_GPUS non valido: usa una lista come 0,1,2\n"); return 2; }
+        g_cuda_enabled=coli_cuda_init(g_cuda_devices,g_cuda_ndev);
+        if(!g_cuda_enabled){ fprintf(stderr,"[CUDA] backend richiesto ma non disponibile\n"); return 2; }
+    }
+    g_cuda_dense=getenv("CUDA_DENSE")?atoi(getenv("CUDA_DENSE")):0;
+    g_cuda_expert_gb=getenv("CUDA_EXPERT_GB")?atof(getenv("CUDA_EXPERT_GB")):0;
+    if((getenv("COLI_GPU")||getenv("COLI_GPUS"))&&!g_cuda_enabled){ fprintf(stderr,"COLI_GPU(S) richiede COLI_CUDA=1\n"); return 2; }
+    if(g_cuda_dense&&!g_cuda_enabled){ fprintf(stderr,"CUDA_DENSE richiede COLI_CUDA=1\n"); return 2; }
+    if(g_cuda_expert_gb>0 && !g_cuda_enabled){ fprintf(stderr,"CUDA_EXPERT_GB richiede COLI_CUDA=1\n"); return 2; }
+    if(g_cuda_enabled) fprintf(stderr,"[CUDA] mode: routed experts%s\n",g_cuda_dense?" + resident dense tensors":" only (resident dense on CPU)");
+#else
+    if((getenv("COLI_CUDA") && atoi(getenv("COLI_CUDA"))) ||
+       getenv("COLI_GPU") || getenv("COLI_GPUS") ||
+       (getenv("CUDA_DENSE") && atoi(getenv("CUDA_DENSE"))) ||
+       (getenv("CUDA_EXPERT_GB") && atof(getenv("CUDA_EXPERT_GB"))>0)){
+        fprintf(stderr,"CUDA richiesto ma questo binario e' CPU-only; ricompila con: make CUDA=1\n");
+        return 2;
+    }
+#endif
     printf("== Motore C GLM (glm_moe_dsa), cache=%d expert/layer | expert@%d-bit densa@%d-bit | idot: " IDOT_KERNEL " ==\n", cap, ebits, dbits);
     g_mem_avail_boot = mem_available_gb();
     Model m; double t0=now_s(); model_init(&m,snap,cap,ebits,dbits);
@@ -1934,11 +2129,23 @@ int main(int argc, char **argv){
     int np,nfull; int *prompt=read_arr(ref,"prompt_ids",&np); int *full=read_arr(ref,"full_ids",&nfull);
     int n_new=nfull-np;
 
+    if(getenv("REPLAY")){
+        run_replay(&m,full,nfull,np);
+        if(stats) stats_dump(&m,stats);
+        return 0;
+    }
+
     if(getenv("TF")){
         int *tf=read_arr(ref,"tf_pred",&(int){0});
-        int *pred=malloc(nfull*sizeof(int)); forward_all(&m, full, nfull, pred);
+        int *pred=malloc(nfull*sizeof(int)); double tt=now_s();
+        forward_all(&m, full, nfull, pred); double tdt=now_s()-tt;
         int ok=0; for(int i=0;i<nfull;i++) ok+=(pred[i]==tf[i]);
-        printf("PREFILL (teacher-forcing) C vs oracolo: %d/%d posizioni\n", ok, nfull);
+        printf("PREFILL (teacher-forcing) C vs oracolo: %d/%d posizioni | %.1f pos/s\n",
+            ok,nfull,nfull/tdt);
+        profile_print(&m,tdt);
+#ifdef COLI_CUDA
+        if(g_cuda_enabled) cuda_stats_print();
+#endif
         return 0;
     }
     int *out=malloc((np+n_new)*sizeof(int));
@@ -1952,6 +2159,12 @@ int main(int argc, char **argv){
         g_draft, m.n_fw?(double)m.n_emit/m.n_fw:1.0, (unsigned long long)m.n_fw, (unsigned long long)m.n_emit);
     printf("Hit-rate cache expert: %.1f%% (hit=%llu miss=%llu) | RSS: %.2f GB | %.1f tok/s\n",
            tot?100.0*m.hits/tot:0.0, (unsigned long long)m.hits, (unsigned long long)m.miss, rss_gb(), n_new/dt);
+    profile_print(&m,dt);
+#ifdef COLI_CUDA
+    if(m.gpu_expert_count) printf("CUDA expert tier: %d residenti (%.2f GB) | %llu chiamate servite da VRAM\n",
+        m.gpu_expert_count,m.gpu_expert_bytes/1e9,(unsigned long long)m.gpu_expert_calls);
+    if(g_cuda_enabled) cuda_stats_print();
+#endif
     if(stats) stats_dump(&m,stats);
     return 0;
 }

c/make_glm_bench_model.py

@@ -0,0 +1,99 @@
+"""Build a deterministic, medium-size GLM-MoE fixture for backend benchmarks.
+
+This is not a useful language model. It preserves the real glm_moe_dsa data
+flow while remaining small enough to generate locally and run repeated CPU/CUDA
+A/B tests without downloading the 379 GB checkpoint.
+"""
+
+import argparse
+import json
+from pathlib import Path
+
+import torch
+from transformers import GlmMoeDsaConfig, GlmMoeDsaForCausalLM
+
+
+def build_config() -> GlmMoeDsaConfig:
+    return GlmMoeDsaConfig(
+        vocab_size=8192,
+        hidden_size=1024,
+        intermediate_size=2048,
+        moe_intermediate_size=512,
+        num_hidden_layers=8,
+        first_k_dense_replace=3,
+        num_attention_heads=16,
+        num_key_value_heads=16,
+        n_routed_experts=32,
+        num_experts_per_tok=8,
+        n_shared_experts=1,
+        q_lora_rank=256,
+        kv_lora_rank=128,
+        qk_nope_head_dim=64,
+        qk_rope_head_dim=32,
+        v_head_dim=64,
+        index_topk=4096,
+        index_head_dim=32,
+        index_n_heads=4,
+        n_group=1,
+        topk_group=1,
+        norm_topk_prob=True,
+        routed_scaling_factor=2.5,
+        rope_parameters={"rope_type": "default", "rope_theta": 10000.0},
+        tie_word_embeddings=False,
+        rms_norm_eps=1e-5,
+        attention_bias=False,
+        max_position_embeddings=4096,
+    )
+
+
+def main() -> None:
+    parser = argparse.ArgumentParser()
+    parser.add_argument("--output", default="glm_bench_medium")
+    parser.add_argument("--device", default="cuda" if torch.cuda.is_available() else "cpu")
+    parser.add_argument("--seed", type=int, default=1234)
+    args = parser.parse_args()
+
+    torch.manual_seed(args.seed)
+    cfg = build_config()
+    cfg._attn_implementation = "eager"
+    model = GlmMoeDsaForCausalLM(cfg).eval()
+    with torch.no_grad():
+        for param in model.parameters():
+            if param.dim() >= 2:
+                param.normal_(0, 0.02)
+        for layer in model.model.layers:
+            if hasattr(layer.mlp, "gate"):
+                layer.mlp.gate.e_score_correction_bias.copy_(
+                    torch.linspace(-0.1, 0.1, cfg.n_routed_experts)
+                )
+
+    output = Path(args.output)
+    output.mkdir(parents=True, exist_ok=True)
+    params = sum(p.numel() for p in model.parameters())
+    model.save_pretrained(output, safe_serialization=True, max_shard_size="4GB")
+
+    model.to(args.device)
+    prompt = [3, 14, 159, 26, 53, 58, 200, 11, 77, 240, 5, 99]
+    ids = torch.tensor([prompt], device=args.device)
+    with torch.inference_mode():
+        full = model.generate(ids, max_new_tokens=8, do_sample=False, use_cache=True)[0]
+        logits = model(full.unsqueeze(0), use_cache=False).logits[0]
+
+    ref = {
+        "prompt_ids": prompt,
+        "full_ids": full.cpu().tolist(),
+        "tf_pred": logits.argmax(-1).cpu().tolist(),
+    }
+    (output / "ref_glm.json").write_text(json.dumps(ref))
+    manifest = {
+        "seed": args.seed,
+        "parameters": params,
+        "parameters_billions": round(params / 1e9, 4),
+        "purpose": "backend benchmark fixture; random weights, not a language model",
+    }
+    (output / "bench_manifest.json").write_text(json.dumps(manifest, indent=2))
+    print(json.dumps(manifest, indent=2))
+
+
+if __name__ == "__main__":
+    main()