diff --git a/.gitignore b/.gitignore index 203a322..7b2ad19 100644 --- a/.gitignore +++ b/.gitignore @@ -1,62 +1 @@ -# ambienti python / cache -c/mio_env/ -__pycache__/ -**/__pycache__/ -*.pyc -web/node_modules/ -web/dist/ -web/*.tsbuildinfo -desktop/src-tauri/target/ -desktop/src-tauri/gen/ - -# binari compilati (si rigenerano con make / coli build) -c/glm -c/glm.exe -c/olmoe -c/olmoe.exe -c/iobench -c/iobench.exe -c/tok_test -c/backend_cuda.o -c/backend_cuda_test -c/backend_loader.o -c/coli_cuda.dll -c/coli_cuda.lib -c/coli_cuda.exp -c/tests/test_json -c/tests/test_json.exe -c/tests/test_st -c/tests/test_st.exe -c/tests/test_tier -c/tests/test_tier.exe -c/tests/test_grammar -c/tests/test_grammar.exe - -# oracoli tiny generati (make_glm_oracle.py) e dati benchmark scaricati -c/glm_tiny/ -c/glm_tiny_i2/ -c/glm_tiny_i4/ -c/glm_tiny_mix/ -c/glm_bench_medium/ -c/bench/ - -# pesi modello / artefatti di run -*.safetensors -stats*.txt -*.log - -# script di ricerca stadio-0 (esperimenti locali, non parte del motore) -/engine.py -/run.py -/validate_ref.py -/s0_costmodel.py -/s1_trace_hitrate.py -/s2_research.py -/quant_test.py -/profile_run.py -/sweep.py -c/backend_metal.o -c/backend_metal_test -c/tests/test_tier -c/mio_env/ -c/bench/ +.coli_kv\n*.o\n*.log\n__pycache__/\n*.pyc\n/tmp/ diff --git a/c/__pycache__/doctor.cpython-314.pyc b/c/__pycache__/doctor.cpython-314.pyc new file mode 100644 index 0000000..29df673 Binary files /dev/null and b/c/__pycache__/doctor.cpython-314.pyc differ diff --git a/c/__pycache__/openai_server.cpython-314.pyc b/c/__pycache__/openai_server.cpython-314.pyc new file mode 100644 index 0000000..e18495b Binary files /dev/null and b/c/__pycache__/openai_server.cpython-314.pyc differ diff --git a/c/__pycache__/resource_plan.cpython-314.pyc b/c/__pycache__/resource_plan.cpython-314.pyc new file mode 100644 index 0000000..8197f26 Binary files /dev/null and b/c/__pycache__/resource_plan.cpython-314.pyc differ diff --git a/c/backend_hip.cpp b/c/backend_hip.cpp new file mode 100644 index 0000000..808e27b --- /dev/null +++ b/c/backend_hip.cpp @@ -0,0 +1,565 @@ +#include "backend_cuda.h" + +#include +// HIP: mma.h not available, TC_INT4 disabled via env + +#include +#include +#include +#include + +struct ColiCudaTensor { + void *weights; + float *scales; + size_t weight_bytes; + int fmt, I, O, device; + int tracked; +}; + +typedef struct { + int device; + float *x, *y, *gate, *up; + size_t x_cap, y_cap, gate_cap, up_cap; + uint8_t *qx; float *qscale; + size_t qx_cap, qscale_cap; + float *host_x,*host_y; size_t host_x_cap,host_y_cap; + float *aq,*al,*ar,*ac; size_t aq_cap,al_cap,ar_cap,ac_cap; + hipStream_t stream; + void *group_desc; size_t group_desc_cap; + size_t tensor_count, tensor_bytes; +} DeviceContext; + +typedef struct { + const void *g,*u,*d; const float *gs,*us,*ds; + int gf,uf,df,rows,offset; +} GroupDesc; + +static DeviceContext g_ctx[COLI_CUDA_MAX_DEVICES]; +static int g_nctx; +static uint64_t g_group_calls,g_group_experts,g_group_rows; +static double g_group_h2d_ms,g_group_kernel_ms,g_group_d2h_ms; +static std::mutex g_group_stats_mu; + +static int cuda_ok(hipError_t err, const char *what) { + if (err == hipSuccess) return 1; + std::fprintf(stderr, "[CUDA] %s: %s\n", what, hipGetErrorString(err)); + return 0; +} + +static DeviceContext *find_ctx(int device) { + for (int i = 0; i < g_nctx; i++) if (g_ctx[i].device == device) return &g_ctx[i]; + return nullptr; +} + +static int select_ctx(DeviceContext *ctx) { + return ctx && cuda_ok(hipSetDevice(ctx->device), "select device"); +} + +__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 == 3) return (size_t)(I + 3) / 4; + return 0; +} + +__device__ static float weight_at(const void *weights, int fmt, size_t row, int i) { + const uint8_t *base = static_cast(weights) + row; + if (fmt == 0) return reinterpret_cast(base)[i]; + if (fmt == 1) return static_cast(reinterpret_cast(base)[i]); + const uint8_t *q = base; + if (fmt == 2) { + uint8_t v = q[i >> 1]; + int n=(i&1)?(v>>4):(v&15); return static_cast(n&8?n-16:n); + } + uint8_t v = q[i >> 2]; + return static_cast(((v >> ((i & 3) * 2)) & 3) - 2); +} + +__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> 1; n; n >>= 1) { + if (threadIdx.x < n) partial[threadIdx.x] += partial[threadIdx.x + n]; + __syncthreads(); + } + if (!threadIdx.x) + y[(size_t)s * O + o] = partial[0] * (fmt ? scales[o] : 1.0f); +} + +__global__ static void silu_mul(float *gate, const float *up, size_t n) { + size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; + if (i < n) { + float v = gate[i]; + gate[i] = (v / (1.0f + expf(-v))) * up[i]; + } +} + +__global__ static void quantize_s4_rows(uint8_t *q,float *scale,const float *x,int S,int K){ + int s=blockIdx.x; if(s>=S)return; const float *xs=x+(size_t)s*K; + float v=0; for(int i=threadIdx.x;i>=1){if(threadIdx.x0?m[0]/7.f:1.f; if(!threadIdx.x)scale[s]=sc; + uint8_t *dst=q+(size_t)s*((K+1)/2); + for(int b=threadIdx.x;b<(K+1)/2;b+=blockDim.x){ + int i=b*2,a=__float2int_rn(xs[i]/sc),c=i+1= 750 + // HIP: nvcuda removed + int warp=threadIdx.x/32,lane=threadIdx.x%32,tile=blockIdx.x*8+warp,c=blockIdx.y; + if(tile*8>=O)return; GroupDesc d=desc[c]; + const void *w=which==0?d.g:(which==1?d.u:d.d); + const float *ws=which==0?d.gs:(which==1?d.us:d.ds); + int fmt=which==0?d.gf:(which==1?d.uf:d.df); + if(fmt!=2)return; + wmma::fragment acc; wmma::fill_fragment(acc,0); + const uint8_t *a=x+(size_t)d.offset*((K+1)/2); + const uint8_t *b=(const uint8_t*)w+(size_t)(tile*8)*((K+1)/2); + for(int k=0;k af; + wmma::fragment bf; + wmma::load_matrix_sync(af,a+k/2,K); + wmma::load_matrix_sync(bf,b+k/2,K); + wmma::mma_sync(acc,af,bf,acc); + } + __shared__ int out[8][64]; wmma::store_matrix_sync(out[warp],acc,8,wmma::mem_row_major); + for(int i=lane;i<64;i+=32){int s=i/8,o=tile*8+i%8; + if(s=d.rows) return; + const void *w=which?d.u:d.g; const float *sc=which?d.us:d.gs; int fmt=which?d.uf:d.gf; + size_t rb=row_bytes(fmt,D),row=(size_t)o*rb; const float *xs=x+(size_t)(d.offset+s)*D; + float sum=0; for(int i=threadIdx.x;i>=1){ if(threadIdx.x=d.rows) return; + size_t rb=row_bytes(d.df,I),row=(size_t)o*rb; const float *xs=x+(size_t)(d.offset+s)*I; + float sum=0; for(int i=threadIdx.x;i>=1){ if(threadIdx.x>4; *lo=(float)(a&8?a-16:a); *hi=(float)(b&8?b-16:b); +} + +/* Exact low-row W4A32 path. It consumes each packed weight byte once instead + * of routing both nibbles through weight_at(), preserving FP32 activations. */ +__global__ static void grouped_hidden_w4(float *y,const float *x,const GroupDesc *desc, + int I,int D,int which){ + int o=blockIdx.x,s=blockIdx.y,c=blockIdx.z;GroupDesc d=desc[c];if(s>=d.rows)return; + const uint8_t *w=(const uint8_t*)(which?d.u:d.g);const float *sc=which?d.us:d.gs; + const uint8_t *row=w+(size_t)o*((D+1)/2);const float *xs=x+(size_t)(d.offset+s)*D; + float sum=0;for(int b=threadIdx.x;b<(D+1)/2;b+=blockDim.x){float a,z;unpack_s4(row[b],&a,&z); + int i=b*2;sum+=xs[i]*a;if(i+1>=1){if(threadIdx.x=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); + 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;ga+=xs[i]*g0;ua+=xs[i]*u0;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); + 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;sum+=xs[i]*a;if(i+1>=1){if(threadIdx.x= bytes) return 1; + if (*ptr) hipFree(*ptr); + *ptr = nullptr; + *cap = 0; + if (!cuda_ok(hipMalloc(ptr, bytes), "scratch allocation")) return 0; + *cap = bytes; + return 1; +} + +static int reserve_bytes(void **ptr,size_t *cap,size_t bytes){ + if(*cap>=bytes) return 1; if(*ptr) hipFree(*ptr); *ptr=nullptr; *cap=0; + if(!cuda_ok(hipMalloc(ptr,bytes),"descriptor allocation")) return 0; *cap=bytes; return 1; +} + +static int reserve_pinned(float **ptr,size_t *cap,size_t bytes){ + if(*cap>=bytes)return 1;if(*ptr)hipFreeHost(*ptr);*ptr=nullptr;*cap=0; + if(!cuda_ok(hipMallocHost((void**)ptr,bytes),"pinned staging allocation"))return 0;*cap=bytes;return 1; +} + +extern "C" int coli_cuda_init(const int *devices, int count) { + int available = 0; + if (!devices || count < 1 || count > COLI_CUDA_MAX_DEVICES) return 0; + if (!cuda_ok(hipGetDeviceCount(&available), "device discovery")) return 0; + g_nctx = 0; + for (int i = 0; i < count; i++) { + int device = devices[i]; + if (device < 0 || device >= available) { + std::fprintf(stderr, "[CUDA] invalid device %d (available: 0..%d)\n", device, available - 1); + g_nctx = 0; + return 0; + } + if (find_ctx(device)) { + std::fprintf(stderr, "[CUDA] duplicate device %d\n", device); + g_nctx = 0; + return 0; + } + DeviceContext *ctx = &g_ctx[g_nctx]; + *ctx = {}; + ctx->device = device; + if (!select_ctx(ctx)) { g_nctx = 0; return 0; } + hipDeviceProp_t prop{}; + if (!cuda_ok(hipGetDeviceProperties(&prop, device), "device properties")) { g_nctx = 0; return 0; } + if(!cuda_ok(hipStreamCreateWithFlags(&ctx->stream,hipStreamNonBlocking),"stream creation")){ + g_nctx=0;return 0; + } + g_nctx++; + std::fprintf(stderr, "[CUDA] device %d: %s, %.1f GB VRAM, sm_%d%d\n", + device, prop.name, prop.totalGlobalMem / 1e9, prop.major, prop.minor); + } + return 1; +} + +extern "C" void coli_cuda_shutdown(void) { + for (int i = 0; i < g_nctx; i++) { + DeviceContext *ctx = &g_ctx[i]; + if (!select_ctx(ctx)) continue; + if (ctx->x) hipFree(ctx->x); + if (ctx->y) hipFree(ctx->y); + if (ctx->gate) hipFree(ctx->gate); + if (ctx->up) hipFree(ctx->up); + if (ctx->qx) hipFree(ctx->qx); + if (ctx->qscale) hipFree(ctx->qscale); + if(ctx->aq)hipFree(ctx->aq);if(ctx->al)hipFree(ctx->al);if(ctx->ar)hipFree(ctx->ar);if(ctx->ac)hipFree(ctx->ac); + if (ctx->host_x) hipFreeHost(ctx->host_x); + if (ctx->host_y) hipFreeHost(ctx->host_y); + if (ctx->stream) hipStreamDestroy(ctx->stream); + if (ctx->group_desc) hipFree(ctx->group_desc); + ctx->x = ctx->y = ctx->gate = ctx->up = nullptr; + ctx->qx=nullptr; ctx->qscale=nullptr; + ctx->aq=ctx->al=ctx->ar=ctx->ac=nullptr; + ctx->host_x=ctx->host_y=nullptr;ctx->stream=nullptr; + ctx->x_cap = ctx->y_cap = ctx->gate_cap = ctx->up_cap = 0; + ctx->qx_cap=ctx->qscale_cap=0; + ctx->aq_cap=ctx->al_cap=ctx->ar_cap=ctx->ac_cap=0; + ctx->host_x_cap=ctx->host_y_cap=0; + ctx->group_desc=nullptr; ctx->group_desc_cap=0; + } + g_nctx = 0; +} + +extern "C" int coli_cuda_device_count(void) { return g_nctx; } + +extern "C" int coli_cuda_device_at(int index) { + return index >= 0 && index < g_nctx ? g_ctx[index].device : -1; +} + +extern "C" int coli_cuda_mem_info(int device, size_t *free_bytes, size_t *total_bytes) { + DeviceContext *ctx = find_ctx(device); + if (!free_bytes || !total_bytes || !select_ctx(ctx)) return 0; + return cuda_ok(hipMemGetInfo(free_bytes, total_bytes), "memory info"); +} + +extern "C" void coli_cuda_stats(int device, size_t *tensor_count, size_t *tensor_bytes) { + size_t count = 0, bytes = 0; + for (int i = 0; i < g_nctx; i++) if (device < 0 || g_ctx[i].device == device) { + count += g_ctx[i].tensor_count; + bytes += g_ctx[i].tensor_bytes; + } + if (tensor_count) *tensor_count = count; + if (tensor_bytes) *tensor_bytes = bytes; +} + +extern "C" void coli_cuda_group_stats(uint64_t *calls, uint64_t *experts, uint64_t *rows, + double *h2d_ms, double *kernel_ms, double *d2h_ms) { + if(calls) *calls=g_group_calls; if(experts) *experts=g_group_experts; if(rows) *rows=g_group_rows; + if(h2d_ms) *h2d_ms=g_group_h2d_ms; if(kernel_ms) *kernel_ms=g_group_kernel_ms; + if(d2h_ms) *d2h_ms=g_group_d2h_ms; +} + +extern "C" int coli_cuda_tensor_upload(ColiCudaTensor **tensor, + const void *weights, const float *scales, + int fmt, int I, int O, int device) { + DeviceContext *ctx = find_ctx(device); + if (!tensor || !weights || I < 1 || O < 1 || !select_ctx(ctx)) return 0; + size_t rb = row_bytes(fmt, I); + if (!rb || (fmt && !scales)) return 0; + if (*tensor) { + ColiCudaTensor *t = *tensor; + return t->fmt == fmt && t->I == I && t->O == O && t->device == device; + } + ColiCudaTensor *t = static_cast(std::calloc(1, sizeof(*t))); + if (!t) return 0; + t->fmt = fmt; t->I = I; t->O = O; t->device = device; t->weight_bytes = rb * (size_t)O; + if (!cuda_ok(hipMalloc(&t->weights, t->weight_bytes), "tensor allocation") || + !cuda_ok(hipMemcpy(t->weights, weights, t->weight_bytes, hipMemcpyHostToDevice), "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(!cuda_ok(hipGetLastError(),"int4 weight conversion")){coli_cuda_tensor_free(t);return 0;}} + if (fmt) { + if (!cuda_ok(hipMalloc(&t->scales, (size_t)O * sizeof(float)), "scale allocation") || + !cuda_ok(hipMemcpy(t->scales, scales, (size_t)O * sizeof(float), hipMemcpyHostToDevice), "scale upload")) { + coli_cuda_tensor_free(t); + return 0; + } + } + t->tracked = 1; + ctx->tensor_count++; + ctx->tensor_bytes += t->weight_bytes + (fmt ? (size_t)O * sizeof(float) : 0); + *tensor = t; + return 1; +} + +extern "C" int coli_cuda_matmul(ColiCudaTensor **tensor, + float *y, const float *x, + const void *weights, const float *scales, + int fmt, int S, int I, int O, int device) { + if (S < 1 || !coli_cuda_tensor_upload(tensor, weights, scales, fmt, I, O, device)) return 0; + ColiCudaTensor *t = *tensor; + DeviceContext *ctx = find_ctx(t->device); + if (!select_ctx(ctx)) return 0; + size_t rb = row_bytes(fmt, I); + size_t xb = (size_t)S * I * sizeof(float), yb = (size_t)S * O * sizeof(float); + if (!reserve(&ctx->x, &ctx->x_cap, xb) || !reserve(&ctx->y, &ctx->y_cap, yb)) return 0; + if (!cuda_ok(hipMemcpy(ctx->x, x, xb, hipMemcpyHostToDevice), "input upload")) return 0; + dim3 grid((unsigned)O, (unsigned)S); + quant_matmul<<>>(ctx->y, ctx->x, t->weights, t->scales, fmt, S, I, O, rb); + if (!cuda_ok(hipGetLastError(), "matmul launch") || + !cuda_ok(hipMemcpy(y, ctx->y, yb, hipMemcpyDeviceToHost), "output download")) return 0; + return 1; +} + +extern "C" int coli_cuda_expert_mlp(ColiCudaTensor *gate, ColiCudaTensor *up, + ColiCudaTensor *down, float *y, + const float *x, int S) { + if (!gate || !up || !down || !x || !y || S < 1 || + gate->device != up->device || gate->device != down->device || + gate->I != up->I || gate->O != up->O || + down->I != gate->O || down->O != gate->I) return 0; + DeviceContext *ctx = find_ctx(gate->device); + if (!select_ctx(ctx)) return 0; + int D = gate->I, I = gate->O; + size_t xb=(size_t)S*D*sizeof(float), ib=(size_t)S*I*sizeof(float); + size_t yb=(size_t)S*D*sizeof(float); + if (!reserve(&ctx->x,&ctx->x_cap,xb) || !reserve(&ctx->y,&ctx->y_cap,yb) || + !reserve(&ctx->gate,&ctx->gate_cap,ib) || !reserve(&ctx->up,&ctx->up_cap,ib)) return 0; + if (!cuda_ok(hipMemcpy(ctx->x,x,xb,hipMemcpyHostToDevice),"expert input upload")) return 0; + dim3 hidden_grid((unsigned)I,(unsigned)S), output_grid((unsigned)D,(unsigned)S); + quant_matmul<<>>(ctx->gate,ctx->x,gate->weights,gate->scales, + gate->fmt,S,D,I,row_bytes(gate->fmt,D)); + quant_matmul<<>>(ctx->up,ctx->x,up->weights,up->scales, + up->fmt,S,D,I,row_bytes(up->fmt,D)); + size_t n=(size_t)S*I; + silu_mul<<<(unsigned)((n+255)/256),256>>>(ctx->gate,ctx->up,n); + quant_matmul<<>>(ctx->y,ctx->gate,down->weights,down->scales, + down->fmt,S,I,D,row_bytes(down->fmt,I)); + if (!cuda_ok(hipGetLastError(),"expert MLP launch") || + !cuda_ok(hipMemcpy(y,ctx->y,yb,hipMemcpyDeviceToHost),"expert output download")) return 0; + return 1; +} + +extern "C" int coli_cuda_expert_group(ColiCudaTensor *const *gates, + ColiCudaTensor *const *ups, + ColiCudaTensor *const *downs, + const int *rows, int count, + float *y, const float *x) { + if (!gates || !ups || !downs || !rows || !x || !y || count < 1) return 0; + ColiCudaTensor *first=gates[0]; + 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; + 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}; + all_s4&=g->fmt==2&&u->fmt==2&&d->fmt==2; + total+=rows[c]; if(rows[c]>max_rows) max_rows=rows[c]; + } + DeviceContext *ctx=find_ctx(device); if(!select_ctx(ctx)) return 0; + size_t xb=(size_t)total*D*sizeof(float), ib=(size_t)total*I*sizeof(float); + if(!reserve(&ctx->x,&ctx->x_cap,xb)||!reserve(&ctx->y,&ctx->y_cap,xb)|| + !reserve(&ctx->gate,&ctx->gate_cap,ib)||!reserve(&ctx->up,&ctx->up_cap,ib)|| + !reserve_bytes(&ctx->group_desc,&ctx->group_desc_cap,(size_t)count*sizeof(GroupDesc))) return 0; + int async=!getenv("COLI_CUDA_ASYNC")||atoi(getenv("COLI_CUDA_ASYNC")); + if(async&&(!reserve_pinned(&ctx->host_x,&ctx->host_x_cap,xb)|| + !reserve_pinned(&ctx->host_y,&ctx->host_y_cap,xb)))return 0; + hipError_t copy_desc=async?hipMemcpyAsync(ctx->group_desc,host,(size_t)count*sizeof(GroupDesc), + hipMemcpyHostToDevice,ctx->stream) + :hipMemcpy(ctx->group_desc,host,(size_t)count*sizeof(GroupDesc),hipMemcpyHostToDevice); + if(!cuda_ok(copy_desc,"expert group descriptors"))return 0; + int profile=getenv("COLI_CUDA_PROFILE")&&atoi(getenv("COLI_CUDA_PROFILE")); + hipEvent_t ev[4]={}; + if(profile) for(int i=0;i<4;i++) if(!cuda_ok(hipEventCreate(&ev[i]),"profile event")) profile=0; + if(profile) hipEventRecord(ev[0],ctx->stream); + if(async)std::memcpy(ctx->host_x,x,xb); + hipError_t copy_x=async?hipMemcpyAsync(ctx->x,ctx->host_x,xb,hipMemcpyHostToDevice,ctx->stream) + :hipMemcpy(ctx->x,x,xb,hipMemcpyHostToDevice); + if(!cuda_ok(copy_x,"expert group input upload")) return 0; + if(profile) hipEventRecord(ev[1],ctx->stream); + GroupDesc *dev=(GroupDesc*)ctx->group_desc; + int tc=getenv("COLI_CUDA_TC_INT4")&&atoi(getenv("COLI_CUDA_TC_INT4")); + tc=tc&&all_s4&&D%32==0&&I%32==0&&D%8==0&&I%8==0; + int tc_min=getenv("COLI_CUDA_TC_MIN_ROWS")?atoi(getenv("COLI_CUDA_TC_MIN_ROWS")):8; + for(int c=0;c=tc_min; + if(tc){ + size_t qb=(size_t)(total+7)*(size_t)(D>I?D:I)/2; + if(!reserve_bytes((void**)&ctx->qx,&ctx->qx_cap,qb)|| + !reserve(&ctx->qscale,&ctx->qscale_cap,(size_t)(total+7)*sizeof(float)))return 0; + hipMemsetAsync(ctx->qx,0,qb,ctx->stream); + quantize_s4_rows<<stream>>>(ctx->qx,ctx->qscale,ctx->x,total,D); + grouped_s4_wmma<<stream>>>(ctx->gate,ctx->qx,ctx->qscale,dev,D,I,0); + grouped_s4_wmma<<stream>>>(ctx->up,ctx->qx,ctx->qscale,dev,D,I,1); + silu_mul<<<(unsigned)(((size_t)total*I+255)/256),256,0,ctx->stream>>>(ctx->gate,ctx->up,(size_t)total*I); + quantize_s4_rows<<stream>>>(ctx->qx,ctx->qscale,ctx->gate,total,I); + grouped_s4_wmma<<stream>>>(ctx->y,ctx->qx,ctx->qscale,dev,I,D,2); + }else if(all_s4&&(!getenv("COLI_CUDA_W4_PACKED")||atoi(getenv("COLI_CUDA_W4_PACKED")))){ + dim3 hg((unsigned)I,(unsigned)max_rows,(unsigned)count),og((unsigned)D,(unsigned)max_rows,(unsigned)count); + int dual=!getenv("COLI_CUDA_DUAL_PROJ")||atoi(getenv("COLI_CUDA_DUAL_PROJ")); + if(dual)grouped_hidden_w4_dual<<stream>>>(ctx->gate,ctx->up,ctx->x,dev,I,D); + else{ + grouped_hidden_w4<<stream>>>(ctx->gate,ctx->x,dev,I,D,0); + grouped_hidden_w4<<stream>>>(ctx->up,ctx->x,dev,I,D,1); + } + 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{ + dim3 hg((unsigned)I,(unsigned)max_rows,(unsigned)count),og((unsigned)D,(unsigned)max_rows,(unsigned)count); + grouped_hidden<<stream>>>(ctx->gate,ctx->x,dev,I,D,0); + grouped_hidden<<stream>>>(ctx->up,ctx->x,dev,I,D,1); + silu_mul<<<(unsigned)(((size_t)total*I+255)/256),256,0,ctx->stream>>>(ctx->gate,ctx->up,(size_t)total*I); + grouped_down<<stream>>>(ctx->y,ctx->gate,dev,D,I); + } + if(profile) hipEventRecord(ev[2],ctx->stream); + if(!async&&!cuda_ok(hipStreamSynchronize(ctx->stream),"expert group synchronize"))return 0; + hipError_t copy_y=async?hipMemcpyAsync(ctx->host_y,ctx->y,xb,hipMemcpyDeviceToHost,ctx->stream) + :hipMemcpy(y,ctx->y,xb,hipMemcpyDeviceToHost); + if(!cuda_ok(hipGetLastError(),"expert group launch")||!cuda_ok(copy_y,"expert group output download"))return 0; + if(async){if(!cuda_ok(hipStreamSynchronize(ctx->stream),"expert group synchronize"))return 0; + std::memcpy(y,ctx->host_y,xb);} + if(profile){ + hipEventRecord(ev[3],ctx->stream); hipEventSynchronize(ev[3]); float a=0,b=0,c=0; + hipEventElapsedTime(&a,ev[0],ev[1]); hipEventElapsedTime(&b,ev[1],ev[2]); + hipEventElapsedTime(&c,ev[2],ev[3]); + { std::lock_guard lock(g_group_stats_mu); + g_group_h2d_ms+=a; g_group_kernel_ms+=b; g_group_d2h_ms+=c; } + for(int i=0;i<4;i++) hipEventDestroy(ev[i]); + } + { std::lock_guard lock(g_group_stats_mu); + g_group_calls++; g_group_experts+=(uint64_t)count; g_group_rows+=(uint64_t)total; } + return 1; +} + +extern "C" int coli_cuda_attention_absorb(ColiCudaTensor *w,float *ctx,const float *q, + const float *latent,const float *rope,int H,int Q, + int R,int V,int K,int T,float scale){ + if(!w||!ctx||!q||!latent||!rope||H<1||Q<1||R<1||V<1||K<1||K>512||T<1||T>4096|| + w->I!=K||w->O!=H*(Q+V))return 0; + DeviceContext *dc=find_ctx(w->device);if(!select_ctx(dc))return 0; + size_t qb=(size_t)H*(Q+R)*sizeof(float),lb=(size_t)T*K*sizeof(float); + size_t rb=(size_t)T*R*sizeof(float),cb=(size_t)H*V*sizeof(float); + if(!reserve(&dc->aq,&dc->aq_cap,qb)||!reserve(&dc->al,&dc->al_cap,lb)|| + !reserve(&dc->ar,&dc->ar_cap,rb)||!reserve(&dc->ac,&dc->ac_cap,cb))return 0; + if(!cuda_ok(hipMemcpyAsync(dc->aq,q,qb,hipMemcpyHostToDevice,dc->stream),"attention q upload")|| + !cuda_ok(hipMemcpyAsync(dc->al,latent,lb,hipMemcpyHostToDevice,dc->stream),"attention latent upload")|| + !cuda_ok(hipMemcpyAsync(dc->ar,rope,rb,hipMemcpyHostToDevice,dc->stream),"attention rope upload"))return 0; + size_t shared=(size_t)(2*K+T)*sizeof(float); + attention_absorb_kernel<<stream>>>(dc->ac,dc->aq,dc->al,dc->ar,w->weights,w->scales, + w->fmt,H,Q,R,V,K,T,scale); + if(!cuda_ok(hipGetLastError(),"attention absorb launch")|| + !cuda_ok(hipMemcpyAsync(ctx,dc->ac,cb,hipMemcpyDeviceToHost,dc->stream),"attention context download")|| + !cuda_ok(hipStreamSynchronize(dc->stream),"attention synchronize"))return 0; + return 1; +} + +extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) { + if (!tensor) return; + DeviceContext *ctx = find_ctx(tensor->device); + if (ctx) select_ctx(ctx); + if (tensor->tracked && ctx) { + size_t bytes = tensor->weight_bytes + (tensor->fmt ? (size_t)tensor->O * sizeof(float) : 0); + if (ctx->tensor_count) ctx->tensor_count--; + if (ctx->tensor_bytes >= bytes) ctx->tensor_bytes -= bytes; + } + if (tensor->weights) hipFree(tensor->weights); + if (tensor->scales) hipFree(tensor->scales); + std::free(tensor); +} + +extern "C" size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor) { + return tensor ? tensor->weight_bytes + (tensor->fmt ? (size_t)tensor->O * sizeof(float) : 0) : 0; +} + +extern "C" int coli_cuda_tensor_device(const ColiCudaTensor *tensor) { + return tensor ? tensor->device : -1; +} diff --git a/c/backend_hip.o b/c/backend_hip.o new file mode 100644 index 0000000..21d7323 Binary files /dev/null and b/c/backend_hip.o differ diff --git a/c/coli b/c/coli index 7ff7812..734a732 100755 --- a/c/coli +++ b/c/coli @@ -110,7 +110,7 @@ def cuda_binary(): if not os.path.exists(GLM) or sys.platform != "linux": return False try: linked=subprocess.run(["ldd",GLM],capture_output=True,text=True,timeout=3) - return any("libcudart" in line and "not found" not in line + return any("libamdhip64" in line and "not found" not in line for line in linked.stdout.splitlines()) except (OSError,subprocess.SubprocessError): return False diff --git a/c/glm b/c/glm new file mode 100755 index 0000000..5e908b2 Binary files /dev/null and b/c/glm differ diff --git a/c/glm_cpu_backup b/c/glm_cpu_backup new file mode 100755 index 0000000..b99dbdf Binary files /dev/null and b/c/glm_cpu_backup differ diff --git a/c/glm_cpu_v1 b/c/glm_cpu_v1 new file mode 100755 index 0000000..b99dbdf Binary files /dev/null and b/c/glm_cpu_v1 differ diff --git a/c/glm_hip b/c/glm_hip new file mode 100755 index 0000000..5e908b2 Binary files /dev/null and b/c/glm_hip differ diff --git a/c/tests/__pycache__/test_benchmark_cuda_fixture.cpython-314.pyc b/c/tests/__pycache__/test_benchmark_cuda_fixture.cpython-314.pyc new file mode 100644 index 0000000..bbc4f12 Binary files /dev/null and b/c/tests/__pycache__/test_benchmark_cuda_fixture.cpython-314.pyc differ diff --git a/c/tests/__pycache__/test_cli_output.cpython-314.pyc b/c/tests/__pycache__/test_cli_output.cpython-314.pyc new file mode 100644 index 0000000..5440a7d Binary files /dev/null and b/c/tests/__pycache__/test_cli_output.cpython-314.pyc differ diff --git a/c/tests/__pycache__/test_doctor.cpython-314.pyc b/c/tests/__pycache__/test_doctor.cpython-314.pyc new file mode 100644 index 0000000..2ada71c Binary files /dev/null and b/c/tests/__pycache__/test_doctor.cpython-314.pyc differ diff --git a/c/tests/__pycache__/test_openai_server.cpython-314.pyc b/c/tests/__pycache__/test_openai_server.cpython-314.pyc new file mode 100644 index 0000000..8f5c530 Binary files /dev/null and b/c/tests/__pycache__/test_openai_server.cpython-314.pyc differ diff --git a/c/tests/__pycache__/test_resource_plan.cpython-314.pyc b/c/tests/__pycache__/test_resource_plan.cpython-314.pyc new file mode 100644 index 0000000..9cc1be9 Binary files /dev/null and b/c/tests/__pycache__/test_resource_plan.cpython-314.pyc differ diff --git a/c/tests/test_decode_batch b/c/tests/test_decode_batch index 989fe42..cdcca76 100755 Binary files a/c/tests/test_decode_batch and b/c/tests/test_decode_batch differ diff --git a/c/tests/test_grammar b/c/tests/test_grammar new file mode 100755 index 0000000..70294a9 Binary files /dev/null and b/c/tests/test_grammar differ diff --git a/c/tests/test_i4_acc512 b/c/tests/test_i4_acc512 index 76f68a3..99bdca5 100755 Binary files a/c/tests/test_i4_acc512 and b/c/tests/test_i4_acc512 differ diff --git a/c/tests/test_idot b/c/tests/test_idot index 74dd03c..a685f5c 100755 Binary files a/c/tests/test_idot and b/c/tests/test_idot differ diff --git a/c/tests/test_json b/c/tests/test_json new file mode 100755 index 0000000..31b1812 Binary files /dev/null and b/c/tests/test_json differ diff --git a/c/tests/test_st b/c/tests/test_st new file mode 100755 index 0000000..f1f381a Binary files /dev/null and b/c/tests/test_st differ diff --git a/c/tests/test_tier b/c/tests/test_tier new file mode 100755 index 0000000..729a4fe Binary files /dev/null and b/c/tests/test_tier differ diff --git a/c/tools/__pycache__/__init__.cpython-314.pyc b/c/tools/__pycache__/__init__.cpython-314.pyc new file mode 100644 index 0000000..dda5755 Binary files /dev/null and b/c/tools/__pycache__/__init__.cpython-314.pyc differ diff --git a/c/tools/__pycache__/benchmark_cuda_fixture.cpython-314.pyc b/c/tools/__pycache__/benchmark_cuda_fixture.cpython-314.pyc new file mode 100644 index 0000000..dff15b6 Binary files /dev/null and b/c/tools/__pycache__/benchmark_cuda_fixture.cpython-314.pyc differ