diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu index 4bab7c0..cbd8d85 100644 --- a/c/backend_cuda.cu +++ b/c/backend_cuda.cu @@ -20,6 +20,8 @@ struct ColiCudaTensor { float *scales; size_t weight_bytes; int fmt, I, O, device; + int gs; /* quant group size; 0 = per-row scales (#334) */ + size_t scale_count; /* floats in `scales`: O per-row, O*ng grouped */ int tracked; RaggedKVEntry ragged[512]; int ragged_count; @@ -43,6 +45,7 @@ typedef struct { typedef struct { const void *g,*u,*d; const float *gs,*us,*ds; int gf,uf,df,rows,offset; + int ggs,ugs,dgs; /* per-tensor quant group size; 0 = per-row scales (#334 fmt=4) */ } GroupDesc; static DeviceContext g_ctx[COLI_CUDA_MAX_DEVICES]; @@ -81,6 +84,7 @@ __host__ __device__ static size_t row_bytes(int fmt, int I) { if (fmt == 1) return (size_t)I; if (fmt == 2) return (size_t)(I + 1) / 2; if (fmt == 3) return (size_t)(I + 3) / 4; + if (fmt == 4) return (size_t)(I + 1) / 2; /* grouped int4: nibbles like fmt 2 */ return 0; } @@ -296,6 +300,42 @@ __global__ static void grouped_down_w4(float *y,const float *x,const GroupDesc * if(!threadIdx.x)y[(size_t)(d.offset+s)*D+o]=p[0]*d.ds[o]; } +/* fmt=4 grouped-int4 variants (#334): identical structure to the w4 kernels, + * but the scale varies along the input dimension — sc[o*ng + i/gs], applied + * per element inside the accumulation (gs is even, so a packed byte never + * straddles a group). gs<=0 degrades to per-row (ng=1), so mixed fmt2/fmt4 + * groups run correctly through this one kernel family. */ +__global__ static void grouped_hidden_g4_dual(float *gate,float *up,const float *x, + const GroupDesc *desc,int I,int D){ + int o=blockIdx.x,s=blockIdx.y,c=blockIdx.z;GroupDesc d=desc[c];if(s>=d.rows)return; + const uint8_t *gr=(const uint8_t*)d.g+(size_t)o*((D+1)/2); + const uint8_t *ur=(const uint8_t*)d.u+(size_t)o*((D+1)/2); + int ggs=d.ggs>0?d.ggs:D, ugs=d.ugs>0?d.ugs:D; + const float *gsc=d.gs+(size_t)o*(size_t)((D+ggs-1)/ggs); + const float *usc=d.us+(size_t)o*(size_t)((D+ugs-1)/ugs); + const float *xs=x+(size_t)(d.offset+s)*D;float ga=0,ua=0; + for(int b=threadIdx.x;b<(D+1)/2;b+=blockDim.x){float g0,g1,u0,u1;unpack_s4(gr[b],&g0,&g1);unpack_s4(ur[b],&u0,&u1); + int i=b*2;float gv=gsc[i/ggs],uv=usc[i/ugs]; + ga+=xs[i]*g0*gv;ua+=xs[i]*u0*uv; + if(i+1>=1){if(threadIdx.x=d.rows)return; + const uint8_t *row=(const uint8_t*)d.d+(size_t)o*((I+1)/2); + int dgs=d.dgs>0?d.dgs:I; + const float *dsc=d.ds+(size_t)o*(size_t)((I+dgs-1)/dgs); + const float *xs=x+(size_t)(d.offset+s)*I;float sum=0; + for(int b=threadIdx.x;b<(I+1)/2;b+=blockDim.x){float a,z;unpack_s4(row[b],&a,&z); + int i=b*2;float sv=dsc[i/dgs]; + sum+=xs[i]*a*sv;if(i+1>=1){if(threadIdx.x(std::calloc(1, sizeof(*t))); if (!t) return 0; t->fmt = fmt; t->I = I; t->O = O; t->device = device; t->weight_bytes = rb * (size_t)O; + 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; if (!cuda_ok(cudaMalloc(&t->weights, t->weight_bytes), "tensor allocation") || !cuda_ok(cudaMemcpy(t->weights, weights, t->weight_bytes, cudaMemcpyHostToDevice), "tensor upload")) { coli_cuda_tensor_free(t); return 0; } - if(fmt==2){offset_to_signed_s4<<<(unsigned)((t->weight_bytes+255)/256),256>>>((uint8_t*)t->weights,t->weight_bytes); + if(fmt==2||fmt==4){ /* same nibble layout: offset-binary -> signed in place */ + offset_to_signed_s4<<<(unsigned)((t->weight_bytes+255)/256),256>>>((uint8_t*)t->weights,t->weight_bytes); if(!cuda_ok(cudaGetLastError(),"int4 weight conversion")){coli_cuda_tensor_free(t);return 0;}} if (fmt) { - if (!cuda_ok(cudaMalloc(&t->scales, (size_t)O * sizeof(float)), "scale allocation") || - !cuda_ok(cudaMemcpy(t->scales, scales, (size_t)O * sizeof(float), cudaMemcpyHostToDevice), "scale upload")) { + if (!cuda_ok(cudaMalloc(&t->scales, t->scale_count * sizeof(float)), "scale allocation") || + !cuda_ok(cudaMemcpy(t->scales, scales, t->scale_count * sizeof(float), cudaMemcpyHostToDevice), "scale upload")) { coli_cuda_tensor_free(t); return 0; } } t->tracked = 1; ctx->tensor_count++; - ctx->tensor_bytes += t->weight_bytes + (fmt ? (size_t)O * sizeof(float) : 0); + ctx->tensor_bytes += t->weight_bytes + (fmt ? t->scale_count * sizeof(float) : 0); *tensor = t; return 1; } +extern "C" int coli_cuda_tensor_upload_g(ColiCudaTensor **tensor, + const void *weights, const float *scales, + int fmt, int I, int O, int device, int gs){ + g_upload_gs = gs>0 ? gs : 0; + int r = coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device); + g_upload_gs = 0; + return r; +} extern "C" int coli_cuda_tensor_update(ColiCudaTensor *tensor, const void *weights, @@ -546,13 +604,14 @@ extern "C" int coli_cuda_tensor_update(ColiCudaTensor *tensor, if (!select_ctx(ctx)) return 0; if (!cuda_ok(cudaMemcpy(tensor->weights,weights,tensor->weight_bytes, cudaMemcpyHostToDevice),"tensor refresh")) return 0; - if(tensor->fmt==2){ + if(tensor->fmt==2||tensor->fmt==4){ offset_to_signed_s4<<<(unsigned)((tensor->weight_bytes+255)/256),256>>>( (uint8_t*)tensor->weights,tensor->weight_bytes); if(!cuda_ok(cudaGetLastError(),"int4 weight refresh")) return 0; } return !tensor->fmt || cuda_ok(cudaMemcpy(tensor->scales,scales, - (size_t)tensor->O*sizeof(float),cudaMemcpyHostToDevice),"scale refresh"); + (tensor->scale_count?tensor->scale_count:(size_t)tensor->O)*sizeof(float), + cudaMemcpyHostToDevice),"scale refresh"); } extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor, @@ -641,14 +700,18 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates, if (!first) return 0; int device=first->device,D=first->I,I=first->O,total=0,max_rows=0; GroupDesc host[64]; if(count>64) return 0; - int all_s4=1; + int all_s4=1,all_q4=1,any_g4=0; for(int c=0;cdevice!=device||u->device!=device||d->device!=device|| g->I!=D||u->I!=D||g->O!=I||u->O!=I||d->I!=I||d->O!=D) return 0; host[c]={g->weights,u->weights,d->weights,g->scales,u->scales,d->scales, - g->fmt,u->fmt,d->fmt,rows[c],total}; + 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; + 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; total+=rows[c]; if(rows[c]>max_rows) max_rows=rows[c]; } DeviceContext *ctx=find_ctx(device); if(!select_ctx(ctx)) return 0; @@ -738,7 +801,18 @@ extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates, } silu_mul<<<(unsigned)(((size_t)total*I+255)/256),256,0,ctx->stream>>>(ctx->gate,ctx->up,(size_t)total*I); grouped_down_w4<<stream>>>(ctx->y,ctx->gate,dev,D,I); + }else if(all_q4&&any_g4){ + /* grouped-int4 (fmt=4) present: per-group scales (#334). fmt=2 members + * ride along as the ng=1 special case. */ + dim3 hg((unsigned)I,(unsigned)max_rows,(unsigned)count),og((unsigned)D,(unsigned)max_rows,(unsigned)count); + grouped_hidden_g4_dual<<stream>>>(ctx->gate,ctx->up,ctx->x,dev,I,D); + silu_mul<<<(unsigned)(((size_t)total*I+255)/256),256,0,ctx->stream>>>(ctx->gate,ctx->up,(size_t)total*I); + grouped_down_g4<<stream>>>(ctx->y,ctx->gate,dev,D,I); }else{ + /* generic path decodes fmt 0/1/2/3 only — a fmt=4 group that slipped the + * gates above (odd gs) must NOT be silently decoded as int2 (#334). */ + for(int c=0;cstream>>>(ctx->gate,ctx->x,dev,I,D,0); grouped_hidden<<stream>>>(ctx->up,ctx->x,dev,I,D,1); diff --git a/c/backend_cuda.h b/c/backend_cuda.h index eb959c2..1d7e642 100644 --- a/c/backend_cuda.h +++ b/c/backend_cuda.h @@ -36,6 +36,9 @@ COLI_CUDA_DLLEXPORT void coli_cuda_group_stats(uint64_t *calls, uint64_t *expert double *h2d_ms, double *kernel_ms, double *d2h_ms); /* Upload without executing, so capacity failures happen during model startup. */ +COLI_CUDA_DLLEXPORT int coli_cuda_tensor_upload_g(ColiCudaTensor **tensor, + const void *weights, const float *scales, + int fmt, int I, int O, int device, int gs); COLI_CUDA_DLLEXPORT int coli_cuda_tensor_upload(ColiCudaTensor **tensor, const void *weights, const float *scales, int fmt, int I, int O, int device); diff --git a/c/backend_loader.c b/c/backend_loader.c index c87fbcd..7b3b1a7 100644 --- a/c/backend_loader.c +++ b/c/backend_loader.c @@ -46,6 +46,7 @@ typedef int (*fn_attention_absorb)(ColiCudaTensor *kv_b, float *ctx, int R, int V, int K, int T, float attention_scale); typedef int (*fn_tensor_upload)(ColiCudaTensor **tensor, const void *weights, const float *scales, int fmt, int I, int O, int device); +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); @@ -103,6 +104,7 @@ static struct { fn_expert_group expert_group; fn_attention_absorb attention_absorb; fn_tensor_upload tensor_upload; + fn_tensor_upload_g tensor_upload_g; fn_matmul matmul; fn_tensor_free tensor_free; fn_tensor_bytes tensor_bytes; @@ -198,6 +200,7 @@ static int coli_cuda_load(void){ RESOLVE(expert_group, fn_expert_group) RESOLVE(attention_absorb, fn_attention_absorb) RESOLVE(tensor_upload, fn_tensor_upload) + RESOLVE(tensor_upload_g, fn_tensor_upload_g) RESOLVE(matmul, fn_matmul) RESOLVE(tensor_free, fn_tensor_free) RESOLVE(tensor_bytes, fn_tensor_bytes) @@ -305,6 +308,11 @@ int coli_cuda_tensor_upload(ColiCudaTensor **tensor, const void *weights, return g_cuda.tensor_upload(tensor, weights, scales, fmt, I, O, device); } +int coli_cuda_tensor_upload_g(ColiCudaTensor **tensor, const void *weights, const float *scales, int fmt, int I, int O, int device, int gs){ + if(!g_cuda.available || !g_cuda.tensor_upload_g){ return 0; } + return g_cuda.tensor_upload_g(tensor, weights, scales, fmt, I, O, device, gs); +} + 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){ diff --git a/c/colibri.c b/c/colibri.c index a5044a5..cb0048d 100644 --- a/c/colibri.c +++ b/c/colibri.c @@ -263,6 +263,11 @@ static void qt_cuda_reset(QT *t){ static int qt_cuda_upload(QT *t){ const void *weights = t->fmt==0 ? (const void*)t->qf : t->fmt==1 ? (const void*)t->q8 : (const void*)t->q4; + if(t->fmt==4) /* grouped int4 (#334): scales are [O, ceil(I/gs)] — the plain + * upload would truncate them to O floats and the group kernels + * would read garbage. An old DLL without the _g symbol returns 0 + * and the tensor simply stays CPU-side. */ + return coli_cuda_tensor_upload_g(&t->cuda,weights,t->s,t->fmt,t->I,t->O,t->cuda_device,t->gs); return coli_cuda_tensor_upload(&t->cuda,weights,t->s,t->fmt,t->I,t->O,t->cuda_device); } static int qt_cuda_update(QT *t){ diff --git a/c/tests/test_grouped_g4_cuda.cu b/c/tests/test_grouped_g4_cuda.cu new file mode 100644 index 0000000..b370cdd --- /dev/null +++ b/c/tests/test_grouped_g4_cuda.cu @@ -0,0 +1,108 @@ +/* Grouped-int4 (fmt=4) CUDA kernel oracle (#334). + * + * Feeds random offset-binary nibble weights + [O, ng] group scales through + * grouped_hidden_g4_dual / grouped_down_g4 and checks against a CPU reference + * that replicates matmul_i4_grouped's semantics (value = nibble - 8, per-group + * partial dot x scale). Covers gs=64, a non-divisible tail group, and a + * per-row (gs=0) member riding in the same launch — the fmt=2-compat case. + * + * The device buffers get the same XOR 0x88 offset->signed conversion the + * upload path applies, so the kernels are exercised exactly as deployed. + * + * Build: nvcc -O2 -std=c++17 -arch=native tests/test_grouped_g4_cuda.cu -o tests/test_grouped_g4 + */ +#include +#include +#include +#include +#include + +#include "../backend_cuda.cu" + +static void cpu_gemv_g4(const uint8_t *q,const float *sc,int K,int O,int gs, + const float *x,float *y){ + int rb=(K+1)/2, ng=gs>0?(K+gs-1)/gs:1, egs=gs>0?gs:K; + for(int o=0;oK) glen=K-base; + double p=0; + for(int i=base;i>1]; int n=(i&1)?(v>>4):(v&15); + p+=(double)x[i]*(n-8); + } + a+=p*scl[g]; + } + y[o]=(float)a; + } +} + +int main(void){ + srand(7); + const int D=200, I=96, gs=64; /* tail group: 200 % 64 = 8 */ + const int COUNT=3; /* expert 0,1: fmt4 gs=64; expert 2: per-row (gs=0) */ + const int rbD=(D+1)/2, rbI=(I+1)/2; + const int ngD=(D+gs-1)/gs, ngI=(I+gs-1)/gs; + int trials=50, bad=0; + for(int t=0;t>>(qg[c],(size_t)I*rbD); + offset_to_signed_s4<<<64,256>>>(qu[c],(size_t)I*rbD); + offset_to_signed_s4<<<64,256>>>(qd[c],(size_t)D*rbI); + cudaMemcpy(sg[c],hgs[c],(size_t)I*cngD*4,cudaMemcpyHostToDevice); + cudaMemcpy(su[c],hus[c],(size_t)I*cngD*4,cudaMemcpyHostToDevice); + cudaMemcpy(sd[c],hds[c],(size_t)D*cngI*4,cudaMemcpyHostToDevice); + host[c]={qg[c],qu[c],qd[c],sg[c],su[c],sd[c],4,4,4,1,c,cgs,cgs,cgs}; + } + for(size_t i=0;i<(size_t)COUNT*D;i++) xs[i]=(rand()/(float)RAND_MAX-.5f)*2.f; + GroupDesc *ddesc; cudaMalloc(&ddesc,sizeof(host)); + cudaMemcpy(ddesc,host,sizeof(host),cudaMemcpyHostToDevice); + dim3 hgd((unsigned)I,1,(unsigned)COUNT),ogd((unsigned)D,1,(unsigned)COUNT); + grouped_hidden_g4_dual<<>>(gate,up,xs,ddesc,I,D); + grouped_down_g4<<>>(y,gate,ddesc,D,I); + if(cudaDeviceSynchronize()!=cudaSuccess){ printf("FAIL cuda\n"); return 1; } + for(int c=0;c1e-3f*(fabsf(rg[o])+1e-3f)|| + fabsf(up[(size_t)c*I+o]-ru[o])>1e-3f*(fabsf(ru[o])+1e-3f)) bad++; + } + cpu_gemv_g4(hd[c],hds[c],I,D,cgs,(float*)gate+(size_t)c*I,ry); + for(int o=0;o1e-3f*(fabsf(ry[o])+1e-3f)) bad++; + } + for(int c=0;c