From 57706a02007d68ec6ffc7f6aa83ab895286395ff Mon Sep 17 00:00:00 2001 From: ZacharyZcR Date: Fri, 10 Jul 2026 13:41:09 +0800 Subject: [PATCH] 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 | 3 + README.md | 64 ++++++++++ c/Makefile | 32 ++++- c/backend_cuda.cu | 230 ++++++++++++++++++++++++++++++++++++ c/backend_cuda.h | 49 ++++++++ c/backend_cuda_test.cu | 81 +++++++++++++ c/benchmark_cuda_fixture.py | 102 ++++++++++++++++ c/glm.c | 229 +++++++++++++++++++++++++++++++++-- c/make_glm_bench_model.py | 99 ++++++++++++++++ 9 files changed, 878 insertions(+), 11 deletions(-) create mode 100644 c/backend_cuda.cu create mode 100644 c/backend_cuda.h create mode 100644 c/backend_cuda_test.cu create mode 100644 c/benchmark_cuda_fixture.py create mode 100644 c/make_glm_bench_model.py diff --git a/.gitignore b/.gitignore index 667044d..2132b6f 100644 --- a/.gitignore +++ b/.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 diff --git a/README.md b/README.md index 2b1a1f4..53fa86d 100644 --- a/README.md +++ b/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. diff --git a/c/Makefile b/c/Makefile index e48d40d..d3cb210 100644 --- a/c/Makefile +++ b/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 - $(CC) $(CFLAGS) glm.c -o glm $(LDFLAGS) +glm: glm.c st.h json.h tok.h tok_unicode.h compat.h $(CUDA_OBJ) + $(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 diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu new file mode 100644 index 0000000..b40d361 --- /dev/null +++ b/c/backend_cuda.cu @@ -0,0 +1,230 @@ +#include "backend_cuda.h" + +#include + +#include +#include + +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(weights) + row; + if (fmt == 0) return reinterpret_cast(base)[i]; + if (fmt == 1) return static_cast(reinterpret_cast(base)[i]); + const uint8_t *q = base; + if (fmt == 2) { + uint8_t v = q[i >> 1]; + return static_cast(((i & 1) ? (v >> 4) : (v & 15)) - 8); + } + uint8_t v = q[i >> 2]; + return static_cast(((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(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<<>>(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; +} diff --git a/c/backend_cuda.h b/c/backend_cuda.h new file mode 100644 index 0000000..5a7377b --- /dev/null +++ b/c/backend_cuda.h @@ -0,0 +1,49 @@ +#ifndef COLIBRI_BACKEND_CUDA_H +#define COLIBRI_BACKEND_CUDA_H + +#include +#include + +#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 diff --git a/c/backend_cuda_test.cu b/c/backend_cuda_test.cu new file mode 100644 index 0000000..cc28ef6 --- /dev/null +++ b/c/backend_cuda_test.cu @@ -0,0 +1,81 @@ +#include "backend_cuda.h" + +#include +#include +#include +#include + +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; +} diff --git a/c/benchmark_cuda_fixture.py b/c/benchmark_cuda_fixture.py new file mode 100644 index 0000000..5b52258 --- /dev/null +++ b/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() diff --git a/c/glm.c b/c/glm.c index 24b6abd..c4ef863 100644 --- a/c/glm.c +++ b/c/glm.c @@ -22,12 +22,17 @@ #include #include #include +#include #include #if defined(__APPLE__) || defined(__linux__) #include /* mlock: inchioda le pagine in RAM / wire pages into RAM */ #endif #include "st.h" #include "tok.h" +#ifdef COLI_CUDA +#include +#include "backend_cuda.h" +#endif #ifdef __AVX2__ #include 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;iINT_MAX||n>=COLI_CUDA_MAX_DEVICES) return 0; + for(int i=0;icuda_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;sg.cuda_eligible) m->gpu_expert_calls++; +#endif for(int r=0;rg, 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;ihits+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; - printf("PROFILO: expert-disk %.1fs | expert-matmul %.1fs | attention %.1fs (di cui kvb %.1fs) | lm_head %.1fs | altro %.1fs\n", - m->t_edisk, m->t_emm, m->t_attn, m->t_kvb, m->t_head, dt-acc); +#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 + 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;isafe_total) budget=safe_total; + for(int a=0;agpu_expert_bytesnpin[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=need && + (best<0||placed_b[i]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;i1?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 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()