Merge #298: per-group scales in CUDA dense/attention kernels (fmt=4 correctness)
Fixes a live correctness bug on dev: with CUDA_DENSE=1 on a g64 (fmt=4) container, the dense matmul and attention-absorb kernels applied scales per-row (scales[o] / wscale[row]) while the uploaded scale array is per-group [O × ceil(I/gs)] — wrong scale for nearly every row, garbage output ('odesk odesk…'). quant_matmul now applies group scales inline (matching CPU matmul_i4_grouped exactly), absorb_scale handles fmt=4 in the attention kernels (the kv_b crash), w4a16/TC fast paths stay correctly gated to fmt=2, and a fused CPU gate+up pair (matmul_i4_grouped_pair, AVX2) lands as a bonus. Also adds the fmt=4→CPU fallback log requested in review.
Credits: @woolcoxm (author, rebase onto post-#391 dev), @mohamedmastouri2000-boop (root-cause isolation + hardware verification on RTX 5080/sm_120: garbage→coherent, 952 dense tensors + 109 experts fully VRAM-resident, zero fallbacks). Verified locally: clean build, token-exact tiny models unchanged; CI 8/8 green.
This commit is contained in:
Vincenzo Fornaro
+76
-40
@@ -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);
|
||||
}
|
||||
@@ -104,20 +105,45 @@ __device__ static float weight_at(const void *weights, int fmt, size_t row, int
|
||||
return static_cast<float>(((v >> ((i & 3) * 2)) & 3) - 2);
|
||||
}
|
||||
|
||||
/* Scale for output `row`, input element `k`. fmt=4 (grouped int4) stores ng
|
||||
* scales per row at scales[row*ng + k/gs]; every other quantized format has
|
||||
* one scale per row at scales[row]. Mirrors quant_matmul's fmt==4 branch so the
|
||||
* attention absorb kernels apply per-group scales instead of the per-row
|
||||
* (fmt=2) semantic that crashed #298's g64 kv_b. */
|
||||
__device__ static float absorb_scale(const float *wscale, int fmt, int gs, int ng, int row, int k) {
|
||||
if (!fmt) return 1.f;
|
||||
if (fmt != 4) return wscale[row];
|
||||
int g = k / gs; if (g >= ng) g = ng - 1; /* tail of the last (partial) group */
|
||||
return wscale[(size_t)row * ng + g];
|
||||
}
|
||||
|
||||
__global__ static void offset_to_signed_s4(uint8_t *q,size_t n){
|
||||
size_t i=(size_t)blockIdx.x*blockDim.x+threadIdx.x;if(i<n)q[i]^=0x88;
|
||||
}
|
||||
|
||||
__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 +153,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) {
|
||||
@@ -341,11 +367,12 @@ __global__ static void grouped_down_g4(float *y,const float *x,const GroupDesc *
|
||||
|
||||
__global__ static void attention_absorb_kernel(float *ctx,const float *q,const float *latent,
|
||||
const float *rope,const void *weights,const float *wscale,
|
||||
int fmt,int H,int Q,int R,int V,int K,int T,float scale){
|
||||
int fmt,int H,int Q,int R,int V,int K,int T,float scale,
|
||||
int gs,int ng){
|
||||
int h=blockIdx.x,tid=threadIdx.x,rbase=h*(Q+V);extern __shared__ float sm[];
|
||||
float *qa=sm,*cl=qa+K,*scores=cl+K;
|
||||
for(int k=tid;k<K;k+=blockDim.x){float a=0;for(int d=0;d<Q;d++)
|
||||
a+=q[(size_t)h*(Q+R)+d]*weight_at(weights,fmt,(size_t)(rbase+d)*row_bytes(fmt,K),k)*(fmt?wscale[rbase+d]:1.f);qa[k]=a;}
|
||||
a+=q[(size_t)h*(Q+R)+d]*weight_at(weights,fmt,(size_t)(rbase+d)*row_bytes(fmt,K),k)*absorb_scale(wscale,fmt,gs,ng,rbase+d,k);qa[k]=a;}
|
||||
__syncthreads();
|
||||
for(int t=tid;t<T;t+=blockDim.x){float a=0;const float *lt=latent+(size_t)t*K,*rt=rope+(size_t)t*R;
|
||||
for(int k=0;k<K;k++)a+=qa[k]*lt[k];for(int d=0;d<R;d++)a+=q[(size_t)h*(Q+R)+Q+d]*rt[d];scores[t]=a*scale;}
|
||||
@@ -356,19 +383,20 @@ __global__ static void attention_absorb_kernel(float *ctx,const float *q,const f
|
||||
for(int k=tid;k<K;k+=blockDim.x){float a=0;for(int t=0;t<T;t++)a+=scores[t]*latent[(size_t)t*K+k];cl[k]=a;}
|
||||
__syncthreads();
|
||||
for(int v=tid;v<V;v+=blockDim.x){int row=rbase+Q+v;float a=0;size_t rb=row_bytes(fmt,K);
|
||||
for(int k=0;k<K;k++)a+=cl[k]*weight_at(weights,fmt,(size_t)row*rb,k);ctx[(size_t)h*V+v]=a*(fmt?wscale[row]:1.f);}
|
||||
for(int k=0;k<K;k++)a+=cl[k]*weight_at(weights,fmt,(size_t)row*rb,k)*absorb_scale(wscale,fmt,gs,ng,row,k);ctx[(size_t)h*V+v]=a;}
|
||||
}
|
||||
|
||||
__global__ static void attention_absorb_batch_kernel(float *ctx,const float *q,
|
||||
const float *latent,const float *rope,const void *weights,const float *wscale,
|
||||
int fmt,int S,int H,int Q,int R,int V,int K,int T,float scale){
|
||||
int fmt,int S,int H,int Q,int R,int V,int K,int T,float scale,
|
||||
int gs,int ng){
|
||||
int s=blockIdx.y,h=blockIdx.x,tid=threadIdx.x,nt=T-S+s+1,rbase=h*(Q+V);
|
||||
if(s>=S||nt<1)return;
|
||||
extern __shared__ float sm[];float *qa=sm,*cl=qa+K,*scores=cl+K,*red=scores+T;
|
||||
const float *qs=q+((size_t)s*H+h)*(Q+R);
|
||||
for(int k=tid;k<K;k+=blockDim.x){float a=0;for(int d=0;d<Q;d++)
|
||||
a+=qs[d]*weight_at(weights,fmt,(size_t)(rbase+d)*row_bytes(fmt,K),k)*
|
||||
(fmt?wscale[rbase+d]:1.f);qa[k]=a;}
|
||||
absorb_scale(wscale,fmt,gs,ng,rbase+d,k);qa[k]=a;}
|
||||
__syncthreads();
|
||||
for(int t=tid;t<nt;t+=blockDim.x){float a=0;const float *lt=latent+(size_t)t*K;
|
||||
const float *rt=rope+(size_t)t*R;for(int k=0;k<K;k++)a+=qa[k]*lt[k];
|
||||
@@ -386,8 +414,8 @@ __global__ static void attention_absorb_batch_kernel(float *ctx,const float *q,
|
||||
a+=scores[t]*latent[(size_t)t*K+k];cl[k]=a;}
|
||||
__syncthreads();
|
||||
for(int v=tid;v<V;v+=blockDim.x){int row=rbase+Q+v;float a=0;size_t rb=row_bytes(fmt,K);
|
||||
for(int k=0;k<K;k++)a+=cl[k]*weight_at(weights,fmt,(size_t)row*rb,k);
|
||||
ctx[((size_t)s*H+h)*V+v]=a*(fmt?wscale[row]:1.f);}
|
||||
for(int k=0;k<K;k++)a+=cl[k]*weight_at(weights,fmt,(size_t)row*rb,k)*absorb_scale(wscale,fmt,gs,ng,row,k);
|
||||
ctx[((size_t)s*H+h)*V+v]=a;}
|
||||
}
|
||||
|
||||
/* Independent device-resident KV sequence per row. lengths selects the valid
|
||||
@@ -564,7 +592,8 @@ 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;
|
||||
int want_gs = (fmt==4 && g_upload_gs>0) ? g_upload_gs : 0;
|
||||
return t->fmt == fmt && t->I == I && t->O == O && t->device == device && t->gs == want_gs;
|
||||
}
|
||||
DeviceContext *ctx = find_ctx(device);
|
||||
if (!weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0;
|
||||
@@ -574,7 +603,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,6 +648,7 @@ 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,
|
||||
(tensor->scale_count?tensor->scale_count:(size_t)tensor->O)*sizeof(float),
|
||||
cudaMemcpyHostToDevice),"scale refresh");
|
||||
@@ -636,9 +667,11 @@ static int fault_injected(void) {
|
||||
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) {
|
||||
int fmt, int S, int I, int O, int device, int gs) {
|
||||
if (fault_injected()) return 0;
|
||||
if (S < 1 || !coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device)) return 0;
|
||||
if (S < 1) return 0;
|
||||
if (gs > 0) { if (!coli_cuda_tensor_upload_g(tensor, weights, scales, fmt, I, O, device, gs)) return 0; }
|
||||
else { if (!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;
|
||||
@@ -647,7 +680,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 +704,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;
|
||||
@@ -805,12 +838,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;
|
||||
}
|
||||
@@ -905,12 +938,12 @@ extern "C" int coli_cuda_expert_group_issue(ColiCudaTensor *const *gates,
|
||||
float *g16=ctx->gate+(size_t)host[c].offset*I,*u16=ctx->up+(size_t)host[c].offset*I;
|
||||
float *x16=ctx->x+(size_t)host[c].offset*D,*y16=ctx->y+(size_t)host[c].offset*D;
|
||||
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);
|
||||
}
|
||||
if(!cuda_ok(cudaGetLastError(),"expert group issue launch")||
|
||||
!cuda_ok(cudaMemcpyAsync(ctx->host_y,ctx->y,xb,cudaMemcpyDeviceToHost,ctx->stream),
|
||||
@@ -947,7 +980,7 @@ extern "C" int coli_cuda_attention_absorb(ColiCudaTensor *w,float *ctx,const flo
|
||||
!cuda_ok(cudaMemcpyAsync(dc->ar,rope,rb,cudaMemcpyHostToDevice,dc->stream),"attention rope upload"))return 0;
|
||||
size_t shared=(size_t)(2*K+T)*sizeof(float);
|
||||
attention_absorb_kernel<<<H,256,shared,dc->stream>>>(dc->ac,dc->aq,dc->al,dc->ar,w->weights,w->scales,
|
||||
w->fmt,H,Q,R,V,K,T,scale);
|
||||
w->fmt,H,Q,R,V,K,T,scale,w->gs,w->ng);
|
||||
if(!cuda_ok(cudaGetLastError(),"attention absorb launch")||
|
||||
!cuda_ok(cudaMemcpyAsync(ctx,dc->ac,cb,cudaMemcpyDeviceToHost,dc->stream),"attention context download")||
|
||||
!cuda_ok(cudaStreamSynchronize(dc->stream),"attention synchronize"))return 0;
|
||||
@@ -970,13 +1003,13 @@ static int attention_absorb_batch_run(ColiCudaTensor *w,ColiCudaTensor *proj,flo
|
||||
!cuda_ok(cudaMemcpyAsync(dc->ar,rope,rb,cudaMemcpyHostToDevice,dc->stream),"attention batch rope upload"))return 0;
|
||||
size_t shared=(size_t)(2*K+T+256)*sizeof(float);
|
||||
attention_absorb_batch_kernel<<<dim3(H,S),256,shared,dc->stream>>>(dc->ac,dc->aq,dc->al,
|
||||
dc->ar,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
|
||||
dc->ar,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale,w->gs,w->ng);
|
||||
if(!cuda_ok(cudaGetLastError(),"attention batch launch"))return 0;
|
||||
const float *src=dc->ac;size_t ob=cb;
|
||||
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),
|
||||
@@ -1076,7 +1109,7 @@ extern "C" int coli_cuda_attention_project_ragged(ColiCudaTensor *w,ColiCudaTens
|
||||
attention_absorb_ragged_kernel<<<dim3(H,S),256,shared,dc->stream>>>(dc->ac,dc->aq,ddl,ddr,
|
||||
dn,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
|
||||
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);
|
||||
return cuda_ok(cudaGetLastError(),"ragged attention launch")&&
|
||||
cuda_ok(cudaMemcpyAsync(out,dc->y,ob,cudaMemcpyDeviceToHost,dc->stream),"ragged output download")&&
|
||||
cuda_ok(cudaStreamSynchronize(dc->stream),"ragged attention synchronize");
|
||||
@@ -1087,7 +1120,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 +1135,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) {
|
||||
@@ -1315,12 +1351,12 @@ extern "C" int coli_cuda_attention_project_batch_dev(ColiCudaTensor *w,ColiCudaT
|
||||
if(!reserve(&dc->ac,&dc->ac_cap,cb))return 0;
|
||||
size_t shared=(size_t)(2*K+T+256)*sizeof(float);
|
||||
attention_absorb_batch_kernel<<<dim3(H,S),256,shared,dc->stream>>>(dc->ac,q_dev,latent_dev,
|
||||
rope_dev,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
|
||||
rope_dev,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale,w->gs,w->ng);
|
||||
if(!cuda_ok(cudaGetLastError(),"pipe attention launch"))return 0;
|
||||
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 +1391,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) */
|
||||
@@ -1379,10 +1415,10 @@ extern "C" int coli_cuda_attention_project_batch_dev_out(ColiCudaTensor *w,ColiC
|
||||
if(!reserve(&dc->ac,&dc->ac_cap,cb))return 0;
|
||||
size_t shared=(size_t)(2*K+T+256)*sizeof(float);
|
||||
attention_absorb_batch_kernel<<<dim3(H,S),256,shared,dc->stream>>>(dc->ac,q_dev,latent_dev,
|
||||
rope_dev,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
|
||||
rope_dev,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale,w->gs,w->ng);
|
||||
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)");
|
||||
}
|
||||
@@ -1398,7 +1434,7 @@ extern "C" int coli_cuda_attention_absorb_batch_dev(ColiCudaTensor *w,float *ctx
|
||||
DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0;
|
||||
size_t shared=(size_t)(2*K+T+256)*sizeof(float);
|
||||
attention_absorb_batch_kernel<<<dim3(H,S),256,shared,dc->stream>>>(ctx_dev,q_dev,latent_dev,
|
||||
rope_dev,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
|
||||
rope_dev,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale,w->gs,w->ng);
|
||||
if(!cuda_ok(cudaGetLastError(),"pipe shard attention launch"))return 0;
|
||||
return cuda_ok(cudaStreamSynchronize(dc->stream),"pipe shard attention sync");
|
||||
}
|
||||
@@ -1416,7 +1452,7 @@ extern "C" int coli_cuda_attention_absorb_kvdev(ColiCudaTensor *w,float *ctx,con
|
||||
if(!cuda_ok(cudaMemcpyAsync(dc->aq,q,qb,cudaMemcpyHostToDevice,dc->stream),"kvdev q upload"))return 0;
|
||||
size_t shared=(size_t)(2*K+T+256)*sizeof(float);
|
||||
attention_absorb_batch_kernel<<<dim3(H,1),256,shared,dc->stream>>>(dc->ac,dc->aq,latent_dev,
|
||||
rope_dev,w->weights,w->scales,w->fmt,1,H,Q,R,V,K,T,scale);
|
||||
rope_dev,w->weights,w->scales,w->fmt,1,H,Q,R,V,K,T,scale,w->gs,w->ng);
|
||||
if(!cuda_ok(cudaGetLastError(),"kvdev absorb launch")||
|
||||
!cuda_ok(cudaMemcpyAsync(ctx,dc->ac,cb,cudaMemcpyDeviceToHost,dc->stream),"kvdev ctx download")||
|
||||
!cuda_ok(cudaStreamSynchronize(dc->stream),"kvdev absorb sync"))return 0;
|
||||
|
||||
+4
-3
@@ -45,14 +45,15 @@ COLI_CUDA_DLLEXPORT int coli_cuda_tensor_upload(ColiCudaTensor **tensor,
|
||||
|
||||
/*
|
||||
* 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.
|
||||
* fmt matches QT in glm.c: 0=f32, 1=int8, 2=int4, 3=int2, 4=grouped int4.
|
||||
* gs is the group size for fmt=4 (0 for all other formats).
|
||||
* The first successful call uploads W and its scales; later calls reuse it.
|
||||
* Returns 1 on success and 0 when CUDA is not initialized or the format is invalid.
|
||||
*/
|
||||
COLI_CUDA_DLLEXPORT 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);
|
||||
int fmt, int S, int I, int O, int device, int gs);
|
||||
|
||||
/* Fused expert pipeline: y = down(silu(gate(x)) * up(x)). All three tensors
|
||||
* must already be resident on one device. Activations cross PCIe once in
|
||||
|
||||
+3
-3
@@ -54,7 +54,7 @@ typedef int (*fn_tensor_upload)(ColiCudaTensor **tensor, const void *
|
||||
typedef int (*fn_tensor_upload_g)(ColiCudaTensor **tensor, const void *weights, const float *scales, int fmt, int I, int O, int device, int gs);
|
||||
typedef int (*fn_matmul)(ColiCudaTensor **tensor, float *y, const float *x,
|
||||
const void *weights, const float *scales,
|
||||
int fmt, int S, int I, int O, int device);
|
||||
int fmt, int S, int I, int O, int device, int gs);
|
||||
typedef void (*fn_tensor_free)(ColiCudaTensor *tensor);
|
||||
typedef size_t (*fn_tensor_bytes)(const ColiCudaTensor *tensor);
|
||||
typedef int (*fn_tensor_device)(const ColiCudaTensor *tensor);
|
||||
@@ -337,9 +337,9 @@ int coli_cuda_tensor_upload_g(ColiCudaTensor **tensor, const void *weights, cons
|
||||
|
||||
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){
|
||||
int fmt, int S, int I, int O, int device, int gs){
|
||||
if(!g_cuda.available) return 0;
|
||||
return g_cuda.matmul(tensor, y, x, weights, scales, fmt, S, I, O, device);
|
||||
return g_cuda.matmul(tensor, y, x, weights, scales, fmt, S, I, O, device, gs);
|
||||
}
|
||||
|
||||
void coli_cuda_tensor_free(ColiCudaTensor *tensor){
|
||||
|
||||
+105
-5
@@ -256,9 +256,17 @@ static inline int spec_pinned(void){ return g_spec_pin && g_spec_live; }
|
||||
static void matmul_qt_ex(float *y, const float *x, QT *w, int S, int allow_idot);
|
||||
static void matmul_qt(float *y, const float *x, QT *w, int S){ matmul_qt_ex(y,x,w,S,1); }
|
||||
|
||||
/* fmt=4 fused gate+up (defined later, after the quant kernels) */
|
||||
static void matmul_i4_grouped_pair(float *yg, float *yu, const float *x,
|
||||
const uint8_t *qg, const float *sg,
|
||||
const uint8_t *qu, const float *su,
|
||||
int S, int I, int O, int gs);
|
||||
|
||||
static void expert_gate_up(float *g,float *u,const float *x,QT *wg,QT *wu,int S){
|
||||
if(!g_no_fused_pair&&!spec_pinned()&&S==1&&wg->fmt==2&&wu->fmt==2&&wg->I==wu->I&&wg->O==wu->O)
|
||||
matmul_i4_pair(g,u,x,wg->q4,wg->s,wu->q4,wu->s,wg->I,wg->O);
|
||||
else if(!g_no_fused_pair&&S==1&&wg->fmt==4&&wu->fmt==4&&wg->I==wu->I&&wg->O==wu->O&&wg->gs==wu->gs)
|
||||
matmul_i4_grouped_pair(g,u,x,wg->q4,wg->s,wu->q4,wu->s,S,wg->I,wg->O,wg->gs);
|
||||
else { matmul_qt(g,x,wg,S); matmul_qt(u,x,wu,S); }
|
||||
}
|
||||
|
||||
@@ -448,6 +456,70 @@ static float *falloc(int64_t n){
|
||||
float *p=malloc((size_t)n*sizeof(float)); if(!p){fprintf(stderr,"OOM\n");exit(1);} return p; }
|
||||
|
||||
|
||||
|
||||
/* Fused gate+up for grouped int4 (fmt=4): computes both yg[S,O] and yu[S,O] from
|
||||
* the same x[S,I], reading x once instead of twice — saves ~33% of expert-matmul time at decode.
|
||||
* The per-group scale logic matches matmul_i4_grouped exactly. */
|
||||
static void matmul_i4_grouped_pair(float *yg, float *yu, const float *x,
|
||||
const uint8_t *qg, const float *sg,
|
||||
const uint8_t *qu, const float *su,
|
||||
int S, int I, int O, int gs){
|
||||
int rb=(I+1)/2; int ng=(I+gs-1)/gs;
|
||||
#pragma omp parallel for schedule(static)
|
||||
for(int o=0;o<O;o++){
|
||||
const uint8_t *wg=qg+(int64_t)o*rb; const uint8_t *wu2=qu+(int64_t)o*rb;
|
||||
const float *sgl=sg+(int64_t)o*ng; const float *sul=su+(int64_t)o*ng;
|
||||
for(int s=0;s<S;s++){
|
||||
const float *xs=x+(int64_t)s*I;
|
||||
float ag=0, au=0;
|
||||
for(int g=0; g*gs<I; g++){
|
||||
int base=g*gs; int glen=gs; if(base+glen>I) glen=I-base;
|
||||
float scg=sgl[g], scu=sul[g];
|
||||
int i=base;
|
||||
#ifdef __AVX2__
|
||||
const __m128i m4=_mm_set1_epi8(0x0F); const __m256i b8=_mm256_set1_epi32(8);
|
||||
__m256 accg=_mm256_setzero_ps(), accu=_mm256_setzero_ps();
|
||||
for(; i+16<=base+glen; i+=16){
|
||||
__m128i byg=_mm_loadl_epi64((const __m128i*)(wg+(i>>1)));
|
||||
__m128i log=_mm_and_si128(byg,m4),hig=_mm_and_si128(_mm_srli_epi16(byg,4),m4);
|
||||
__m128i nibg=_mm_unpacklo_epi8(log,hig);
|
||||
__m256 w0g=_mm256_cvtepi32_ps(_mm256_sub_epi32(_mm256_cvtepu8_epi32(nibg),b8));
|
||||
__m256 w1g=_mm256_cvtepi32_ps(_mm256_sub_epi32(_mm256_cvtepu8_epi32(_mm_srli_si128(nibg,8)),b8));
|
||||
accg=_mm256_fmadd_ps(_mm256_loadu_ps(xs+i), w0g, accg);
|
||||
accg=_mm256_fmadd_ps(_mm256_loadu_ps(xs+i+8), w1g, accg);
|
||||
__m128i byu=_mm_loadl_epi64((const __m128i*)(wu2+(i>>1)));
|
||||
__m128i lou=_mm_and_si128(byu,m4),hiu=_mm_and_si128(_mm_srli_epi16(byu,4),m4);
|
||||
__m128i nibu=_mm_unpacklo_epi8(lou,hiu);
|
||||
__m256 w0u=_mm256_cvtepi32_ps(_mm256_sub_epi32(_mm256_cvtepu8_epi32(nibu),b8));
|
||||
__m256 w1u=_mm256_cvtepi32_ps(_mm256_sub_epi32(_mm256_cvtepu8_epi32(_mm_srli_si128(nibu,8)),b8));
|
||||
accu=_mm256_fmadd_ps(_mm256_loadu_ps(xs+i), w0u, accu);
|
||||
accu=_mm256_fmadd_ps(_mm256_loadu_ps(xs+i+8), w1u, accu);
|
||||
}
|
||||
ag+=hsum256(accg)*scg; au+=hsum256(accu)*scu;
|
||||
#endif
|
||||
for(; i+1<base+glen; i+=2){
|
||||
uint8_t bg=wg[i>>1], bu=wu2[i>>1];
|
||||
ag+=(xs[i]*(float)((int)(bg&0xF)-8)+xs[i+1]*(float)((int)(bg>>4)-8))*scg;
|
||||
au+=(xs[i]*(float)((int)(bu&0xF)-8)+xs[i+1]*(float)((int)(bu>>4)-8))*scu;
|
||||
}
|
||||
if(i<base+glen){
|
||||
uint8_t bg=wg[i>>1], bu=wu2[i>>1];
|
||||
ag+=xs[i]*(float)((int)(bg&0xF)-8)*scg;
|
||||
au+=xs[i]*(float)((int)(bu&0xF)-8)*scu;
|
||||
}
|
||||
}
|
||||
yg[(int64_t)s*O+o]=ag; yu[(int64_t)s*O+o]=au;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/* allow_idot=0: forza il kernel int4/int8 ESATTO (attivazioni f32). Serve alle proiezioni di
|
||||
* attenzione: sono sensibili alla quantizzazione int8 delle attivazioni dell'IDOT. Misurato su
|
||||
* GLM-5.2 int4, 1023 token, log-lik -5040.33 (esatto) -> -5160.47 (IDOT) = +0.117 nat/token,
|
||||
* ~+12% perplexity. Gli altri matmul del prefill (o_proj, kv_b, expert) tengono l'IDOT.
|
||||
* EN: allow_idot=0 forces the EXACT int4/int8 kernel (f32 activations). The attention
|
||||
* projections need it: IDOT's int8 activation quantization costs +0.117 nats/token there
|
||||
* (~+12% perplexity), measured. Every other prefill matmul keeps IDOT as before. */
|
||||
static void matmul_qt_ex(float *y, const float *x, QT *w, int S, int allow_idot){
|
||||
#ifdef COLI_METAL
|
||||
if(g_metal_enabled && S>=g_metal_gemm_min && !spec_pinned() && (w->fmt==1||w->fmt==2) && !omp_in_parallel()){
|
||||
@@ -459,7 +531,7 @@ static void matmul_qt_ex(float *y, const float *x, QT *w, int S, int allow_idot)
|
||||
if(g_cuda_enabled && w->cuda_eligible && !w->cuda_failed && w->fmt!=5 && !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;
|
||||
if(coli_cuda_matmul(&w->cuda,y,x,weights,w->s,w->fmt,S,w->I,w->O,w->cuda_device,w->gs)) return;
|
||||
w->cuda_failed=1;
|
||||
fprintf(stderr,"[CUDA] tensor [%d,%d] on device %d disabled after an error; falling back to CPU\n",
|
||||
w->O,w->I,w->cuda_device);
|
||||
@@ -934,13 +1006,14 @@ static void qt_cuda_colocate(QT *dst,const QT *src){
|
||||
}
|
||||
static void layer_cuda_shard_kvb(Layer *l,int H,int Q,int V){
|
||||
if(!g_cuda_enabled||!g_cuda_dense||g_cuda_ndev<2||l->kv_b.fmt==0)return;
|
||||
int rb=l->kv_b.fmt==1?l->kv_b.I:(l->kv_b.fmt==2?(l->kv_b.I+1)/2:(l->kv_b.I+3)/4);
|
||||
int rb=l->kv_b.fmt==1?l->kv_b.I:
|
||||
(l->kv_b.fmt==2||l->kv_b.fmt==4)?(l->kv_b.I+1)/2:(l->kv_b.I+3)/4;
|
||||
const uint8_t *weights=l->kv_b.fmt==1?(const uint8_t*)l->kv_b.q8:l->kv_b.q4;
|
||||
for(int d=0,h0=0;d<g_cuda_ndev;d++){
|
||||
int hn=H/g_cuda_ndev+(d<H%g_cuda_ndev),rows=hn*(Q+V);
|
||||
const void *part=weights+(int64_t)h0*(Q+V)*rb;
|
||||
const float *scale=l->kv_b.s+(int64_t)h0*(Q+V);
|
||||
if(!coli_cuda_tensor_upload(&l->kv_b_shard[d],part,scale,l->kv_b.fmt,l->kv_b.I,rows,g_cuda_devices[d]))return;
|
||||
const float *scale=l->kv_b.s+(int64_t)h0*(Q+V)*(l->kv_b.gs>0?(l->kv_b.I+l->kv_b.gs-1)/l->kv_b.gs:1);
|
||||
if(!coli_cuda_tensor_upload_g(&l->kv_b_shard[d],part,scale,l->kv_b.fmt,l->kv_b.I,rows,g_cuda_devices[d],l->kv_b.gs))return;
|
||||
l->shard_h0[d]=h0;l->shard_hn[d]=hn;l->n_kv_b_shard++;h0+=hn;
|
||||
}
|
||||
int old=-1;for(int i=0;i<g_cuda_ndev;i++)if(g_cuda_devices[i]==l->kv_b.cuda_device)old=i;
|
||||
@@ -1122,6 +1195,14 @@ static void model_init(Model *m, const char *snap, int cap, int ebits, int dbits
|
||||
static void embed_row(Model *m, int tok, float *x){
|
||||
int D=m->c.hidden; QT *e=&m->embed;
|
||||
if(e->fmt==0){ memcpy(x, e->qf+(int64_t)tok*D, D*sizeof(float)); return; }
|
||||
if(e->fmt==4){ /* grouped int4: per-group scale (embed/lm_head at io_bits, usually fmt 0/1) */
|
||||
const uint8_t *q=e->q4+(int64_t)tok*((D+1)/2); int gs=e->gs,ng=(D+gs-1)/gs;
|
||||
const float *scl=e->s+(int64_t)tok*ng;
|
||||
for(int g=0;g*gs<D;g++){ int base=g*gs,glen=gs; if(base+glen>D)glen=D-base; float s=scl[g];
|
||||
for(int i=base;i+1<base+glen;i+=2){ uint8_t byte=q[i>>1]; x[i]=(float)((int)(byte&0xF)-8)*s;
|
||||
x[i+1]=(float)((int)(byte>>4)-8)*s; }
|
||||
if(glen&1){ uint8_t byte=q[(base+glen-1)>>1]; x[base+glen-1]=(float)((int)(byte&0xF)-8)*s; } }
|
||||
return; }
|
||||
if(e->fmt==1){ const int8_t *q=e->q8+(int64_t)tok*D; float s=e->s[tok];
|
||||
for(int i=0;i<D;i++) x[i]=(float)q[i]*s; return; }
|
||||
if(e->fmt==2){ const uint8_t *q=e->q4+(int64_t)tok*((D+1)/2); float s=e->s[tok]; /* int4 */
|
||||
@@ -1885,6 +1966,13 @@ static void qt_matvec_rows(const QT *t, int r0, int n, const float *x, float *y)
|
||||
int I=t->I;
|
||||
for(int j=0;j<n;j++){ int row=r0+j; double a=0;
|
||||
if(t->fmt==0){ const float *w=t->qf+(int64_t)row*I; for(int i=0;i<I;i++) a+=(double)w[i]*x[i]; }
|
||||
else if(t->fmt==4){ /* grouped int4: per-group scale */
|
||||
const uint8_t *w=t->q4+(int64_t)row*((I+1)/2); int gs=t->gs,ng=(I+gs-1)/gs;
|
||||
const float *scl=t->s+(int64_t)row*ng;
|
||||
for(int g=0;g*gs<I;g++){ int base=g*gs,glen=gs; if(base+glen>I)glen=I-base; float sc=scl[g]; float acc=0;
|
||||
for(int i=base;i+1<base+glen;i+=2){ uint8_t b=w[i>>1]; acc+=((int)(b&0xF)-8)*x[i]+((int)(b>>4)-8)*x[i+1]; }
|
||||
if(glen&1){ uint8_t b=w[(base+glen-1)>>1]; acc+=((int)(b&0xF)-8)*x[base+glen-1]; }
|
||||
a+=acc*sc; } }
|
||||
else if(t->fmt==1){ const int8_t *w=t->q8+(int64_t)row*I; float s=t->s[row];
|
||||
float acc=0; for(int i=0;i<I;i++) acc+=(float)w[i]*x[i]; a=acc*s; }
|
||||
else if(t->fmt==2){ const uint8_t *w=t->q4+(int64_t)row*((I+1)/2); float s=t->s[row]; float acc=0;
|
||||
@@ -4479,9 +4567,21 @@ static void run_text(Model *m, const char *snap, const char *prompt, int ngen){
|
||||
grammar_setup(&T); /* metodo F: GRAMMAR=file.gbnf (#48) */
|
||||
if(g_temp<0) g_temp=0.7f; /* auto: 0.7, NON l'1.0 ufficiale — la coda della
|
||||
* distribuzione int4 e' rumore di quantizzazione */
|
||||
int cap=(int)strlen(prompt)+16; int *pids=malloc(cap*sizeof(int));
|
||||
int cap=(int)strlen(prompt)+16; int *pids=malloc((cap+2)*sizeof(int));
|
||||
int np=tok_encode(&T,prompt,(int)strlen(prompt),pids,cap);
|
||||
if(np<1){ fprintf(stderr,"prompt is empty after tokenization\n"); return; }
|
||||
/* GLM prefix (#108): every GLM training sequence starts [gMASK]<sop>; without it the
|
||||
* model is out-of-distribution and generates garbage. Serve/SCORE modes prepend it;
|
||||
* run_text must too. CHAT_TEMPLATE=0 disables (raw prompt, for debugging/non-GLM).
|
||||
* A prompt that already starts with the prefix passes through intact. */
|
||||
int templ=getenv("CHAT_TEMPLATE")?atoi(getenv("CHAT_TEMPLATE")):1;
|
||||
if(templ){
|
||||
int gmask=tok_id_of(&T,"[gMASK]"), sop=tok_id_of(&T,"<sop>");
|
||||
if(gmask>=0 && sop>=0 && (np<2 || pids[0]!=gmask || pids[1]!=sop)){
|
||||
memmove(pids+2,pids,np*sizeof(int));
|
||||
pids[0]=gmask; pids[1]=sop; np+=2;
|
||||
}
|
||||
}
|
||||
printf("prompt: %d tokens | generating up to %d (EOS stop=%d) | n-gram draft=%d\n", np, ngen, eos, g_draft);
|
||||
fputs(prompt,stdout); fflush(stdout);
|
||||
kv_alloc(m, np+ngen+g_draft+2);
|
||||
|
||||
@@ -30,9 +30,9 @@ int main(){
|
||||
for(int i=0;i<D;i++)ds[i]=0.006f+(i%7)*0.0002f;
|
||||
for(size_t i=0;i<x.size();i++)x[i]=std::sin((float)(i+1)*0.013f)*2.f;
|
||||
ColiCudaTensor *g=nullptr,*u=nullptr,*d=nullptr;
|
||||
if(!coli_cuda_tensor_upload(&g,hidden.data(),hs.data(),2,D,I,device)||
|
||||
!coli_cuda_tensor_upload(&u,hidden.data(),hs.data(),2,D,I,device)||
|
||||
!coli_cuda_tensor_upload(&d,down.data(),ds.data(),2,I,D,device))return 2;
|
||||
if(!coli_cuda_tensor_upload(&g,hidden.data(),hs.data(),2,D,I,device,0)||
|
||||
!coli_cuda_tensor_upload(&u,hidden.data(),hs.data(),2,D,I,device,0)||
|
||||
!coli_cuda_tensor_upload(&d,down.data(),ds.data(),2,I,D,device,0))return 2;
|
||||
for(int rows: {1,2,4,8}){
|
||||
double scalar=run(g,u,d,x.data(),a.data(),rows,3,0);
|
||||
double packed=run(g,u,d,x.data(),b.data(),rows,3,1);
|
||||
|
||||
@@ -99,7 +99,7 @@ int main(int argc, char **argv) {
|
||||
const int8_t q8b[8]={-1,-2,-3,-4, 1,-2,3,-4};
|
||||
const float s8b[2]={1.f,.5f},want8b[4]={10.f,15.f,-3.f,-2.5f};
|
||||
if(!coli_cuda_tensor_update(t8,q8b,s8b)||
|
||||
!coli_cuda_matmul(&t8,got,x,q8b,s8b,1,2,4,2,d0)||
|
||||
!coli_cuda_matmul(&t8,got,x,q8b,s8b,1,2,4,2,d0,0)||
|
||||
!close_enough(got,want8b,4))return 1;
|
||||
|
||||
/* Rows [-8,-1,0,7] and [1,2,3,4], packed low nibble first. */
|
||||
@@ -107,26 +107,26 @@ int main(int argc, char **argv) {
|
||||
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;
|
||||
if (!coli_cuda_matmul(&t4, got, x, q4, s4, 2, 1, 4, 2, d1, 0) || !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;
|
||||
if (!coli_cuda_matmul(&t2, got, x, q2, s2, 3, 1, 4, 2, d1, 0) || !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;
|
||||
if (!coli_cuda_matmul(&tf, got, x, wf, nullptr, 0, 1, 4, 2, d0, 0) || !close_enough(got, wantf, 2)) return 1;
|
||||
|
||||
const float eg[8] = {1,0,0,0, 0,1,0,0};
|
||||
const float eu[8] = {1,0,0,0, 0,1,0,0};
|
||||
const float ed[8] = {1,0, 0,1, 1,1, 1,-1};
|
||||
ColiCudaTensor *tg=nullptr,*tu=nullptr,*td=nullptr;
|
||||
if (!coli_cuda_tensor_upload(&tg,eg,nullptr,0,4,2,d0) ||
|
||||
!coli_cuda_tensor_upload(&tu,eu,nullptr,0,4,2,d0) ||
|
||||
!coli_cuda_tensor_upload(&td,ed,nullptr,0,2,4,d0)) return 1;
|
||||
if (!coli_cuda_tensor_upload(&tg,eg,nullptr,0,4,2,d0,0) ||
|
||||
!coli_cuda_tensor_upload(&tu,eu,nullptr,0,4,2,d0,0) ||
|
||||
!coli_cuda_tensor_upload(&td,ed,nullptr,0,2,4,d0,0)) return 1;
|
||||
float expert[8], want_expert[8];
|
||||
for(int s=0;s<2;s++){
|
||||
float a=x[s*4], b=x[s*4+1];
|
||||
@@ -144,7 +144,7 @@ int main(int argc, char **argv) {
|
||||
const float aw[16]={1,0,0,0, 0,1,0,0, 0,0,1,0, 0,0,0,1};
|
||||
const float aq[4]={1,2,.5f,-.5f},al[12]={1,0,0,0, 0,1,0,0, 0,0,1,0};
|
||||
const float ar[6]={1,0, 0,1, 1,1};float actx[2],aref[2];
|
||||
ColiCudaTensor *at=nullptr;if(!coli_cuda_tensor_upload(&at,aw,nullptr,0,4,4,d0))return 1;
|
||||
ColiCudaTensor *at=nullptr;if(!coli_cuda_tensor_upload(&at,aw,nullptr,0,4,4,d0,0))return 1;
|
||||
float score[3];for(int t=0;t<3;t++)score[t]=aq[0]*al[t*4]+aq[1]*al[t*4+1]+aq[2]*ar[t*2]+aq[3]*ar[t*2+1];
|
||||
float mx=score[0],z=0;for(int t=1;t<3;t++)mx=score[t]>mx?score[t]:mx;
|
||||
for(int t=0;t<3;t++){score[t]=std::exp(score[t]-mx);z+=score[t];}for(int t=0;t<3;t++)score[t]/=z;
|
||||
@@ -163,9 +163,9 @@ int main(int argc, char **argv) {
|
||||
for(int i=0;i<32;i++)ws4[i]=0.01f+(i%5)*0.002f;
|
||||
for(int i=0;i<64;i++)gx4[i]=std::sin((float)(i+1)*0.17f)*2.f;
|
||||
ColiCudaTensor *g4=nullptr,*u4=nullptr,*d4=nullptr;
|
||||
if(!coli_cuda_tensor_upload(&g4,w4,ws4,2,32,32,d0)||
|
||||
!coli_cuda_tensor_upload(&u4,w4,ws4,2,32,32,d0)||
|
||||
!coli_cuda_tensor_upload(&d4,w4,ws4,2,32,32,d0))return 1;
|
||||
if(!coli_cuda_tensor_upload(&g4,w4,ws4,2,32,32,d0,0)||
|
||||
!coli_cuda_tensor_upload(&u4,w4,ws4,2,32,32,d0,0)||
|
||||
!coli_cuda_tensor_upload(&d4,w4,ws4,2,32,32,d0,0))return 1;
|
||||
ColiCudaTensor *gg4[2]={g4,g4},*ug4[2]={u4,u4},*dg4[2]={d4,d4};
|
||||
if(!coli_cuda_expert_group(gg4,ug4,dg4,group_rows,2,scalar4,gx4))return 1;
|
||||
setenv("COLI_CUDA_TC_INT4","1",1);
|
||||
|
||||
@@ -119,7 +119,7 @@ int main(void){
|
||||
for(int i=0;i<O;i++) sc[i]=0.01f+0.001f*(i%7);
|
||||
for(size_t i=0;i<(size_t)S*K;i++) x[i]=rndf();
|
||||
ColiCudaTensor *t=NULL;
|
||||
ok&=coli_cuda_matmul(&t,ref,x,w4,sc,2,S,K,O,0); /* host path = reference */
|
||||
ok&=coli_cuda_matmul(&t,ref,x,w4,sc,2,S,K,O,0,0); /* host path = reference */
|
||||
float *xd=(float*)coli_cuda_pipe_alloc(0,(size_t)S*K*4);
|
||||
float *yd=(float*)coli_cuda_pipe_alloc(0,(size_t)S*O*4);
|
||||
ok&=coli_cuda_pipe_upload(0,xd,x,(size_t)S*K*4);
|
||||
|
||||
@@ -0,0 +1,704 @@
|
||||
#!/usr/bin/env python3
|
||||
"""
|
||||
diag_harness.py — Comprehensive model diagnostic harness for the colibri GLM-5.2 engine.
|
||||
|
||||
Runs a full campaign of tests against any model snapshot:
|
||||
Phase 0 (system): startup telemetry — GPU, RAM, cache cap, load time, MTP, idot kernel
|
||||
Phase 1 (smoke): correctness — curated prompts, coherence checks, corruption detection
|
||||
Phase 2 (diagnostic): deep x-ray — full PROFILE breakdown, routing, MTP, disk-split, CUDA tier
|
||||
Phase 3 (quality): benchmark accuracy — hellaswag/arc_challenge/mmlu via eval_glm.py SCORE
|
||||
Phase 4 (throughput): tok/s with and without MTP speculation
|
||||
Phase 5 (report): structured JSON + human-readable Markdown summary
|
||||
|
||||
Usage:
|
||||
python tools/diag_harness.py --snap /path/to/model --phase all
|
||||
python tools/diag_harness.py --snap /path/to/model --phase smoke --ngen 64
|
||||
python tools/diag_harness.py --snap /path/to/model --phase quality --quality-limit 200
|
||||
|
||||
Output goes to --out (default ./diag_results/<timestamp>/). Each phase writes a raw log
|
||||
(<phase>_<run>.log) and all metrics are collected into report.json + report.md.
|
||||
|
||||
Design notes:
|
||||
- stdout and stderr are captured separately via subprocess.PIPE. The engine streams
|
||||
generated text to stdout (interleaved with prompt + PROFILE stats), but TOKENS=1 dumps
|
||||
clean token-id lists to stderr — that is the primary text-capture path.
|
||||
- Every regex is anchored to the exact printf format strings in glm.c (verified against
|
||||
profile_print line 3853, run_text line 3948, the banner line 5299, etc.).
|
||||
- Subprocess calls have a hard timeout (default 600s) and are killed cleanly on expiry.
|
||||
- A single phase can be run standalone; results accumulate in the output dir.
|
||||
"""
|
||||
import os, sys, re, json, time, argparse, subprocess, signal, traceback
|
||||
from datetime import datetime
|
||||
from pathlib import Path
|
||||
|
||||
# ---------------------------------------------------------------------------
|
||||
# PROMPT SUITE — curated across categories. "expect" is a case-insensitive
|
||||
# substring checked against the generated text (None = coherence-only check).
|
||||
# ---------------------------------------------------------------------------
|
||||
PROMPTS = [
|
||||
{"id":"fact_capital", "cat":"factual", "prompt":"The capital of France is",
|
||||
"expect":"Paris", "note":"basic world knowledge"},
|
||||
{"id":"fact_boiling", "cat":"factual", "prompt":"What is the boiling point of water in Celsius?",
|
||||
"expect":"100", "note":"basic science"},
|
||||
{"id":"fact_planet", "cat":"factual", "prompt":"What planet is closest to the Sun?",
|
||||
"expect":"Mercury","note":"astronomy fact"},
|
||||
{"id":"math_mult", "cat":"math", "prompt":"What is 15 times 12?",
|
||||
"expect":"180", "note":"2-digit multiplication"},
|
||||
{"id":"math_add", "cat":"math", "prompt":"What is 847 plus 153?",
|
||||
"expect":"1000", "note":"3-digit addition"},
|
||||
{"id":"reason_train", "cat":"reasoning", "prompt":"If a train travels 60 mph for 2.5 hours, how far does it go?",
|
||||
"expect":"150", "note":"rate-time-distance"},
|
||||
{"id":"code_factorial","cat":"code", "prompt":"Write a Python function that computes the factorial of a number.",
|
||||
"expect":"def", "note":"code generation"},
|
||||
{"id":"explain_nn", "cat":"explanation","prompt":"Explain what a neural network is in one sentence.",
|
||||
"expect":None, "note":"coherence check"},
|
||||
{"id":"creative_story","cat":"creative", "prompt":"Write a one-sentence story about a lighthouse.",
|
||||
"expect":None, "note":"creative coherence"},
|
||||
{"id":"edge_hello", "cat":"edge", "prompt":"Hello, how are you today?",
|
||||
"expect":None, "note":"conversational opener"},
|
||||
{"id":"edge_the", "cat":"edge", "prompt":"The weather today is",
|
||||
"expect":None, "note":"simple continuation"},
|
||||
{"id":"edge_repeat", "cat":"edge", "prompt":"The quick brown fox jumps over the lazy dog. The quick brown fox",
|
||||
"expect":None, "note":"repetition-bait"},
|
||||
]
|
||||
|
||||
# ---------------------------------------------------------------------------
|
||||
# METRIC EXTRACTION — each parser is (name, compiled_regex, group_index, cast).
|
||||
# Regexes match the EXACT printf format strings in glm.c.
|
||||
# ---------------------------------------------------------------------------
|
||||
def _f1(g): return float(g)
|
||||
def _i1(g): return int(g)
|
||||
|
||||
# Build parsers as (regex, lambda(match)->value) for clarity
|
||||
RX = {
|
||||
# stdout: run_text summary line (line 3948)
|
||||
"decode_toks": (re.compile(r"decode (\d+) tokens in"), lambda m: int(m.group(1))),
|
||||
"decode_secs": (re.compile(r"decode \d+ tokens in ([\d.]+)s"), lambda m: float(m.group(1))),
|
||||
"decode_tps": (re.compile(r"decode \d+ tokens in [\d.]+s \(([\d.]+) tok/s\)"), lambda m: float(m.group(1))),
|
||||
"prefill_toks": (re.compile(r"prefill (\d+) tokens in"), lambda m: int(m.group(1))),
|
||||
"prefill_secs": (re.compile(r"prefill \d+ tokens in ([\d.]+)s"), lambda m: float(m.group(1))),
|
||||
"hit_rate": (re.compile(r"expert hit rate ([\d.]+)%"), lambda m: float(m.group(1))),
|
||||
"rss_gb": (re.compile(r"RSS ([\d.]+) GB"), lambda m: float(m.group(1))),
|
||||
"experts_per_tok":(re.compile(r"experts loaded/token: ([\d.]+)"), lambda m: float(m.group(1))),
|
||||
"mtp_accept": (re.compile(r"MTP acceptance ([\d.]+)%"), lambda m: float(m.group(1))),
|
||||
"mtp_acc_cnt": (re.compile(r"MTP acceptance \d+% \((\d+)/(\d+)\)"), lambda m: (int(m.group(1)), int(m.group(2)))),
|
||||
"spec_tok_per_fw":(re.compile(r"speculation: ([\d.]+) tokens/forward"),lambda m: float(m.group(1))),
|
||||
# stdout: PROFILE lines (profile_print, line 3853-3864)
|
||||
"prof_expert_disk": (re.compile(r"expert-disk ([\d.]+)s service"), lambda m: float(m.group(1))),
|
||||
"prof_expert_wait": (re.compile(r"service / ([\d.]+)s wait"), lambda m: float(m.group(1))),
|
||||
"prof_expert_mm": (re.compile(r"expert-matmul ([\d.]+)s"), lambda m: float(m.group(1))),
|
||||
"prof_attention": (re.compile(r"\| attention ([\d.]+)s"), lambda m: float(m.group(1))),
|
||||
"prof_kvb": (re.compile(r"including kvb ([\d.]+)s"), lambda m: float(m.group(1))),
|
||||
"prof_lm_head": (re.compile(r"lm_head ([\d.]+)s"), lambda m: float(m.group(1))),
|
||||
"prof_other": (re.compile(r"\| other ([\d.]+)s"), lambda m: float(m.group(1))),
|
||||
# stdout: banner (line 5299)
|
||||
"load_secs": (re.compile(r"loaded in ([\d.]+)s"), lambda m: float(m.group(1))),
|
||||
"resident_mb": (re.compile(r"resident dense: ([\d.]+) MB"), lambda m: float(m.group(1))),
|
||||
"mtp_status": (re.compile(r"MTP (ACTIVE|absent)"), lambda m: m.group(1)),
|
||||
"n_layers": (re.compile(r"layers=(\d+) experts=(\d+)"), lambda m: int(m.group(1))),
|
||||
"n_experts": (re.compile(r"layers=(\d+) experts=(\d+)"), lambda m: int(m.group(2))),
|
||||
# stdout: banner idot kernel
|
||||
"idot_kernel": (re.compile(r"idot: (\S+) =="), lambda m: m.group(1)),
|
||||
"cache_cap": (re.compile(r"cache=(\d+) experts/layer"), lambda m: int(m.group(1))),
|
||||
# stderr: RAM_GB (line 5086-5112)
|
||||
"ram_budget": (re.compile(r"\[RAM_GB=([\d.]+)"), lambda m: float(m.group(1))),
|
||||
"cap_lowered": (re.compile(r"cap lowered (\d+)->(\d+)"), lambda m: (int(m.group(1)), int(m.group(2)))),
|
||||
"cap_raised": (re.compile(r"cap raised (\d+)->(\d+)"), lambda m: (int(m.group(1)), int(m.group(2)))),
|
||||
"cap_ok": (re.compile(r"cap=(\d+) ok"), lambda m: int(m.group(1))),
|
||||
# stderr: CUDA (backend_cuda.cu:389)
|
||||
"cuda_device": (re.compile(r"\[CUDA\] device (\d+): (.*?), ([\d.]+) GB VRAM, sm_(\d)(\d)"),
|
||||
lambda m: {"id":int(m.group(1)),"name":m.group(2).strip(),"vram_gb":float(m.group(3)),
|
||||
"sm":f"{m.group(4)}.{m.group(5)}"}),
|
||||
"cuda_mode": (re.compile(r"\[CUDA\] mode: (.+)"), lambda m: m.group(1)),
|
||||
"cuda_tier": (re.compile(r"CUDA expert tier: (\d+) resident experts \(([\d.]+) GB\)"),
|
||||
lambda m: {"resident":int(m.group(1)),"vram_gb":float(m.group(2))}),
|
||||
# stderr: TOKENS dump (line 4010-4012)
|
||||
"tokens_dump": (re.compile(r"^\[TOKENS\] (\d+) generated:(.*)$", re.MULTILINE),
|
||||
lambda m: [int(x) for x in m.group(2).split()]),
|
||||
# stderr: per-16-token progress (emit_stream line 3742)
|
||||
"progress_tps": (re.compile(r"t=(\d+)\s+RSS ([\d.]+) GB\s+hit ([\d.]+)%\s+([\d.]+) tok/s\s+([\d.]+) tok/fw"),
|
||||
lambda m: {"tok":int(m.group(1)),"rss":float(m.group(2)),"hit":float(m.group(3)),
|
||||
"tps":float(m.group(4)),"tpf":float(m.group(5))}),
|
||||
# stderr: DSA, USAGE, KV startup lines
|
||||
"usage_loaded": (re.compile(r"\[USAGE\].*?(\d+) selections"), lambda m: int(m.group(1))),
|
||||
"kv_slots": (re.compile(r"\[KV\].*?(\d+) context slots"), lambda m: int(m.group(1))),
|
||||
}
|
||||
|
||||
def extract_metrics(stdout: str, stderr: str) -> dict:
|
||||
"""Parse all metrics from the engine's stdout+stderr output."""
|
||||
text_out = stdout or ""
|
||||
text_err = stderr or ""
|
||||
metrics = {}
|
||||
# For most parsers we search BOTH streams (engine is inconsistent about which
|
||||
# channel a given line lands on). Some are stream-specific (noted below).
|
||||
combined = text_out + "\n" + text_err
|
||||
for name, (rx, fn) in RX.items():
|
||||
m = rx.search(combined)
|
||||
if m:
|
||||
try: metrics[name] = fn(m)
|
||||
except (ValueError, IndexError): pass
|
||||
# TOKENS dump is stderr-only and may appear once; grab it explicitly
|
||||
m = RX["tokens_dump"][0].search(text_err)
|
||||
if m:
|
||||
try: metrics["tokens_dump"] = RX["tokens_dump"][1](m)
|
||||
except: pass
|
||||
# Multiple CUDA devices — collect all
|
||||
cuda_devs = []
|
||||
for m in re.finditer(r"\[CUDA\] device (\d+): (.*?), ([\d.]+) GB VRAM, sm_(\d)(\d)", text_err):
|
||||
cuda_devs.append({"id":int(m.group(1)),"name":m.group(2).strip(),
|
||||
"vram_gb":float(m.group(3)),"sm":f"{m.group(4)}.{m.group(5)}"})
|
||||
if cuda_devs: metrics["cuda_devices"] = cuda_devs
|
||||
# Multiple progress checkpoints — collect the full curve
|
||||
progress = []
|
||||
for m in RX["progress_tps"][0].finditer(text_err):
|
||||
try:
|
||||
progress.append(RX["progress_tps"][1](m))
|
||||
except: pass
|
||||
if progress: metrics["progress_curve"] = progress
|
||||
# Multiple PROFILE lines — there are two (prefill + decode). Keep both.
|
||||
prof_lines = [(i, line) for i, line in enumerate(text_out.splitlines()) if line.startswith("PROFILE:")]
|
||||
for idx, (line_no, line) in enumerate(prof_lines):
|
||||
label = "prefill_profile" if idx == 0 else "decode_profile"
|
||||
p = {}
|
||||
for pname, (rx, fn) in RX.items():
|
||||
if not pname.startswith("prof_"): continue
|
||||
mm = rx.search(line)
|
||||
if mm:
|
||||
try: p[pname] = fn(mm)
|
||||
except: pass
|
||||
if p: metrics[label] = p
|
||||
return metrics
|
||||
|
||||
|
||||
# ---------------------------------------------------------------------------
|
||||
# ENGINE RUNNER — subprocess wrapper with timeout, signal handling, logging.
|
||||
# ---------------------------------------------------------------------------
|
||||
class EngineRunner:
|
||||
def __init__(self, glm_path, snap, out_dir, default_env=None, timeout=600):
|
||||
self.glm = str(glm_path)
|
||||
self.snap = str(snap)
|
||||
self.out_dir = Path(out_dir)
|
||||
self.out_dir.mkdir(parents=True, exist_ok=True)
|
||||
self.default_env = default_env or {}
|
||||
self.timeout = timeout
|
||||
self.default_cap = 75 # production default (matches bench_full.sh)
|
||||
|
||||
def run_prompt(self, prompt, ngen=64, env_extra=None, log_name=None, cap=None):
|
||||
"""Run the engine in PROMPT mode and return (stdout, stderr, returncode, elapsed)."""
|
||||
env = dict(os.environ, SNAP=self.snap, PROMPT=prompt, NGEN=str(ngen))
|
||||
env.update(self.default_env)
|
||||
if env_extra: env.update(env_extra)
|
||||
env["TOKENS"] = "1" # always capture token ids for reliable text extraction
|
||||
cmd = [self.glm, str(cap if cap is not None else self.default_cap)]
|
||||
return self._exec(cmd, env, log_name)
|
||||
|
||||
def run_score(self, score_file, cap=None, env_extra=None, log_name=None):
|
||||
"""Run the engine in SCORE (log-likelihood) mode."""
|
||||
env = dict(os.environ, SNAP=self.snap, SCORE=score_file)
|
||||
env.update(self.default_env)
|
||||
if env_extra: env.update(env_extra)
|
||||
cmd = [self.glm, str(cap if cap is not None else self.default_cap)]
|
||||
return self._exec(cmd, env, log_name)
|
||||
|
||||
def _exec(self, cmd, env, log_name):
|
||||
t0 = time.time()
|
||||
log_path = self.out_dir / (log_name or f"run_{int(t0)}.log")
|
||||
try:
|
||||
proc = subprocess.Popen(cmd, env=env, stdout=subprocess.PIPE, stderr=subprocess.PIPE,
|
||||
text=True, encoding="utf-8", errors="replace")
|
||||
try:
|
||||
stdout, stderr = proc.communicate(timeout=self.timeout)
|
||||
rc = proc.returncode
|
||||
except subprocess.TimeoutExpired:
|
||||
proc.kill()
|
||||
stdout, stderr = proc.communicate()
|
||||
rc = -1
|
||||
stderr = (stderr or "") + f"\n[TIMEOUT after {self.timeout}s]\n"
|
||||
except Exception as e:
|
||||
stdout, stderr, rc = "", f"[EXCEPTION] {e}\n{traceback.format_exc()}", -2
|
||||
elapsed = time.time() - t0
|
||||
# Write raw log (both streams, clearly delimited)
|
||||
with open(log_path, "w", encoding="utf-8") as f:
|
||||
f.write(f"=== CMD: {' '.join(cmd)}\n=== ELAPSED: {elapsed:.1f}s\n=== RC: {rc}\n\n")
|
||||
f.write("--- STDOUT ---\n"); f.write(stdout or ""); f.write("\n")
|
||||
f.write("--- STDERR ---\n"); f.write(stderr or ""); f.write("\n")
|
||||
return stdout or "", stderr or "", rc, elapsed
|
||||
|
||||
|
||||
# ---------------------------------------------------------------------------
|
||||
# TEXT RECOVERY — decode the [TOKENS] ids back to text using the tokenizer.
|
||||
# Falls back to stdout extraction if tokenizer/tokens unavailable.
|
||||
# ---------------------------------------------------------------------------
|
||||
def recover_text(stdout, stderr, prompt, tokenizer=None):
|
||||
"""Extract generated text. Primary: TOKENS dump decoded. Fallback: stdout parse."""
|
||||
# Method 1: decode TOKENS ids via tokenizer
|
||||
if tokenizer:
|
||||
m = re.search(r"^\[TOKENS\] \d+ generated:(.*)$", stderr, re.MULTILINE)
|
||||
if m:
|
||||
ids = [int(x) for x in m.group(1).split()]
|
||||
if ids:
|
||||
try:
|
||||
return tokenizer.decode(ids), "tokens"
|
||||
except Exception: pass
|
||||
# Method 2: stdout — text is between the prompt string and "PROFILO" or "\n---"
|
||||
text = stdout or ""
|
||||
# Find the prompt in stdout, take everything after it up to PROFILO
|
||||
idx = text.find(prompt)
|
||||
if idx >= 0:
|
||||
after = text[idx + len(prompt):]
|
||||
cut = after.find("PROFILO")
|
||||
if cut >= 0:
|
||||
return after[:cut].strip(), "stdout"
|
||||
cut = after.find("\n---")
|
||||
if cut >= 0:
|
||||
return after[:cut].strip(), "stdout"
|
||||
return "", "none"
|
||||
|
||||
|
||||
# ---------------------------------------------------------------------------
|
||||
# CORRUPTION / COHERENCE CHECKS
|
||||
# ---------------------------------------------------------------------------
|
||||
def check_repetition(token_ids):
|
||||
"""Detect degenerate repetition: same token 3+ consecutive times."""
|
||||
if not token_ids or len(token_ids) < 6:
|
||||
return False, 0
|
||||
max_run = 1; cur = 1
|
||||
for i in range(1, len(token_ids)):
|
||||
if token_ids[i] == token_ids[i-1]: cur += 1; max_run = max(max_run, cur)
|
||||
else: cur = 1
|
||||
return max_run >= 3, max_run
|
||||
|
||||
def check_expected(text, expect):
|
||||
"""Check if expected substring appears case-insensitively in the first 200 chars."""
|
||||
if not expect or not text: return None
|
||||
return expect.lower() in text[:200].lower()
|
||||
|
||||
|
||||
# ---------------------------------------------------------------------------
|
||||
# PHASES
|
||||
# ---------------------------------------------------------------------------
|
||||
class DiagnosticHarness:
|
||||
def __init__(self, args):
|
||||
self.args = args
|
||||
self.snap = args.snap
|
||||
self.glm = args.glm
|
||||
ts = datetime.now().strftime("%Y%m%d_%H%M%S")
|
||||
self.out_dir = Path(args.out or f"./diag_results/{ts}")
|
||||
self.out_dir.mkdir(parents=True, exist_ok=True)
|
||||
# Base env for all runs
|
||||
base_env = {"TEMP": "0", "PIPE": "1", "PIPE_WORKERS": "8", "DIRECT": "1"}
|
||||
if args.ram: base_env["RAM_GB"] = str(args.ram)
|
||||
if args.cuda:
|
||||
base_env.update({
|
||||
"COLI_CUDA": "1",
|
||||
"COLI_GPU": str(args.gpu) if args.gpu is not None else "0",
|
||||
"CUDA_DENSE": "1",
|
||||
"COLI_CUDA_ATTN": "1",
|
||||
"COLI_CUDA_PIPE": "2",
|
||||
"COLI_CUDA_PIPE_S_MIN": "1",
|
||||
})
|
||||
self.runner = EngineRunner(self.glm, self.snap, self.out_dir, base_env, timeout=args.timeout)
|
||||
self.runner.default_cap = args.cap
|
||||
# Load tokenizer for text recovery
|
||||
self.tokenizer = None
|
||||
tok_path = os.path.join(self.snap, "tokenizer.json")
|
||||
if os.path.exists(tok_path):
|
||||
try:
|
||||
from tokenizers import Tokenizer
|
||||
self.tokenizer = Tokenizer.from_file(tok_path)
|
||||
except Exception as e:
|
||||
print(f"[warn] could not load tokenizer: {e}", file=sys.stderr)
|
||||
self.results = {"meta": {"snap": self.snap, "glm": self.glm,
|
||||
"timestamp": ts, "args": vars(args)},
|
||||
"phases": {}}
|
||||
|
||||
def phase_system(self):
|
||||
"""Phase 0: minimal run to capture startup telemetry."""
|
||||
print("\n" + "="*60 + "\nPHASE 0: SYSTEM PROBE\n" + "="*60)
|
||||
stdout, stderr, rc, elapsed = self.runner.run_prompt(
|
||||
"Hello", ngen=1, log_name="system_probe.log")
|
||||
metrics = extract_metrics(stdout, stderr)
|
||||
result = {
|
||||
"rc": rc, "elapsed": elapsed,
|
||||
"load_secs": metrics.get("load_secs"),
|
||||
"resident_mb": metrics.get("resident_mb"),
|
||||
"n_layers": metrics.get("n_layers"),
|
||||
"n_experts": metrics.get("n_experts"),
|
||||
"mtp_status": metrics.get("mtp_status"),
|
||||
"idot_kernel": metrics.get("idot_kernel"),
|
||||
"cache_cap_requested": metrics.get("cache_cap"),
|
||||
"ram_budget_gb": metrics.get("ram_budget"),
|
||||
"cap_lowered": metrics.get("cap_lowered"),
|
||||
"cap_raised": metrics.get("cap_raised"),
|
||||
"cap_final": metrics.get("cap_ok"),
|
||||
"cuda_devices": metrics.get("cuda_devices", []),
|
||||
"cuda_mode": metrics.get("cuda_mode"),
|
||||
}
|
||||
# Print summary
|
||||
print(f" load time: {result['load_secs']:.2f}s" if result['load_secs'] else " load time: FAILED")
|
||||
print(f" resident: {result['resident_mb']:.0f} MB" if result['resident_mb'] else "")
|
||||
print(f" layers/exp: {result['n_layers']}/{result['n_experts']}" if result['n_layers'] else "")
|
||||
print(f" MTP: {result['mtp_status']}")
|
||||
print(f" idot kernel: {result['idot_kernel']}")
|
||||
print(f" RAM budget: {result['ram_budget_gb']} GB" if result['ram_budget_gb'] else " RAM budget: auto")
|
||||
if result['cap_lowered']:
|
||||
print(f" cache cap: {result['cap_lowered'][0]} -> {result['cap_lowered'][1]} (RAM-lowered)")
|
||||
elif result['cap_final']:
|
||||
print(f" cache cap: {result['cap_final']} (ok)")
|
||||
else:
|
||||
print(f" cache cap: {result['cache_cap_requested']}")
|
||||
for d in result['cuda_devices']:
|
||||
print(f" GPU {d['id']}: {d['name']}, {d['vram_gb']:.1f} GB, sm_{d['sm']}")
|
||||
if rc != 0 and rc != -1:
|
||||
print(f" WARNING: engine returned rc={rc}")
|
||||
self.results["phases"]["system"] = result
|
||||
return result
|
||||
|
||||
def phase_smoke(self):
|
||||
"""Phase 1: correctness smoke test across curated prompts."""
|
||||
print("\n" + "="*60 + "\nPHASE 1: CORRECTNESS SMOKE TEST\n" + "="*60)
|
||||
ngen = self.args.ngen
|
||||
prompt_results = []
|
||||
pass_count = 0; total = len(PROMPTS)
|
||||
for p in PROMPTS:
|
||||
stdout, stderr, rc, elapsed = self.runner.run_prompt(
|
||||
p["prompt"], ngen=ngen, log_name=f"smoke_{p['id']}.log")
|
||||
metrics = extract_metrics(stdout, stderr)
|
||||
token_ids = metrics.get("tokens_dump", [])
|
||||
text, method = recover_text(stdout, stderr, p["prompt"], self.tokenizer)
|
||||
is_rep, max_run = check_repetition(token_ids)
|
||||
expect_ok = check_expected(text, p.get("expect"))
|
||||
# Verdict: pass if text is non-empty, no severe repetition, and (if factual) expected found
|
||||
has_text = bool(text.strip())
|
||||
ok = has_text and not is_rep
|
||||
if p.get("expect") and expect_ok is False: ok = False
|
||||
if ok: pass_count += 1
|
||||
# Diagnose failure mode for reporting
|
||||
if not has_text:
|
||||
fail_reason = "no tokens generated (immediate EOS)"
|
||||
elif is_rep:
|
||||
fail_reason = f"repetition loop (max run={max_run})"
|
||||
elif p.get("expect") and expect_ok is False:
|
||||
fail_reason = f"expected '{p['expect']}' not found"
|
||||
else:
|
||||
fail_reason = ""
|
||||
prompt_results.append({
|
||||
"id": p["id"], "cat": p["cat"], "prompt": p["prompt"],
|
||||
"expect": p.get("expect"), "generated": text[:300],
|
||||
"expect_match": expect_ok, "repetition": is_rep, "max_run": max_run,
|
||||
"toks_generated": len(token_ids), "decode_tps": metrics.get("decode_tps"),
|
||||
"hit_rate": metrics.get("hit_rate"), "rc": rc, "elapsed": elapsed,
|
||||
"text_method": method, "pass": ok, "fail_reason": fail_reason,
|
||||
})
|
||||
status = "PASS" if ok else "FAIL"
|
||||
extra = f" expect={'Y' if expect_ok else ('N' if expect_ok is False else '-')}" if p.get("expect") else ""
|
||||
tps = f" {metrics.get('decode_tps',0):.2f}t/s" if metrics.get("decode_tps") else ""
|
||||
print(f" [{status}] {p['id']:<22} ({p['cat']:<11}) rep={'Y' if is_rep else 'N'}{extra}{tps}")
|
||||
if text:
|
||||
preview = text.replace("\n", " ")[:80]
|
||||
print(f" -> {preview}")
|
||||
elif rc != 0:
|
||||
print(f" -> [engine rc={rc}]")
|
||||
else:
|
||||
print(f" -> [{fail_reason}]")
|
||||
result = {"prompts": prompt_results, "pass": pass_count, "total": total,
|
||||
"pass_rate": 100.0 * pass_count / total if total else 0}
|
||||
print(f"\n SMOKE SUMMARY: {pass_count}/{total} passed ({result['pass_rate']:.0f}%)")
|
||||
self.results["phases"]["smoke"] = result
|
||||
return result
|
||||
|
||||
def phase_diagnostic(self):
|
||||
"""Phase 2: single deep-instrumented run — full PROFILE + routing + MTP."""
|
||||
print("\n" + "="*60 + "\nPHASE 2: FULL SYSTEM DIAGNOSTIC\n" + "="*60)
|
||||
diag_env = {
|
||||
"LOOKA": "1", "DISK_SPLIT": "1", "ROUTE_AGREE": "1",
|
||||
}
|
||||
if self.args.cuda:
|
||||
diag_env["COLI_CUDA_PROFILE"] = "1"
|
||||
prompt = ("Write a short paragraph explaining how photosynthesis works. "
|
||||
"Include the roles of sunlight, water, and carbon dioxide.")
|
||||
stdout, stderr, rc, elapsed = self.runner.run_prompt(
|
||||
prompt, ngen=self.args.ngen, env_extra=diag_env, log_name="diagnostic.log")
|
||||
metrics = extract_metrics(stdout, stderr)
|
||||
text, _ = recover_text(stdout, stderr, prompt, self.tokenizer)
|
||||
result = {
|
||||
"rc": rc, "elapsed": elapsed,
|
||||
"generated_text": text[:500],
|
||||
"prefill_profile": metrics.get("prefill_profile", {}),
|
||||
"decode_profile": metrics.get("decode_profile", {}),
|
||||
"decode_tps": metrics.get("decode_tps"),
|
||||
"prefill_secs": metrics.get("prefill_secs"),
|
||||
"hit_rate": metrics.get("hit_rate"),
|
||||
"rss_gb": metrics.get("rss_gb"),
|
||||
"experts_per_tok": metrics.get("experts_per_tok"),
|
||||
"mtp_accept": metrics.get("mtp_accept"),
|
||||
"mtp_counts": metrics.get("mtp_acc_cnt"),
|
||||
"spec_tok_per_fw": metrics.get("spec_tok_per_fw"),
|
||||
"cuda_tier": metrics.get("cuda_tier"),
|
||||
"progress_curve": metrics.get("progress_curve", []),
|
||||
}
|
||||
# Print the PROFILE breakdown
|
||||
dp = result["decode_profile"]
|
||||
print(f"\n DECODE TIMING BREAKDOWN (per-bucket, seconds):")
|
||||
print(f" expert-disk: {dp.get('prof_expert_disk', '?'):>8}")
|
||||
print(f" expert-matmul: {dp.get('prof_expert_mm', '?'):>8}")
|
||||
print(f" attention: {dp.get('prof_attention', '?'):>8}")
|
||||
print(f" lm_head: {dp.get('prof_lm_head', '?'):>8}")
|
||||
print(f" other: {dp.get('prof_other', '?'):>8}")
|
||||
print(f"\n PERFORMANCE:")
|
||||
print(f" prefill: {result['prefill_secs']:.2f}s" if result['prefill_secs'] else " prefill: ?")
|
||||
print(f" decode: {result['decode_tps']:.3f} tok/s" if result['decode_tps'] else " decode: ?")
|
||||
print(f" hit rate: {result['hit_rate']:.1f}%" if result.get('hit_rate') is not None else " hit rate: ?")
|
||||
print(f" RSS: {result['rss_gb']:.1f} GB" if result.get('rss_gb') else " RSS: ?")
|
||||
print(f" exp/tok: {result['experts_per_tok']:.1f}" if result.get('experts_per_tok') else " exp/tok: ?")
|
||||
print(f" MTP: {result['mtp_accept']:.0f}% accept" if result.get('mtp_accept') is not None else " MTP: ?")
|
||||
if result.get('cuda_tier'):
|
||||
ct = result['cuda_tier']
|
||||
print(f" CUDA tier: {ct['resident']} experts, {ct['vram_gb']:.1f} GB VRAM")
|
||||
# Show generated text preview
|
||||
if text:
|
||||
print(f"\n GENERATED TEXT (first 200 chars):")
|
||||
print(f" {text[:200].replace(chr(10), ' ')}")
|
||||
else:
|
||||
print(f"\n GENERATED TEXT: [none recovered]")
|
||||
self.results["phases"]["diagnostic"] = result
|
||||
return result
|
||||
|
||||
def phase_quality(self):
|
||||
"""Phase 3: benchmark accuracy via eval_glm.py SCORE mode."""
|
||||
print("\n" + "="*60 + "\nPHASE 3: QUALITY BENCHMARKS\n" + "="*60)
|
||||
eval_script = os.path.join(os.path.dirname(__file__), "eval_glm.py")
|
||||
bench_dir = os.path.join(os.path.dirname(os.path.dirname(__file__)), "bench")
|
||||
if not os.path.exists(eval_script):
|
||||
print(f" [SKIP] eval_glm.py not found at {eval_script}")
|
||||
self.results["phases"]["quality"] = {"error": "eval_glm.py not found"}
|
||||
return None
|
||||
tasks = ["hellaswag", "arc_challenge", "mmlu"]
|
||||
missing = [t for t in tasks if not os.path.exists(os.path.join(bench_dir, f"{t}.jsonl"))]
|
||||
if missing:
|
||||
print(f" [SKIP] benchmark data missing: {missing}")
|
||||
print(f" run: python tools/fetch_benchmarks.py --out {bench_dir}")
|
||||
self.results["phases"]["quality"] = {"error": f"missing benchmark data: {missing}"}
|
||||
return None
|
||||
limit = self.args.quality_limit
|
||||
py = sys.executable
|
||||
cmd = [py, eval_script, "--snap", self.snap, "--glm", self.glm,
|
||||
"--data", bench_dir, "--tasks", ",".join(tasks), "--limit", str(limit)]
|
||||
env = dict(os.environ)
|
||||
if self.args.ram: env["RAM_GB"] = str(self.args.ram)
|
||||
print(f" Running eval_glm.py (tasks={tasks}, n={limit})...")
|
||||
print(f" This takes ~{limit*3*4/0.05:.0f}s at 0.05 tok/s (worst case)...")
|
||||
t0 = time.time()
|
||||
log_path = self.out_dir / "quality_eval.log"
|
||||
try:
|
||||
with open(log_path, "w", encoding="utf-8") as logf:
|
||||
proc = subprocess.run(cmd, env=env, capture_output=True, text=True, timeout=self.args.timeout*3,
|
||||
encoding="utf-8", errors="replace")
|
||||
logf.write(proc.stdout); logf.write("\n--- STDERR ---\n"); logf.write(proc.stderr)
|
||||
elapsed = time.time() - t0
|
||||
# Parse the acc/acc_norm table from eval_glm.py output
|
||||
scores = {}
|
||||
for line in proc.stdout.splitlines():
|
||||
# lines like: "hellaswag 40 45.0% 50.0%"
|
||||
m = re.match(r"(\w+)\s+(\d+)\s+([\d.]+)%\s+([\d.]+)%", line.strip())
|
||||
if m:
|
||||
scores[m.group(1)] = {"n": int(m.group(2)), "acc": float(m.group(3)),
|
||||
"acc_norm": float(m.group(4))}
|
||||
mean_m = re.search(r"MEAN acc_norm:\s*([\d.]+)%", proc.stdout)
|
||||
result = {"scores": scores, "mean_acc_norm": float(mean_m.group(1)) if mean_m else None,
|
||||
"elapsed": elapsed, "rc": proc.returncode}
|
||||
print(f" Completed in {elapsed:.0f}s\n")
|
||||
print(f" {'task':<18} {'n':>4} {'acc':>7} {'acc_norm':>9}")
|
||||
for t, s in scores.items():
|
||||
print(f" {t:<18} {s['n']:>4} {s['acc']:>6.1f}% {s['acc_norm']:>8.1f}%")
|
||||
if result["mean_acc_norm"] is not None:
|
||||
print(f"\n MEAN acc_norm: {result['mean_acc_norm']:.1f}%")
|
||||
except subprocess.TimeoutExpired:
|
||||
result = {"error": f"eval timed out after {self.args.timeout*3}s"}
|
||||
print(f" [TIMEOUT] eval_glm.py exceeded {self.args.timeout*3}s")
|
||||
except Exception as e:
|
||||
result = {"error": str(e)}
|
||||
print(f" [ERROR] {e}")
|
||||
self.results["phases"]["quality"] = result
|
||||
return result
|
||||
|
||||
def phase_throughput(self):
|
||||
"""Phase 4: tok/s comparison — MTP on vs off."""
|
||||
print("\n" + "="*60 + "\nPHASE 4: THROUGHPUT BENCHMARK\n" + "="*60)
|
||||
prompt = "Summarize the plot of Romeo and Juliet in three sentences."
|
||||
ngen = self.args.ngen
|
||||
results = {}
|
||||
# Run 1: with MTP (default)
|
||||
print(f" [1/2] MTP ON (draft=3)...")
|
||||
stdout, stderr, rc, t = self.runner.run_prompt(
|
||||
prompt, ngen=ngen, log_name="throughput_mtp_on.log")
|
||||
m_on = extract_metrics(stdout, stderr)
|
||||
results["mtp_on"] = {"tps": m_on.get("decode_tps"), "hit": m_on.get("hit_rate"),
|
||||
"mtp_accept": m_on.get("mtp_accept"), "elapsed": t}
|
||||
print(f" {m_on.get('decode_tps',0):.3f} tok/s | hit {m_on.get('hit_rate',0):.1f}% | "
|
||||
f"MTP {m_on.get('mtp_accept',0):.0f}%")
|
||||
# Run 2: MTP off
|
||||
print(f" [2/2] MTP OFF (MTP=0)...")
|
||||
stdout, stderr, rc, t = self.runner.run_prompt(
|
||||
prompt, ngen=ngen, env_extra={"MTP": "0"}, log_name="throughput_mtp_off.log")
|
||||
m_off = extract_metrics(stdout, stderr)
|
||||
results["mtp_off"] = {"tps": m_off.get("decode_tps"), "hit": m_off.get("hit_rate"),
|
||||
"elapsed": t}
|
||||
print(f" {m_off.get('decode_tps',0):.3f} tok/s | hit {m_off.get('hit_rate',0):.1f}%")
|
||||
# Compute speedup
|
||||
if results["mtp_on"]["tps"] and results["mtp_off"]["tps"] and results["mtp_off"]["tps"] > 0:
|
||||
sp = results["mtp_on"]["tps"] / results["mtp_off"]["tps"]
|
||||
results["mtp_speedup"] = sp
|
||||
print(f"\n MTP speedup: {sp:.2f}x ({'MTP helps' if sp > 1.05 else 'MTP hurts' if sp < 0.95 else 'no effect'})")
|
||||
else:
|
||||
print(f"\n MTP speedup: (insufficient data)")
|
||||
self.results["phases"]["throughput"] = results
|
||||
return results
|
||||
|
||||
def write_report(self):
|
||||
"""Write report.json and report.md."""
|
||||
ts = self.results["meta"]["timestamp"]
|
||||
json_path = self.out_dir / "report.json"
|
||||
md_path = self.out_dir / "report.md"
|
||||
# JSON
|
||||
with open(json_path, "w", encoding="utf-8") as f:
|
||||
json.dump(self.results, f, indent=2, default=str)
|
||||
# Markdown
|
||||
lines = [
|
||||
f"# Diagnostic Report — {ts}",
|
||||
f"",
|
||||
f"**Model:** `{self.snap}`",
|
||||
f"**Engine:** `{self.glm}`",
|
||||
f"**Date:** {datetime.now().strftime('%Y-%m-%d %H:%M:%S')}",
|
||||
f"",
|
||||
]
|
||||
phases = self.results["phases"]
|
||||
# System
|
||||
if "system" in phases:
|
||||
s = phases["system"]
|
||||
lines += ["## Phase 0: System Probe", ""]
|
||||
lines += [f"- Load time: **{s.get('load_secs','?')}s**"]
|
||||
lines += [f"- Layers/experts: {s.get('n_layers','?')}/{s.get('n_experts','?')}"]
|
||||
lines += [f"- MTP: {s.get('mtp_status','?')}"]
|
||||
lines += [f"- idot kernel: `{s.get('idot_kernel','?')}`"]
|
||||
lines += [f"- RAM budget: {s.get('ram_budget_gb','auto')} GB"]
|
||||
if s.get("cap_lowered"):
|
||||
lines += [f"- Cache cap: {s['cap_lowered'][0]}→{s['cap_lowered'][1]} (RAM-lowered)"]
|
||||
elif s.get("cap_final"):
|
||||
lines += [f"- Cache cap: {s['cap_final']}"]
|
||||
for d in s.get("cuda_devices", []):
|
||||
lines += [f"- GPU {d['id']}: {d['name']}, {d['vram_gb']:.1f} GB, sm_{d['sm']}"]
|
||||
lines.append("")
|
||||
# Smoke
|
||||
if "smoke" in phases:
|
||||
sm = phases["smoke"]
|
||||
lines += [f"## Phase 1: Correctness Smoke", "",
|
||||
f"**{sm['pass']}/{sm['total']} prompts passed ({sm['pass_rate']:.0f}%)**", "",
|
||||
"| ID | Category | Pass | Expect | Repetition | tok/s | Generated (first 80 chars) |",
|
||||
"|---|---|---|---|---|---|---|"]
|
||||
for p in sm["prompts"]:
|
||||
gen = p["generated"][:80].replace("|", "\\|").replace("\n", " ") if p["generated"] else ""
|
||||
exp = "Y" if p.get("expect_match") else ("N" if p.get("expect_match") is False else "-")
|
||||
tps = f"{p.get('decode_tps',0):.2f}" if p.get("decode_tps") else "-"
|
||||
lines.append(f"| {p['id']} | {p['cat']} | {'✅' if p['pass'] else '❌'} | {exp} | "
|
||||
f"{'⚠️' if p['repetition'] else '-'} (run={p.get('max_run',0)}) | {tps} | {gen} |")
|
||||
lines.append("")
|
||||
# Diagnostic
|
||||
if "diagnostic" in phases:
|
||||
d = phases["diagnostic"]
|
||||
lines += ["## Phase 2: Full System Diagnostic", ""]
|
||||
dp = d.get("decode_profile", {})
|
||||
lines += ["### Decode Timing Breakdown", "",
|
||||
"| Bucket | Seconds |", "|---|---|"]
|
||||
for k, label in [("prof_expert_disk","expert-disk"),("prof_expert_mm","expert-matmul"),
|
||||
("prof_attention","attention"),("prof_lm_head","lm_head"),
|
||||
("prof_other","other")]:
|
||||
v = dp.get(k, "?")
|
||||
lines.append(f"| {label} | {v} |")
|
||||
lines += [f"", f"### Performance", f"- Decode: **{d.get('decode_tps','?')} tok/s**",
|
||||
f"- Prefill: {d.get('prefill_secs','?')}s",
|
||||
f"- Expert hit rate: {d.get('hit_rate','?')}%",
|
||||
f"- RSS: {d.get('rss_gb','?')} GB",
|
||||
f"- Experts/token: {d.get('experts_per_tok','?')}",
|
||||
f"- MTP acceptance: {d.get('mtp_accept','?')}%", ""]
|
||||
if d.get("generated_text"):
|
||||
lines += [f"### Generated Text", f"```\n{d['generated_text'][:300]}\n```", ""]
|
||||
# Quality
|
||||
if "quality" in phases:
|
||||
q = phases["quality"]
|
||||
lines += ["## Phase 3: Quality Benchmarks", ""]
|
||||
if "error" in q:
|
||||
lines += [f"⚠️ {q['error']}", ""]
|
||||
else:
|
||||
lines += [f"**Mean acc_norm: {q.get('mean_acc_norm','?')}%**", "",
|
||||
"| Task | n | acc | acc_norm |", "|---|---|---|---|"]
|
||||
for t, s in q.get("scores", {}).items():
|
||||
lines.append(f"| {t} | {s['n']} | {s['acc']:.1f}% | {s['acc_norm']:.1f}% |")
|
||||
lines.append("")
|
||||
# Throughput
|
||||
if "throughput" in phases:
|
||||
th = phases["throughput"]
|
||||
lines += ["## Phase 4: Throughput", "",
|
||||
"| Mode | tok/s | hit% | MTP accept |", "|---|---|---|---|"]
|
||||
on = th.get("mtp_on", {}); off = th.get("mtp_off", {})
|
||||
lines.append(f"| MTP ON | {on.get('tps','?')} | {on.get('hit','?')}% | {on.get('mtp_accept','?')}% |")
|
||||
lines.append(f"| MTP OFF | {off.get('tps','?')} | {off.get('hit','?')}% | — |")
|
||||
if th.get("mtp_speedup"):
|
||||
lines.append(f"\n**MTP speedup: {th['mtp_speedup']:.2f}x**")
|
||||
lines.append("")
|
||||
with open(md_path, "w", encoding="utf-8") as f:
|
||||
f.write("\n".join(lines))
|
||||
print(f"\n{'='*60}")
|
||||
print(f"Report written:")
|
||||
print(f" JSON: {json_path}")
|
||||
print(f" MD: {md_path}")
|
||||
print(f" Logs: {self.out_dir}/*.log")
|
||||
print(f"{'='*60}")
|
||||
|
||||
def run(self):
|
||||
phase = self.args.phase
|
||||
if phase == "all":
|
||||
self.phase_system()
|
||||
self.phase_smoke()
|
||||
self.phase_diagnostic()
|
||||
self.phase_quality()
|
||||
self.phase_throughput()
|
||||
elif phase == "system": self.phase_system()
|
||||
elif phase == "smoke": self.phase_smoke()
|
||||
elif phase == "diagnostic": self.phase_diagnostic()
|
||||
elif phase == "quality": self.phase_quality()
|
||||
elif phase == "throughput": self.phase_throughput()
|
||||
else:
|
||||
print(f"Unknown phase: {phase}", file=sys.stderr); sys.exit(1)
|
||||
self.write_report()
|
||||
|
||||
|
||||
def main():
|
||||
ap = argparse.ArgumentParser(description="Comprehensive model diagnostic harness for colibri GLM-5.2")
|
||||
ap.add_argument("--snap", required=True, help="model snapshot directory")
|
||||
ap.add_argument("--glm", default=None, help="engine binary path (default: ./glm.exe or ./glm)")
|
||||
ap.add_argument("--phase", default="all",
|
||||
choices=["all","system","smoke","diagnostic","quality","throughput"],
|
||||
help="which test phase to run")
|
||||
ap.add_argument("--out", default=None, help="output directory (default: ./diag_results/<timestamp>)")
|
||||
ap.add_argument("--ngen", type=int, default=64, help="generation length for smoke/throughput (default 64)")
|
||||
ap.add_argument("--quality-limit", type=int, default=40, help="questions per benchmark task (default 40)")
|
||||
ap.add_argument("--ram", type=float, default=0, help="RAM_GB override (0=auto)")
|
||||
ap.add_argument("--cuda", action="store_true", help="enable COLI_CUDA GPU tier")
|
||||
ap.add_argument("--gpu", type=int, default=None, help="GPU device ordinal (with --cuda)")
|
||||
ap.add_argument("--cap", type=int, default=75, help="experts-per-layer cache cap (default 75)")
|
||||
ap.add_argument("--timeout", type=int, default=600, help="per-run timeout in seconds (default 600)")
|
||||
a = ap.parse_args()
|
||||
# Auto-detect engine binary
|
||||
if a.glm is None:
|
||||
for cand in ["./glm.exe", "./glm", "../glm.exe", os.path.join(os.path.dirname(__file__), "..", "glm.exe")]:
|
||||
if os.path.exists(cand): a.glm = os.path.abspath(cand); break
|
||||
if a.glm is None:
|
||||
print("ERROR: could not find glm/glm.exe. Specify with --glm", file=sys.stderr); sys.exit(1)
|
||||
if not os.path.isdir(a.snap):
|
||||
print(f"ERROR: snapshot dir does not exist: {a.snap}", file=sys.stderr); sys.exit(1)
|
||||
harness = DiagnosticHarness(a)
|
||||
harness.run()
|
||||
|
||||
if __name__ == "__main__":
|
||||
main()
|
||||
+57
-9
@@ -22,7 +22,7 @@ USO:
|
||||
# leve di ricerca: passate al motore via env
|
||||
TOPP=0.9 python3 tools/eval_glm.py --snap /home/vincenzo/glm52_i4 --data ./bench --tasks mmlu --ram 15
|
||||
"""
|
||||
import os, sys, subprocess, argparse, random, json, tempfile, time
|
||||
import os, sys, subprocess, argparse, random, json, tempfile, time, threading
|
||||
|
||||
# mini-set OFFLINE per testare la meccanica (NON misura qualita': domande banali)
|
||||
SMOKE = [
|
||||
@@ -114,6 +114,7 @@ def main():
|
||||
ap.add_argument("--seed", type=int, default=1234)
|
||||
ap.add_argument("--dry", action="store_true", help="build requests and stop without running the engine")
|
||||
ap.add_argument("--selftest", action="store_true", help="verify the scoring calculations")
|
||||
ap.add_argument("--out", default="", help="write incremental results CSV here (one row per request, flushed as it lands)")
|
||||
a = ap.parse_args()
|
||||
|
||||
if a.selftest: # acc/acc_norm con logprob sintetici
|
||||
@@ -143,15 +144,62 @@ def main():
|
||||
if a.ram: env["RAM_GB"] = str(a.ram)
|
||||
cmd = [a.glm, str(a.cap)] + a.bits.split()
|
||||
print("running:", " ".join(cmd), file=sys.stderr)
|
||||
|
||||
# Stream results line-by-line so a crash at request N keeps 1..N-1 and shows
|
||||
# exactly where it stopped. The engine prints "<lp> <contlen> <greedy>" per
|
||||
# request to stdout and "[score N req | ...]" progress to stderr; buffering
|
||||
# both until exit (the old subprocess.run) wastes the whole run on a crash.
|
||||
out_f = open(a.out, "a") if a.out else None
|
||||
if out_f:
|
||||
out_f.write(f"# eval_glm snap={a.snap} tasks={a.tasks} limit={a.limit} seed={a.seed} started={time.strftime('%Y-%m-%dT%H:%M:%S')}\n")
|
||||
out_f.write("req_idx,task,qi,oi,contlen,contchars,gold,logprob,greedy\n")
|
||||
out_f.flush()
|
||||
t0 = time.time()
|
||||
proc = subprocess.run(cmd, env=env, capture_output=True, text=True)
|
||||
if proc.returncode != 0:
|
||||
print("ENGINE ERROR:\n", proc.stderr[-2000:], file=sys.stderr); sys.exit(1)
|
||||
lines = [l for l in proc.stdout.strip().splitlines() if l and l[0] in "-0123456789"]
|
||||
if len(lines) != len(reqs):
|
||||
print(f"WARNING: {len(lines)} outputs for {len(reqs)} requests", file=sys.stderr)
|
||||
lp = [float(l.split()[0]) for l in lines]
|
||||
print(f"(engine: {time.time()-t0:.0f}s){proc.stderr.strip().splitlines()[-1] if proc.stderr.strip() else ''}", file=sys.stderr)
|
||||
proc = subprocess.Popen(cmd, env=env, stdout=subprocess.PIPE, stderr=subprocess.PIPE,
|
||||
text=True, bufsize=1) # line-buffered
|
||||
lp = [None] * len(reqs)
|
||||
n_done = 0
|
||||
# Drain stderr (engine progress lines) to console live on a background thread
|
||||
# so the [score N req] heartbeat is visible while stdout is consumed below.
|
||||
def _drain_stderr():
|
||||
for line in proc.stderr:
|
||||
print(f" [engine] {line.rstrip()}", file=sys.stderr)
|
||||
threading.Thread(target=_drain_stderr, daemon=True).start()
|
||||
for line in proc.stdout:
|
||||
line = line.strip()
|
||||
if not line or line[0] not in "-0123456789": continue
|
||||
parts = line.split()
|
||||
if n_done >= len(reqs): break
|
||||
try: logprob = float(parts[0])
|
||||
except (ValueError, IndexError): continue
|
||||
lp[n_done] = logprob
|
||||
greedy = parts[2] if len(parts) > 2 else "?"
|
||||
t, qi, oi, clen, cchars, gold = meta[n_done]
|
||||
if out_f:
|
||||
out_f.write(f"{n_done},{t},{qi},{oi},{clen},{cchars},{gold},{logprob:.6f},{greedy}\n")
|
||||
out_f.flush()
|
||||
n_done += 1
|
||||
if n_done % 5 == 0 or n_done == len(reqs):
|
||||
elapsed = time.time() - t0
|
||||
rate = n_done / elapsed if elapsed > 0 else 0
|
||||
eta = (len(reqs) - n_done) / rate if rate > 0 else 0
|
||||
print(f"[progress] {n_done}/{len(reqs)} requests scored | {elapsed:.0f}s elapsed | "
|
||||
f"{rate:.2f} req/s | ETA {eta:.0f}s | last: {t} q{qi} opt{oi} lp={logprob:.3f}",
|
||||
file=sys.stderr)
|
||||
proc.wait()
|
||||
elapsed = time.time() - t0
|
||||
if out_f:
|
||||
out_f.write(f"# finished: {n_done}/{len(reqs)} in {elapsed:.0f}s, exit={proc.returncode}\n")
|
||||
out_f.close()
|
||||
if proc.returncode != 0 and n_done == 0:
|
||||
print(f"ENGINE ERROR (exit {proc.returncode})", file=sys.stderr); sys.exit(1)
|
||||
if n_done != len(reqs):
|
||||
print(f"WARNING: only {n_done}/{len(reqs)} requests scored (engine exited {proc.returncode}); "
|
||||
f"scoring partial results.", file=sys.stderr)
|
||||
# Fill any unscored slots with -inf so argmax never picks them
|
||||
for i in range(len(lp)):
|
||||
if lp[i] is None: lp[i] = float("-inf")
|
||||
print(f"(engine: {elapsed:.0f}s, {n_done}/{len(reqs)} scored, exit {proc.returncode})", file=sys.stderr)
|
||||
score_accuracy(tasks, meta, perq, lp)
|
||||
print("\nNOTE: compare acc_norm with GLM-5.2's PUBLISHED model-card score. A close result"
|
||||
"\n indicates that int4 quantization preserved quality. (Fill REFERENCE in tools/eval_glm.py.)")
|
||||
|
||||
Reference in New Issue
Block a user