url/colibri
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness

Browse Source 
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)
This commit is contained in:
woolcoxm
2026-07-16 01:09:47 -04:00
parent 86987154f5
commit 6e29260635
8 changed files with 1524 additions and 49 deletions
Download Patch File Download Diff File Expand all files Collapse all files
c/backend_cuda.cu
+56 -21
View File
@@ -21,6 +21,7 @@ struct ColiCudaTensor {
size_t weight_bytes;
int fmt, I, O, device;
int gs; /* quant group size; 0 = per-row scales (#334) */
int ng; /* number of scale groups per row = ceil(I/gs) for fmt=4 */
size_t scale_count; /* floats in `scales`: O per-row, O*ng grouped */
int tracked;
RaggedKVEntry ragged[512];
@@ -85,7 +86,7 @@ static int select_ctx(DeviceContext *ctx) {
__host__ __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 == 2 || fmt == 4) return (size_t)(I + 1) / 2; /* fmt=4: same packed int4 */
if (fmt == 3) return (size_t)(I + 3) / 4;
if (fmt == 4) return (size_t)(I + 1) / 2; /* grouped int4: nibbles like fmt 2 */
return 0;
@@ -96,7 +97,7 @@ __device__ static float weight_at(const void *weights, int fmt, size_t row, int
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) {
if (fmt == 2 || fmt == 4) { /* fmt=4: same nibble layout */
uint8_t v = q[i >> 1];
int n=(i&1)?(v>>4):(v&15); return static_cast<float>(n&8?n-16:n);
}
@@ -110,14 +111,27 @@ __global__ static void offset_to_signed_s4(uint8_t *q,size_t n){
__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) {
size_t rb, int gs, int ng) {
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);
if (fmt == 4) {
/* Grouped int4: one f32 scale per gs elements along I (ng groups per row).
* Scale layout: scales[o*ng + g]. Each thread strides through I, applying
* the appropriate group scale as it crosses group boundaries. This matches
* the CPU matmul_i4_grouped accumulation exactly. */
const float *scl = scales + (size_t)o * ng;
for (int i = threadIdx.x; i < I; i += blockDim.x) {
int g = i / gs;
if (g >= ng) g = ng - 1; /* tail elements in the last (partial) group */
sum += xs[i] * weight_at(weights, fmt, row, i) * scl[g];
}
} else {
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;
@@ -127,7 +141,7 @@ __global__ static void quant_matmul(float *y, const float *x, const void *weight
__syncthreads();
}
if (!threadIdx.x)
y[(size_t)s * O + o] = partial[0] * (fmt ? scales[o] : 1.0f);
y[(size_t)s * O + o] = (fmt && fmt != 4) ? partial[0] * scales[o] : partial[0];
}
__global__ static void silu_mul(float *gate, const float *up, size_t n) {
@@ -564,7 +578,7 @@ extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
* on such a slot failed here — the GPU tier silently never computed
* for host-released slab experts. */
ColiCudaTensor *t = *tensor;
return t->fmt == fmt && t->I == I && t->O == O && t->device == device;
return t->fmt == fmt && t->I == I && t->O == O && t->device == device && t->gs == gs;
}
DeviceContext *ctx = find_ctx(device);
if (!weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0;
@@ -574,7 +588,8 @@ extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
if (!t) return 0;
t->fmt = fmt; t->I = I; t->O = O; t->device = device; t->weight_bytes = rb * (size_t)O;
t->gs = (fmt==4 && g_upload_gs>0) ? g_upload_gs : 0;
t->scale_count = t->gs ? (size_t)O * (size_t)((I + t->gs - 1) / t->gs) : (size_t)O;
t->ng = t->gs ? (I + t->gs - 1) / t->gs : 1;
t->scale_count = t->gs ? (size_t)O * (size_t)t->ng : (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);
@@ -618,7 +633,9 @@ extern "C" int coli_cuda_tensor_update(ColiCudaTensor *tensor,
(uint8_t*)tensor->weights,tensor->weight_bytes);
if(!cuda_ok(cudaGetLastError(),"int4 weight refresh")) return 0;
}
int ng = tensor->ng > 0 ? tensor->ng : 1;
return !tensor->fmt || cuda_ok(cudaMemcpy(tensor->scales,scales,
<<<<<<< HEAD
(tensor->scale_count?tensor->scale_count:(size_t)tensor->O)*sizeof(float),
cudaMemcpyHostToDevice),"scale refresh");
}
@@ -631,14 +648,22 @@ static long g_gpu_calls;
static int fault_injected(void) {
const char *fa = std::getenv("COLI_GPU_FAIL_AFTER");
return fa && g_gpu_calls++ >= std::atol(fa);
=======
(size_t)tensor->O*ng*sizeof(float),cudaMemcpyHostToDevice),"scale refresh");
>>>>>>> 1a69113 (cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness)
}
extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor,
float *y, const float *x,
const void *weights, const float *scales,
<<<<<<< HEAD
int fmt, int S, int I, int O, int device) {
if (fault_injected()) return 0;
if (S < 1 || !coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device)) return 0;
=======
int fmt, int S, int I, int O, int device, int gs) {
if (S < 1 || !coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device, gs)) return 0;
>>>>>>> 1a69113 (cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness)
ColiCudaTensor *t = *tensor;
DeviceContext *ctx = find_ctx(t->device);
if (!select_ctx(ctx)) return 0;
@@ -647,7 +672,7 @@ extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor,
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);
quant_matmul<<<grid, 256>>>(ctx->y, ctx->x, t->weights, t->scales, fmt, S, I, O, rb, t->gs, t->ng);
if (!cuda_ok(cudaGetLastError(), "matmul launch") ||
!cuda_ok(cudaMemcpy(y, ctx->y, yb, cudaMemcpyDeviceToHost), "output download")) return 0;
return 1;
@@ -671,13 +696,13 @@ extern "C" int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up,
if (!cuda_ok(cudaMemcpy(ctx->x,x,xb,cudaMemcpyHostToDevice),"expert input upload")) return 0;
dim3 hidden_grid((unsigned)I,(unsigned)S), output_grid((unsigned)D,(unsigned)S);
quant_matmul<<<hidden_grid,256>>>(ctx->gate,ctx->x,gate->weights,gate->scales,
gate->fmt,S,D,I,row_bytes(gate->fmt,D));
gate->fmt,S,D,I,row_bytes(gate->fmt,D),gate->gs,gate->ng);
quant_matmul<<<hidden_grid,256>>>(ctx->up,ctx->x,up->weights,up->scales,
up->fmt,S,D,I,row_bytes(up->fmt,D));
up->fmt,S,D,I,row_bytes(up->fmt,D),up->gs,up->ng);
size_t n=(size_t)S*I;
silu_mul<<<(unsigned)((n+255)/256),256>>>(ctx->gate,ctx->up,n);
quant_matmul<<<output_grid,256>>>(ctx->y,ctx->gate,down->weights,down->scales,
down->fmt,S,I,D,row_bytes(down->fmt,I));
down->fmt,S,I,D,row_bytes(down->fmt,I),down->gs,down->ng);
if (!cuda_ok(cudaGetLastError(),"expert MLP launch") ||
!cuda_ok(cudaMemcpy(y,ctx->y,yb,cudaMemcpyDeviceToHost),"expert output download")) return 0;
return 1;
@@ -732,9 +757,16 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates,
g->fmt,u->fmt,d->fmt,rows[c],total,
g->gs,u->gs,d->gs};
all_s4&=g->fmt==2&&u->fmt==2&&d->fmt==2;
<<<<<<< HEAD
all_q4&=(g->fmt==2||g->fmt==4)&&(u->fmt==2||u->fmt==4)&&(d->fmt==2||d->fmt==4)&&
!(g->gs&1)&&!(u->gs&1)&&!(d->gs&1); /* even gs: a packed byte never straddles groups */
any_g4|=g->fmt==4||u->fmt==4||d->fmt==4;
=======
/* fmt=4 (grouped int4) experts have per-group scales that the grouped_hidden/
* grouped_down kernels don't handle. Fall back to the per-expert path
* (coli_cuda_expert_mlp), which uses quant_matmul with correct gs/ng. */
if(g->fmt==4||u->fmt==4||d->fmt==4) return 0;
>>>>>>> 1a69113 (cuda+engine: full fmt=4 (grouped int4 gs=64) support + diagnostic harness)
total+=rows[c]; if(rows[c]>max_rows) max_rows=rows[c];
}
DeviceContext *ctx=find_ctx(device); if(!select_ctx(ctx)) return 0;
@@ -805,12 +837,12 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates,
/* piccoli batch: tile TC quasi vuoti + overhead di lancio — il
* kernel naive per-elemento resta piu' veloce (misurato in decode) */
quant_matmul<<<dim3((unsigned)I,(unsigned)r),256,0,ctx->stream>>>(g16,x16,
host[c].g,host[c].gs,host[c].gf,r,D,I,row_bytes(host[c].gf,D));
host[c].g,host[c].gs,host[c].gf,r,D,I,row_bytes(host[c].gf,D),0,1);
quant_matmul<<<dim3((unsigned)I,(unsigned)r),256,0,ctx->stream>>>(u16,x16,
host[c].u,host[c].us,host[c].uf,r,D,I,row_bytes(host[c].uf,D));
host[c].u,host[c].us,host[c].uf,r,D,I,row_bytes(host[c].uf,D),0,1);
silu_mul<<<(unsigned)(((size_t)r*I+255)/256),256,0,ctx->stream>>>(g16,u16,(size_t)r*I);
quant_matmul<<<dim3((unsigned)D,(unsigned)r),256,0,ctx->stream>>>(y16,g16,
host[c].d,host[c].ds,host[c].df,r,I,D,row_bytes(host[c].df,I));
host[c].d,host[c].ds,host[c].df,r,I,D,row_bytes(host[c].df,I),0,1);
}
off16+=r;
}
@@ -976,7 +1008,7 @@ static int attention_absorb_batch_run(ColiCudaTensor *w,ColiCudaTensor *proj,flo
if(proj){
ob=(size_t)S*proj->O*sizeof(float);if(!reserve(&dc->y,&dc->y_cap,ob))return 0;
quant_matmul<<<dim3(proj->O,S),256,0,dc->stream>>>(dc->y,dc->ac,proj->weights,
proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I));
proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I),proj->gs,proj->ng);
if(!cuda_ok(cudaGetLastError(),"attention o_proj launch"))return 0;src=dc->y;
}
if(!cuda_ok(cudaMemcpyAsync(out,src,ob,cudaMemcpyDeviceToHost,dc->stream),
@@ -1087,7 +1119,8 @@ extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) {
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);
int ng = tensor->ng > 0 ? tensor->ng : 1;
size_t bytes = tensor->weight_bytes + (tensor->fmt ? (size_t)tensor->O * ng * sizeof(float) : 0);
if (ctx->tensor_count) ctx->tensor_count--;
if (ctx->tensor_bytes >= bytes) ctx->tensor_bytes -= bytes;
}
@@ -1101,7 +1134,9 @@ extern "C" void coli_cuda_tensor_free(ColiCudaTensor *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;
if (!tensor) return 0;
int ng = tensor->ng > 0 ? tensor->ng : 1;
return tensor->weight_bytes + (tensor->fmt ? (size_t)tensor->O * ng * sizeof(float) : 0);
}
extern "C" int coli_cuda_tensor_device(const ColiCudaTensor *tensor) {
@@ -1320,7 +1355,7 @@ extern "C" int coli_cuda_attention_project_batch_dev(ColiCudaTensor *w,ColiCudaT
size_t ob=(size_t)S*proj->O*sizeof(float);
if(!reserve(&dc->y,&dc->y_cap,ob))return 0;
quant_matmul<<<dim3(proj->O,S),256,0,dc->stream>>>(dc->y,dc->ac,proj->weights,
proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I));
proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I),proj->gs,proj->ng);
if(!cuda_ok(cudaGetLastError(),"pipe o_proj launch"))return 0;
if(!cuda_ok(cudaMemcpyAsync(out,dc->y,ob,cudaMemcpyDeviceToHost,dc->stream),"pipe attention download")||
!cuda_ok(cudaStreamSynchronize(dc->stream),"pipe attention sync"))return 0;
@@ -1355,7 +1390,7 @@ extern "C" int coli_cuda_pipe_gemm(ColiCudaTensor *t,float *y_dev,const float *x
DeviceContext *ctx=find_ctx(t->device); if(!select_ctx(ctx)) return 0;
dim3 grid((unsigned)t->O,(unsigned)S);
quant_matmul<<<grid,256>>>(y_dev,x_dev,t->weights,t->scales,t->fmt,S,t->I,t->O,
row_bytes(t->fmt,t->I));
row_bytes(t->fmt,t->I),t->gs,t->ng);
return cuda_ok(cudaGetLastError(),"pipe gemm");
}
/* copia diretta scheda->scheda (P2P se disponibile, altrimenti staging driver) */
@@ -1382,7 +1417,7 @@ extern "C" int coli_cuda_attention_project_batch_dev_out(ColiCudaTensor *w,ColiC
rope_dev,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
if(!cuda_ok(cudaGetLastError(),"pipe attention launch (dev out)"))return 0;
quant_matmul<<<dim3(proj->O,S),256,0,dc->stream>>>(out_dev,dc->ac,proj->weights,
proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I));
proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I),proj->gs,proj->ng);
if(!cuda_ok(cudaGetLastError(),"pipe o_proj launch (dev out)"))return 0;
return cuda_ok(cudaStreamSynchronize(dc->stream),"pipe attention sync (dev out)");
}