diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml new file mode 100644 index 0000000..e471fb9 --- /dev/null +++ b/.github/workflows/release.yml @@ -0,0 +1,104 @@ +name: Release + +on: + push: + tags: ['v*'] + +permissions: + contents: write + +jobs: + build: + strategy: + matrix: + include: + - os: ubuntu-latest + name: linux-x86_64 + ext: "" + make_args: "ARCH=x86-64-v3" + - os: macos-latest + name: macos-arm64 + ext: "" + make_args: "" + - os: windows-latest + name: windows-x86_64 + ext: ".exe" + make_args: "" + # GitHub shells are format strings: 'msys2 {0}', never bare 'msys2' + # (the v1.0.0 tag build failed on exactly this). Same as ci.yml. + shell: "msys2 {0}" + runs-on: ${{ matrix.os }} + defaults: + run: + shell: ${{ matrix.shell || 'bash' }} + steps: + - uses: actions/checkout@v4 + + - if: matrix.os == 'windows-latest' + uses: msys2/setup-msys2@v2 + with: + msystem: UCRT64 + update: false + # inherit: the Package step calls the runner's 7z.exe, which lives on + # the Windows PATH; the default minimal path would not see it. + path-type: inherit + install: >- + make + mingw-w64-ucrt-x86_64-gcc + + - if: matrix.os == 'macos-latest' + run: brew install libomp + + - name: Build engine + run: | + cd c + make glm ${{ matrix.make_args }} + ls -lh glm${{ matrix.ext }} + + - name: Package + run: | + TAG=${GITHUB_REF#refs/tags/} + mkdir -p dist + cp c/glm${{ matrix.ext }} dist/colibri-${TAG}-${{ matrix.name }}${{ matrix.ext }} + cp c/coli dist/ + cp c/version.py dist/ + cp c/openai_server.py dist/ + cp c/resource_plan.py dist/ + cp c/doctor.py dist/ + cp LICENSE dist/ + cd dist + if [ "${{ matrix.ext }}" = ".exe" ]; then + 7z a colibri-${TAG}-${{ matrix.name }}.zip * + else + tar czf colibri-${TAG}-${{ matrix.name }}.tar.gz * + fi + + - uses: actions/upload-artifact@v4 + with: + name: colibri-${{ matrix.name }} + path: dist/colibri-*.* + + release: + needs: build + runs-on: ubuntu-latest + steps: + - uses: actions/checkout@v4 + + - uses: actions/download-artifact@v4 + with: + path: artifacts + merge-multiple: true + + - name: Create GitHub Release + env: + GH_TOKEN: ${{ github.token }} + run: | + TAG=${GITHUB_REF#refs/tags/} + # Extract the latest version section from CHANGELOG + NOTES=$(awk "/^## \\[${TAG#v}\\]/{found=1;next} /^## \\[/{if(found)exit} found{print}" CHANGELOG.md) + if [ -z "$NOTES" ]; then + NOTES="Release ${TAG}" + fi + gh release create "$TAG" artifacts/* \ + --title "colibrì ${TAG}" \ + --notes "$NOTES" diff --git a/CHANGELOG.md b/CHANGELOG.md new file mode 100644 index 0000000..6177966 --- /dev/null +++ b/CHANGELOG.md @@ -0,0 +1,53 @@ +# Changelog + +All notable changes to colibrì are documented here. +Format follows [Keep a Changelog](https://keepachangelog.com/). + +## [1.0.0] — 2026-07-19 + +First tagged release. The engine has been running in production since late June +2026; this tag marks the baseline for semantic versioning going forward. + +### Highlights + +- **GLM-5.2 (744B MoE)** runs on ~25 GB RAM in pure C, streaming experts from disk +- **Three-tier placement**: VRAM (hot) / RAM (warm) / NVMe (cold), with a learning + cache that pins your workload's hottest experts automatically +- **CUDA backend**: multi-GPU expert tier, dense tensor distribution, batched + ragged attention, resident pipeline (`COLI_CUDA_PIPE=2`) +- **Metal backend** (Apple Silicon): batched expert SwiGLU + fused decode attention + on unified memory GPU +- **MTP speculation**: native GLM-5.2 draft heads, grammar-forced drafts, kernel- + pinned verification (`SPEC_PIN=1`) +- **OpenAI-compatible API**: `coli serve` with SSE streaming, KV slots, bounded + queue, web dashboard (`coli web`) +- **Web UI**: chat with live metrics, expert cortex brain page, profiling breakdown, + expert atlas 3-D galaxy +- **Cross-platform**: Linux, macOS, Windows 11 (native MinGW), PowerPC; CI on all three +- **Auto-tune**: `coli plan --auto-tier` classifies the bottleneck and derives + MTP/PIPE/NUMA/PIN settings with explanations + +### Engine + +- Token-exact validation against `transformers` oracle (teacher-forcing 32/32) +- Compressed MLA KV cache (576 floats/token, 57× smaller), persisted across + restarts (`.coli_kv`, zero re-prefill) +- DSA sparse attention (lightning indexer), faithfully implemented +- Router-lookahead prefetch (`PILOT=1`, 71.6% predictive) +- Async expert I/O pool (`PIPE=1`), io_uring batching (`URING=1`) +- NUMA-aware expert placement (`COLI_NUMA=1`, +13–40% on multi-socket) +- AVX2 / AVX-512 / AVX-VNNI / ARM NEON / NEON-i8mm / POWER VSX kernels +- int4 / int8 / int2 / grouped-int4 (fmt=4) quantization formats + +### Tools + +- `coli convert` — FP8→int4 one-shard-at-a-time converter +- `coli doctor` — read-only setup diagnostics +- `coli plan` — resource planner with auto-tune prescription +- `coli bench` — MMLU / HellaSwag / ARC quality benchmarks +- Expert atlas (`tools/analyze.py --web`) — measured topic affinity for 19,456 experts + +### Community + +- 30+ hardware datapoints in the benchmark tracker +- Contributions from 20+ authors across engine, docs, tooling, and ports diff --git a/c/Makefile b/c/Makefile index feb3045..1f8583b 100644 --- a/c/Makefile +++ b/c/Makefile @@ -166,7 +166,7 @@ else PYTHON ?= python3 endif CUDA_OBJ = -TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) +TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_sample_nan$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE) ifneq (,$(LINUX)) TEST_BINS += tests/test_uring$(EXE) endif @@ -328,9 +328,15 @@ tests/test_topp$(EXE): tests/test_topp.c glm.c st.h uring.h json.h tok.h tok_uni tests/bench_topp$(EXE): tests/bench_topp.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +tests/test_sample_nan$(EXE): tests/test_sample_nan.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + tests/test_kv_alloc$(EXE): tests/test_kv_alloc.c glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) +tests/test_logit_nan$(EXE): tests/test_logit_nan.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h + $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) + tests/test_i4_acc512$(EXE): tests/test_i4_acc512.c $(CC) $(CFLAGS) $< -o $@ $(LDFLAGS) diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu index 69a224f..d808545 100644 --- a/c/backend_cuda.cu +++ b/c/backend_cuda.cu @@ -8,12 +8,21 @@ #include #include +struct RaggedKVEntry { + const void *key; + const float *host_l,*host_r; + float *latent,*rope; + int length,capacity,K,R; +}; + struct ColiCudaTensor { void *weights; float *scales; size_t weight_bytes; int fmt, I, O, device; int tracked; + RaggedKVEntry ragged[512]; + int ragged_count; }; typedef struct { @@ -23,7 +32,7 @@ typedef struct { 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 *host_x,*host_y,*host_kv; size_t host_x_cap,host_y_cap,host_kv_cap; float *aq,*al,*ar,*ac; size_t aq_cap,al_cap,ar_cap,ac_cap; float *pipe_buf[24]; size_t pipe_cap[24]; /* scratch persistenti del resident pipeline */ cudaStream_t stream; @@ -340,6 +349,49 @@ __global__ static void attention_absorb_batch_kernel(float *ctx,const float *q, ctx[((size_t)s*H+h)*V+v]=a*(fmt?wscale[row]:1.f);} } +/* Independent device-resident KV sequence per row. lengths selects the valid + * prefix; latent/rope point at paged caches updated by the host wrapper. */ +__global__ static void attention_absorb_ragged_kernel(float *ctx,const float *q, + const float *const *latent,const float *const *rope,const int *lengths, + 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 s=blockIdx.y,h=blockIdx.x,tid=threadIdx.x,nt=lengths[s],rbase=h*(Q+V); + if(s>=S||nt<1||nt>T)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); + const float *ls=latent[s],*rs=rope[s]; + for(int k=tid;k>1;n;n>>=1){if(tid>1;n;n>>=1){if(tid= bytes) return 1; if (*ptr) cudaFree(*ptr); @@ -408,16 +460,17 @@ extern "C" void coli_cuda_shutdown(void) { for(int b=0;b<24;b++) if(ctx->pipe_buf[b]) cudaFree(ctx->pipe_buf[b]); if (ctx->host_x) cudaFreeHost(ctx->host_x); if (ctx->host_y) cudaFreeHost(ctx->host_y); + if (ctx->host_kv) cudaFreeHost(ctx->host_kv); if (ctx->stream) cudaStreamDestroy(ctx->stream); if (ctx->group_desc) cudaFree(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->host_x=ctx->host_y=ctx->host_kv=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->host_x_cap=ctx->host_y_cap=ctx->host_kv_cap=0; ctx->group_desc=nullptr; ctx->group_desc_cap=0; } g_nctx = 0; @@ -795,6 +848,89 @@ extern "C" int coli_cuda_attention_project_batch(ColiCudaTensor *w,ColiCudaTenso return attention_absorb_batch_run(w,proj,out,q,latent,rope,S,H,Q,R,V,K,T,scale); } +extern "C" int coli_cuda_attention_project_ragged(ColiCudaTensor *w,ColiCudaTensor *proj, + float *out,const float *q,const void *const *keys, + const float *const *latent,const float *const *rope, + const int *lengths,int S,int H,int Q,int R,int V,int K,int T,float scale){ + if(!w||!proj||!out||!q||!keys||!latent||!rope||!lengths||S<1||S>512||T<1||T>8192|| + H<1||Q<1||R<1||V<1||K<1||K>512||w->I!=K||w->O!=H*(Q+V)|| + proj->device!=w->device||proj->I!=H*V)return 0; + DeviceContext *dc=find_ctx(w->device); + if(!select_ctx(dc))return 0; + float **dl=(float**)std::malloc((size_t)S*sizeof(*dl)); + float **dr=(float**)std::malloc((size_t)S*sizeof(*dr)); + int *old=(int*)std::malloc((size_t)S*sizeof(*old)); + int *add=(int*)std::malloc((size_t)S*sizeof(*add)); + int *off=(int*)std::malloc((size_t)S*sizeof(*off));int packed_n=0; + if(!dl||!dr||!old||!add||!off){std::free(dl);std::free(dr);std::free(old);std::free(add);std::free(off);return 0;} + for(int s=0;sT){std::free(dl);std::free(dr);std::free(old);std::free(add);std::free(off);return 0;} + RaggedKVEntry *e=nullptr; + for(int i=0;iragged_count;i++)if(w->ragged[i].key==keys[s]){e=&w->ragged[i];break;} + if(!e){ + if(w->ragged_count>=512){std::free(dl);std::free(dr);std::free(old);std::free(add);std::free(off);return 0;} + e=&w->ragged[w->ragged_count++];std::memset(e,0,sizeof(*e));e->key=keys[s]; + } + if(e->K!=K||e->R!=R||e->host_l!=latent[s]||e->host_r!=rope[s]||lengths[s]length){ + if(e->latent)cudaFree(e->latent);if(e->rope)cudaFree(e->rope); + e->latent=e->rope=nullptr;e->length=e->capacity=0; + e->K=K;e->R=R;e->host_l=latent[s];e->host_r=rope[s]; + } + if(lengths[s]>e->capacity){ + int cap=(lengths[s]+63)&~63;float *nl=nullptr,*nr=nullptr; + if(!cuda_ok(cudaMalloc(&nl,(size_t)cap*K*sizeof(float)),"ragged KV latent page")|| + !cuda_ok(cudaMalloc(&nr,(size_t)cap*R*sizeof(float)),"ragged KV rope page")){ + if(nl)cudaFree(nl);if(nr)cudaFree(nr);std::free(dl);std::free(dr);std::free(old);std::free(add);std::free(off);return 0; + } + if(e->length){ + cudaMemcpyAsync(nl,e->latent,(size_t)e->length*K*sizeof(float),cudaMemcpyDeviceToDevice,dc->stream); + cudaMemcpyAsync(nr,e->rope,(size_t)e->length*R*sizeof(float),cudaMemcpyDeviceToDevice,dc->stream); + } + if(e->latent)cudaFree(e->latent);if(e->rope)cudaFree(e->rope); + e->latent=nl;e->rope=nr;e->capacity=cap; + } + dl[s]=e->latent;dr[s]=e->rope;old[s]=e->length;add[s]=lengths[s]-e->length; + off[s]=packed_n;packed_n+=add[s]*(K+R); + } + size_t qb=(size_t)S*H*(Q+R)*sizeof(float); + size_t cb=(size_t)S*H*V*sizeof(float),ob=(size_t)S*proj->O*sizeof(float); + size_t pb=(size_t)packed_n*sizeof(float); + size_t desc=(size_t)S*(2*sizeof(float*)+4*sizeof(int)); + int ok=reserve(&dc->aq,&dc->aq_cap,qb)&&reserve(&dc->ac,&dc->ac_cap,cb)&& + reserve(&dc->y,&dc->y_cap,ob)&&reserve_bytes(&dc->group_desc,&dc->group_desc_cap,desc)&& + (!pb||(reserve(&dc->al,&dc->al_cap,pb)&&reserve_pinned(&dc->host_kv,&dc->host_kv_cap,pb))); + char *db=(char*)dc->group_desc;float **ddl=(float**)db,**ddr=ddl+S; + int *dn=(int*)(ddr+S),*dold=dn+S,*dadd=dold+S,*doff=dadd+S; + if(ok&&pb){ + for(int s=0;shost_kv+off[s]; + std::memcpy(p,latent[s]+(size_t)old[s]*K,(size_t)add[s]*K*sizeof(float)); + std::memcpy(p+(size_t)add[s]*K,rope[s]+(size_t)old[s]*R,(size_t)add[s]*R*sizeof(float)); + } + ok=cuda_ok(cudaMemcpyAsync(dc->al,dc->host_kv,pb,cudaMemcpyHostToDevice,dc->stream),"ragged KV append upload"); + } + if(ok)ok=cuda_ok(cudaMemcpyAsync(dc->aq,q,qb,cudaMemcpyHostToDevice,dc->stream),"ragged q upload")&& + cuda_ok(cudaMemcpyAsync(ddl,dl,(size_t)S*sizeof(float*),cudaMemcpyHostToDevice,dc->stream),"ragged latent pointers")&& + cuda_ok(cudaMemcpyAsync(ddr,dr,(size_t)S*sizeof(float*),cudaMemcpyHostToDevice,dc->stream),"ragged rope pointers")&& + cuda_ok(cudaMemcpyAsync(dn,lengths,(size_t)S*sizeof(int),cudaMemcpyHostToDevice,dc->stream),"ragged lengths upload")&& + cuda_ok(cudaMemcpyAsync(dold,old,(size_t)S*sizeof(int),cudaMemcpyHostToDevice,dc->stream),"ragged old lengths")&& + cuda_ok(cudaMemcpyAsync(dadd,add,(size_t)S*sizeof(int),cudaMemcpyHostToDevice,dc->stream),"ragged append lengths")&& + cuda_ok(cudaMemcpyAsync(doff,off,(size_t)S*sizeof(int),cudaMemcpyHostToDevice,dc->stream),"ragged append offsets"); + if(ok&&pb)ragged_kv_append<<stream>>>(ddl,ddr,dc->al,dold,dadd,doff,K,R); + if(ok)for(int s=0;sragged_count;i++)if(w->ragged[i].key==keys[s]){w->ragged[i].length=lengths[s];break;} + } + std::free(dl);std::free(dr);std::free(old);std::free(add);std::free(off);if(!ok)return 0; + size_t shared=(size_t)(2*K+T+256)*sizeof(float); + attention_absorb_ragged_kernel<<stream>>>(dc->ac,dc->aq,ddl,ddr, + dn,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale); + quant_matmul<<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)); + 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"); +} + extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) { if (!tensor) return; DeviceContext *ctx = find_ctx(tensor->device); @@ -806,6 +942,10 @@ extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) { } if (tensor->weights) cudaFree(tensor->weights); if (tensor->scales) cudaFree(tensor->scales); + for(int i=0;iragged_count;i++){ + if(tensor->ragged[i].latent)cudaFree(tensor->ragged[i].latent); + if(tensor->ragged[i].rope)cudaFree(tensor->ragged[i].rope); + } std::free(tensor); } diff --git a/c/backend_cuda.h b/c/backend_cuda.h index acbe4bc..63c1f11 100644 --- a/c/backend_cuda.h +++ b/c/backend_cuda.h @@ -14,6 +14,7 @@ #define COLI_CUDA_DLLEXPORT #endif + #ifdef __cplusplus extern "C" { #endif @@ -92,6 +93,11 @@ COLI_CUDA_DLLEXPORT int coli_cuda_attention_project_batch(ColiCudaTensor *kv_b,C const float *rope,int S,int H,int Q,int R, int V,int K,int T,float attention_scale); +COLI_CUDA_DLLEXPORT int coli_cuda_attention_project_ragged(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, + float *out,const float *q,const void *const *keys, + const float *const *latent,const float *const *rope, + const int *lengths,int S,int H,int Q,int R,int V,int K,int max_t,float attention_scale); + COLI_CUDA_DLLEXPORT void coli_cuda_tensor_free(ColiCudaTensor *tensor); COLI_CUDA_DLLEXPORT size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor); COLI_CUDA_DLLEXPORT int coli_cuda_tensor_device(const ColiCudaTensor *tensor); @@ -143,4 +149,3 @@ COLI_CUDA_DLLEXPORT int coli_cuda_pipe_sync(int device); #endif #endif - diff --git a/c/backend_loader.c b/c/backend_loader.c index b743d1b..eedbd50 100644 --- a/c/backend_loader.c +++ b/c/backend_loader.c @@ -61,6 +61,10 @@ typedef int (*fn_attention_absorb_batch)(ColiCudaTensor *kv_b,float *ctx,const f typedef int (*fn_attention_absorb_batch_dev)(ColiCudaTensor *kv_b_shard,float *ctx_dev, const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale); typedef int (*fn_attention_absorb_kvdev)(ColiCudaTensor *kv_b,float *ctx,const float *q, const float *latent_dev,const float *rope_dev,int H,int Q,int R,int V,int K,int T, float scale); typedef int (*fn_attention_project_batch)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q,const float *latent, const float *rope,int S,int H,int Q,int R, int V,int K,int T,float attention_scale); +typedef int (*fn_attention_project_ragged)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, + float *out,const float *q,const void *const *keys, + const float *const *latent,const float *const *rope, + const int *lengths,int S,int H,int Q,int R,int V,int K,int max_t,float attention_scale); typedef int (*fn_attention_project_batch_dev)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale); typedef int (*fn_attention_project_batch_dev_out)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out_dev,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale); typedef int (*fn_pipe_add)(int device,float *x_dev,const float *t_dev,size_t n); @@ -107,6 +111,7 @@ static struct { fn_attention_absorb_batch_dev attention_absorb_batch_dev; fn_attention_absorb_kvdev attention_absorb_kvdev; fn_attention_project_batch attention_project_batch; + fn_attention_project_ragged attention_project_ragged; fn_attention_project_batch_dev attention_project_batch_dev; fn_attention_project_batch_dev_out attention_project_batch_dev_out; fn_pipe_add pipe_add; @@ -200,6 +205,7 @@ static int coli_cuda_load(void){ RESOLVE(attention_absorb_batch_dev, fn_attention_absorb_batch_dev) RESOLVE(attention_absorb_kvdev, fn_attention_absorb_kvdev) RESOLVE(attention_project_batch, fn_attention_project_batch) + RESOLVE(attention_project_ragged, fn_attention_project_ragged) RESOLVE(attention_project_batch_dev, fn_attention_project_batch_dev) RESOLVE(attention_project_batch_dev_out, fn_attention_project_batch_dev_out) RESOLVE(pipe_add, fn_pipe_add) @@ -342,6 +348,15 @@ int coli_cuda_attention_project_batch(ColiCudaTensor *kv_b,ColiCudaTensor *o_pro return g_cuda.attention_project_batch(kv_b, o_proj, out, q, latent, rope, S, H, Q, R, V, K, T, attention_scale); } +int coli_cuda_attention_project_ragged(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, + float *out,const float *q,const void *const *keys, + const float *const *latent,const float *const *rope, + const int *lengths,int S,int H,int Q,int R,int V,int K,int max_t,float attention_scale){ + if(!coli_cuda_load()) return 0; + return g_cuda.attention_project_ragged(kv_b,o_proj,out,q,keys,latent,rope,lengths, + S,H,Q,R,V,K,max_t,attention_scale); +} + int coli_cuda_attention_project_batch_dev(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale){ if(!g_cuda.available){ return 0; } return g_cuda.attention_project_batch_dev(kv_b, o_proj, out, q_dev, latent_dev, rope_dev, S, H, Q, R, V, K, T, scale); diff --git a/c/coli b/c/coli index 0dd9de3..803c22b 100755 --- a/c/coli +++ b/c/coli @@ -40,6 +40,8 @@ if sys.platform == "win32": except (AttributeError, OSError): pass HERE = os.path.dirname(os.path.abspath(__file__)) +sys.path.insert(0, HERE) +from version import __version__ as _version # Run-in-place (source checkout, "cd c && ./coli ..."): the engine, the # support modules (resource_plan.py, doctor.py, openai_server.py) and @@ -115,7 +117,7 @@ def sprite_lines(): def banner(sub=""): sp=sprite_lines() txt=[ - f"{C.teal}{C.b}colibrì{C.r} {C.dim}v1.0{C.r}", + f"{C.teal}{C.b}colibrì{C.r} {C.dim}v{_version}{C.r}", f"{C.dim}tiny engine, immense model{C.r}", f"{C.gray}GLM-5.2 · 744B MoE · int4 · streaming CPU{C.r}", f"{C.dgray}{sub}{C.r}" if sub else "", @@ -141,12 +143,24 @@ def need_model(model): sys.exit(f"{C.yel}engine is not built.{C.r} Run: coli build") 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 - for line in linked.stdout.splitlines()) - except (OSError,subprocess.SubprocessError): return False + if not os.path.exists(GLM): return False + if sys.platform == "linux": + try: + linked=subprocess.run(["ldd",GLM],capture_output=True,text=True,timeout=3) + return any("libcudart" in line and "not found" not in line + for line in linked.stdout.splitlines()) + except (OSError,subprocess.SubprocessError): return False + if sys.platform == "win32": + # Windows CUDA_DLL=1 builds never link libcudart directly: glm.exe loads + # coli_cuda.dll at runtime via LoadLibrary (backend_loader.c), so there's no + # import-table entry for ldd/dumpbin to see. Detect the COLI_CUDA build via a + # marker string baked into glm.c's #ifdef COLI_CUDA block instead, and require + # coli_cuda.dll to actually sit next to glm.exe (else CUDA init fails at startup). + try: + with open(GLM,"rb") as f: built=b"[CUDA] mode: routed experts" in f.read() + except OSError: return False + return built and os.path.exists(os.path.join(os.path.dirname(GLM),"coli_cuda.dll")) + return False def resource_request(a, env): ctx=a.ctx or int(env.get("CTX",4096)) @@ -480,8 +494,11 @@ def cmd_run(a): need_model(a.model) prompt=" ".join(a.prompt) if a.prompt else sys.exit('usage: coli run "your prompt"') banner("run") - # template ufficiale GLM-5.2: niente \n dopo i ruoli; = risposta diretta (nothink) - e=env_for(a); e["PROMPT"]=f"[gMASK]<|user|>{prompt}<|assistant|>" + # template ufficiale GLM-5.2: niente \n dopo i ruoli; = risposta diretta (nothink). + # THINK=1 lascia aperto, stessa convenzione del serve mode (glm.c). EN: THINK=1 leaves + # open so the engine emits its reasoning block; the default stays nothink. + tk="" if os.environ.get("THINK","0")=="1" else "" + e=env_for(a); e["PROMPT"]=f"[gMASK]<|user|>{prompt}<|assistant|>{tk}" sys.exit(subprocess.call([GLM, str(a.cap)], env=e)) def server_probe(base, api_key=None, timeout=1.5): @@ -890,6 +907,7 @@ def main(): common.add_argument("--topp", type=float, default=0); common.add_argument("--topk", type=int, default=0) common.add_argument("--temp", type=float, default=None) # temperatura token (0=greedy, default 1.0+nucleus .95) ap=argparse.ArgumentParser(prog="coli", parents=[common], description="colibrì — run GLM-5.2 locally") + ap.add_argument("--version", action="version", version=f"colibrì {_version}") sub=ap.add_subparsers(dest="cmd") sub.add_parser("build", parents=[common]); sub.add_parser("info", parents=[common]) pp=sub.add_parser("plan",parents=[common]) diff --git a/c/compat.h b/c/compat.h index ce21df2..1b56f5f 100644 --- a/c/compat.h +++ b/c/compat.h @@ -80,6 +80,7 @@ static inline int compat_open_direct(const char *path){ #endif #include #include +#include /* _mkdir (for the mkdtemp shim below) */ #include #include #include @@ -313,8 +314,52 @@ static inline int compat_setenv(const char *name, const char *value, int overwri } #define setenv(name,value,overwrite) compat_setenv(name,value,overwrite) +/* --- getenv_utf8: read an env var as UTF-8, not through the ANSI codepage --- + * Plain getenv()/_environ are populated by the CRT from the ANSI-codepage view + * of the process environment block, not UTF-8. A parent that hands the child a + * Unicode value via CreateProcessW's wide env block (e.g. Python's subprocess + * module, which coli uses to pass the chat prompt) round-trips correctly only + * through GetEnvironmentVariableW; going through narrow getenv() re-encodes it + * via CP_ACP first, so any non-ASCII prompt text (Cyrillic, CJK, ...) comes out + * corrupted before the byte-level tokenizer ever sees it. Read the wide value + * directly and convert straight to UTF-8, bypassing the ANSI codepage entirely. + * Returned buffer is intentionally leaked: called a handful of times at + * startup, lives for the process. */ +static inline const char *compat_getenv_utf8(const char *name){ + wchar_t wname[64]; + if(MultiByteToWideChar(CP_UTF8, 0, name, -1, wname, 64) <= 0) return getenv(name); + DWORD need = GetEnvironmentVariableW(wname, NULL, 0); + if(!need) return NULL; + wchar_t *wval = (wchar_t*)malloc(need * sizeof(wchar_t)); + if(!wval) return NULL; + GetEnvironmentVariableW(wname, wval, need); + int blen = WideCharToMultiByte(CP_UTF8, 0, wval, -1, NULL, 0, NULL, NULL); + char *val = blen>0 ? (char*)malloc((size_t)blen) : NULL; + if(val) WideCharToMultiByte(CP_UTF8, 0, wval, -1, val, blen, NULL, NULL); + free(wval); + return val; +} +#define getenv_utf8(name) compat_getenv_utf8(name) + +/* --- mkdtemp -> _mktemp + _mkdir (POSIX mkdtemp assente su Windows) --- + * Test binaries (test_stops.c) create a scratch dir in the CWD via a + * "name_XXXXXX" template; POSIX mkdtemp fills the X's and mkdirs 0700. The + * Windows CRT has _mktemp (in-place, same XXXXXX contract) so we compose it. + * Returns the template pointer on success, NULL on failure — matching POSIX. */ +static inline char *compat_mkdtemp(char *tmpl){ + if(!tmpl) return NULL; + if(!_mktemp(tmpl)) return NULL; /* fills the trailing X's in place */ + if(_mkdir(tmpl) != 0) return NULL; /* EEXIST is impossible post-_mktemp */ + return tmpl; +} +#define mkdtemp(tmpl) compat_mkdtemp(tmpl) + #endif /* _WIN32 */ +#ifndef getenv_utf8 +#define getenv_utf8(name) getenv(name) +#endif + /* --- compat_aligned_free su piattaforme diverse da Windows --- * Su Linux/macOS, posix_memalign usa free() normale. */ #ifndef compat_aligned_free diff --git a/c/doctor.py b/c/doctor.py index 4cb5a42..cbe9c87 100644 --- a/c/doctor.py +++ b/c/doctor.py @@ -19,16 +19,33 @@ def _check(identifier, status, summary, **details): def cuda_linkage(engine_path): """Return CUDA linkage state without loading the executable or CUDA runtime.""" - if not Path(engine_path).is_file() or os.name != "posix": + engine = Path(engine_path) + if not engine.is_file(): return {"linked": False, "missing": False} - try: - result = subprocess.run(["ldd", str(engine_path)], capture_output=True, text=True, - timeout=3, check=False) - except (OSError, subprocess.SubprocessError): - return {"linked": False, "missing": False} - lines = [line for line in result.stdout.splitlines() if "libcudart" in line] - return {"linked": any("not found" not in line for line in lines), - "missing": any("not found" in line for line in lines)} + if os.name == "posix": + try: + result = subprocess.run(["ldd", str(engine)], capture_output=True, text=True, + timeout=3, check=False) + except (OSError, subprocess.SubprocessError): + return {"linked": False, "missing": False} + lines = [line for line in result.stdout.splitlines() if "libcudart" in line] + return {"linked": any("not found" not in line for line in lines), + "missing": any("not found" in line for line in lines)} + if sys.platform == "win32": + # Windows CUDA_DLL=1 builds never link libcudart directly: glm.exe loads + # coli_cuda.dll at runtime via LoadLibrary (backend_loader.c), so there's no + # import-table entry for ldd/dumpbin to see. Detect the COLI_CUDA build via a + # marker string baked into glm.c's #ifdef COLI_CUDA block instead, and require + # coli_cuda.dll to actually sit next to glm.exe (else CUDA init fails at startup). + try: + built = b"[CUDA] mode: routed experts" in engine.read_bytes() + except OSError: + return {"linked": False, "missing": False} + if not built: + return {"linked": False, "missing": False} + dll_present = (engine.parent / "coli_cuda.dll").is_file() + return {"linked": dll_present, "missing": not dll_present} + return {"linked": False, "missing": False} def run_doctor(model, ram_gb=0, context=4096, gpu_indices=None, vram_gb=0, *, diff --git a/c/glm.c b/c/glm.c index 92a4d78..a3e7ab7 100644 --- a/c/glm.c +++ b/c/glm.c @@ -196,12 +196,16 @@ typedef struct { uint64_t mtp_prop, mtp_acc; /* statistica acceptance */ int **eroute; int *enr; /* metodo C: routing dell'ULTIMO token per layer */ uint64_t eclock, hits, miss, ereq; + uint64_t hit_pin, hit_ecache; /* split di hits per tier (#336): pin vs LRU ecache */ uint64_t gpu_expert_calls; int gpu_expert_count; int64_t gpu_expert_bytes; uint64_t n_fw, n_emit; /* metodo E: forward di decode / token emessi */ uint64_t route_slots, route_swaps; /* CACHE_ROUTE: slots chosen / substituted vs true top-K */ uint64_t route_agree_hit, route_agree_tot; /* ROUTE_AGREE: |chosen ∩ true top-K| / K */ double route_kl_sum; uint64_t route_kl_n; /* mean KL(true||chosen) on gate mass */ - double t_ewait, t_emm, t_attn, t_kvb, t_head;/* profiling: dove va il tempo (wall del + double t_ewait, t_emm, t_ecpu, t_egpu, t_route, t_p2p, t_attn, t_kvb, t_head; + uint64_t n_p2p; /* P0 execution profile: tier split + residual hops */ + uint64_t cpu_expert_rows; int64_t cpu_expert_bytes; + /* profiling: dove va il tempo (wall del * thread di compute; il servizio disco * overlappato vive in g_edisk_ns) */ double t_aproj,t_acore,t_aout; /* attention breakdown */ @@ -307,15 +311,19 @@ static uint64_t g_prof_nlat; /* forwards recorded (monotonic static void prof_lat(double s){ g_prof_lat[g_prof_nlat++ % PROF_LAT_CAP]=s; } /* snapshot for windowed reports (serve mode: one report per turn) */ typedef struct { - double edisk,ewait,emm,attn,head; - int64_t io; uint64_t hits,miss,ereq,n_fw,n_emit,nlat; + double edisk,ewait,emm,ecpu,egpu,route,p2p,attn,head; + int64_t io,cpu_bytes; uint64_t hits,miss,ereq,n_fw,n_emit,nlat,n_p2p,cpu_rows; + uint64_t hit_pin,hit_ecache; } ProfBase; static void prof_base(Model *m, ProfBase *b){ b->edisk=edisk_s(); b->ewait=m->t_ewait; b->emm=m->t_emm; + b->ecpu=m->t_ecpu; b->egpu=m->t_egpu; b->route=m->t_route; b->p2p=m->t_p2p; b->attn=m->t_attn; b->head=m->t_head; b->io=atomic_load_explicit(&g_prof_io,memory_order_relaxed); b->hits=m->hits; b->miss=m->miss; b->ereq=m->ereq; - b->n_fw=m->n_fw; b->n_emit=m->n_emit; b->nlat=g_prof_nlat; + b->hit_pin=m->hit_pin; b->hit_ecache=m->hit_ecache; + b->n_fw=m->n_fw; b->n_emit=m->n_emit; b->nlat=g_prof_nlat; b->n_p2p=m->n_p2p; + b->cpu_bytes=m->cpu_expert_bytes;b->cpu_rows=m->cpu_expert_rows; } static float *falloc(int64_t n){ @@ -1649,7 +1657,7 @@ static void model_init(Model *m, const char *snap, int cap, int ebits, int dbits c->index_topk, c->index_topk); } } - m->hlast=falloc(D); m->h_all=falloc((int64_t)64*D); + m->hlast=falloc(D); m->h_all=falloc((int64_t)512*D); /* byte della parte DENSA residente (embed+lm_head+attn+mlp densa+shared+norme) */ int64_t rb=qt_bytes(&m->embed)+qt_bytes(&m->lm_head); @@ -2284,6 +2292,18 @@ static void expert_prefetch(Model *m, int layer, int eid){ static void qt_addrow(const QT *t, int row, float coef, float *acc){ int I=t->I; if(t->fmt==0){ const float *w=t->qf+(int64_t)row*I; for(int i=0;ifmt==4){ 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 i=0;i+1>1]; + acc[i] +=coef*scl[i/gs] *((int)(b&0xF)-8); + acc[i+1]+=coef*scl[(i+1)/gs]*((int)(b>>4)-8); } + if(I&1){ uint8_t b=w[I>>1]; acc[I-1]+=coef*scl[(I-1)/gs]*((int)(b&0xF)-8); } return; } float c=coef*t->s[row]; if(t->fmt==1){ const int8_t *w=t->q8+(int64_t)row*I; for(int i=0;ifmt==2){ const uint8_t *w=t->q4+(int64_t)row*((I+1)/2); @@ -2302,6 +2322,13 @@ static void qt_matvec_rows(const QT *t, int r0, int n, const float *x, float *y) 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; for(int i=0;i+1>1]; acc+=((int)(b&0xF)-8)*x[i]+((int)(b>>4)-8)*x[i+1]; } if(I&1){ uint8_t b=w[I>>1]; acc+=((int)(b&0xF)-8)*x[I-1]; } a=acc*s; } + else if(t->fmt==4){ /* per-gruppo, come matmul_i4_grouped / per-group, as matmul_i4_grouped */ + 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*gsI?I:base+gs; float acc=0; + for(int i=base;i>1]; + acc+=(float)((i&1)?((int)(b>>4)-8):((int)(b&0xF)-8))*x[i]; } + a+=(double)acc*scl[g]; } } else { const uint8_t *w=t->q4+(int64_t)row*((I+3)/4); float s=t->s[row]; float acc=0; for(int i=0;i>2]; acc+=((int)((b>>((i&3)*2))&3)-2)*x[i]; } a=acc*s; } y[j]=(float)a; @@ -2686,7 +2713,25 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p float *sc_all = falloc((int64_t)omp_get_max_threads()*sc_cap); int cuda_core=0,cuda_projected=0; #ifdef COLI_CUDA - if(cuda_absorb&&l->n_kv_b_shard>1){ + if(kvs&&g_cuda_enabled&&getenv("COLI_CUDA_ATTN")&&atoi(getenv("COLI_CUDA_ATTN"))&& + !dnsel&&l->kv_b.cuda_eligible&&l->o.cuda_eligible&& + qt_cuda_upload(&l->kv_b)&&qt_cuda_upload(&l->o)){ + const float **rl=malloc((size_t)S*sizeof(*rl)),**rr=malloc((size_t)S*sizeof(*rr)); + const void **rk=malloc((size_t)S*sizeof(*rk)); + int *rn=malloc((size_t)S*sizeof(*rn)); int mt=0; + if(rk&&rl&&rr&&rn){ + for(int s=0;skv_start[layer]; rn[s]=pos+1-st0; + rk[s]=kvs[s]; + rl[s]=coli_kv_row(kvs[s]->Lc[layer],st0,kvl); + rr[s]=coli_kv_row(kvs[s]->Rc[layer],st0,c->qk_rope); + if(rn[s]>mt)mt=rn[s]; + } + cuda_core=cuda_projected=coli_cuda_attention_project_ragged(l->kv_b.cuda,l->o.cuda, + out,Q,rk,rl,rr,rn,S,H,c->qk_nope,c->qk_rope,vh,kvl,mt,c->attn_scale); + } + free(rk);free(rl);free(rr);free(rn); + } else if(cuda_absorb&&l->n_kv_b_shard>1){ int n=l->n_kv_b_shard,st0=m->kv_start[layer],nt=pos_base+S-st0,ok=1; float *qs=falloc((int64_t)S*H*qh),*cs=falloc((int64_t)S*H*vh); for(int d=0;dt_route+=now_s()-route_t0; if(g_route_fp) g_route_call++; if(g_couple && cp_pred && S<=8) for(int s2=0;s2pin[layer]; - for(int z=0;znpin[layer];z++) if(P[z].eid==eid){ m->hits++; use[j]=&P[z]; break; } + for(int z=0;znpin[layer];z++) if(P[z].eid==eid){ m->hits++; m->hit_pin++; use[j]=&P[z]; break; } if(!use[j]){ ESlot *Sl=m->ecache[layer]; int nn=m->ecn[layer]; - for(int z=0;zhits++; Sl[z].used=(uint64_t)__atomic_add_fetch(&m->eclock,1,__ATOMIC_RELAXED); use[j]=&Sl[z]; break; } } + for(int z=0;zhits++; m->hit_ecache++; Sl[z].used=(uint64_t)__atomic_add_fetch(&m->eclock,1,__ATOMIC_RELAXED); use[j]=&Sl[z]; break; } } if(!use[j]){ qof[j]=nmiss; use[j]=&m->ws[nmiss]; missk[nmiss++]=j; m->miss++; if(g_disk_split){ if(m->ld_ctx==1) m->miss_draft++; else if(m->ld_ctx==2) m->miss_absorb++; } } } @@ -3263,7 +3310,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int coli_cuda_expert_mlp(e->g.cuda,e->u.cuda,e->d.cuda,hh,xg,nr)){ for(int r=0;rt_emm+=now_s()-t0; continue; + double dt=now_s()-t0;m->t_emm+=dt;if(g_prof)m->t_egpu+=dt;continue; } if(!e->slab) expert_host_ensure(m,layer,e); #endif @@ -3272,7 +3319,9 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int matmul_qt(hh, gg, &e->d, nr); for(int r=0;rt_emm += now_s()-t0; + double dt=now_s()-t0;m->t_emm+=dt;if(g_prof){m->t_ecpu+=dt; + m->cpu_expert_bytes+=qt_bytes(&e->g)+qt_bytes(&e->u)+qt_bytes(&e->d); + m->cpu_expert_rows+=(uint64_t)nr;} } #ifdef COLI_CUDA ColiCudaTensor *dev_g[COLI_CUDA_MAX_DEVICES][64],*dev_u[COLI_CUDA_MAX_DEVICES][64]; @@ -3280,6 +3329,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int int dev_rows[COLI_CUDA_MAX_DEVICES][64],dev_which[COLI_CUDA_MAX_DEVICES][64]; int dev_nc[COLI_CUDA_MAX_DEVICES]={0},dev_total[COLI_CUDA_MAX_DEVICES]={0}; int dev_off[COLI_CUDA_MAX_DEVICES]={0},dev_ok[COLI_CUDA_MAX_DEVICES]={0}; + double dev_time[COLI_CUDA_MAX_DEVICES]={0}; for(int di=0;dig.cuda_device==g_cuda_devices[di]) dev_total[di]+=group_n[q]; for(int di=1;di1) schedule(static) - for(int di=0;dig.cuda,e->u.cuda,e->d.cuda,hh,xg,nr)){ expert_host_ensure(m,layer,e); expert_gate_up(gg,uu,xg,&e->g,&e->u,nr); for(int64_t z=0;z<(int64_t)nr*I;z++) gg[z]=siluf(gg[z])*uu[z]; matmul_qt(hh,gg,&e->d,nr); + if(g_prof){m->cpu_expert_bytes+=qt_bytes(&e->g)+qt_bytes(&e->u)+qt_bytes(&e->d); + m->cpu_expert_rows+=(uint64_t)nr;} } + if(g_prof)m->t_ecpu+=now_s()-tc; } float *src=dev_ok[di]?group_y+(int64_t)off*D:hh; for(int r=0;rmx)mx=dev_time[di];m->t_egpu+=mx;} m->t_emm+=now_s()-tg; #endif /* No drain barrier: the per-expert pipe_wait(qof[j]) above (issued for every @@ -3925,7 +3983,11 @@ static void layers_forward_rows(Model *m, float *x, int S, int pos_base, float *dst=coli_cuda_pipe_scratch(dev,15,xb); if(dst){ if(x_dev_on<0) ok=coli_cuda_pipe_upload(dev,dst,x,xb); - else if(x_dev_on!=dev){ ok=coli_cuda_pipe_peer_copy(dev,dst,x_dev_on,x_dev,xb); } + else if(x_dev_on!=dev){ + double tp=g_prof?now_s():0; + ok=coli_cuda_pipe_peer_copy(dev,dst,x_dev_on,x_dev,xb); + if(g_prof){m->t_p2p+=now_s()-tp;m->n_p2p++;} + } else dst=x_dev; if(ok){ x_dev=dst; x_dev_on=dev; @@ -4021,7 +4083,7 @@ static float *step_all(Model *m, const int *ids, int S, int pos_base){ float *x=falloc((int64_t)S*D); for(int s=0;sh_all) memcpy(m->h_all, x, (int64_t)S*D*sizeof(float)); /* hidden di TUTTE le pos (S<=64) */ + if(m->h_all) memcpy(m->h_all, x, (int64_t)S*D*sizeof(float)); /* hidden di TUTTE le pos (S<=512) */ if(m->hlast) memcpy(m->hlast, x+(int64_t)(S-1)*D, D*sizeof(float)); float *lo=falloc((int64_t)S*c->vocab), *row=falloc(D); for(int s=0;sfinal_norm, D, c->eps); @@ -4034,8 +4096,8 @@ static float *step_all(Model *m, const int *ids, int S, int pos_base){ static float *step_decode_batch(Model *m, const DecodeRow *rows, int S){ Cfg *c=&m->c; int D=c->hidden; /* Ragged KV currently uses MLA absorption; the stack kernel is sized to 512. */ - if(!rows || S<1 || S>64 || c->kv_lora>512) return NULL; - KVState *kvs[64]; int positions[64]; + if(!rows || S<1 || S>512 || c->kv_lora>512) return NULL; + KVState *kvs[512]; int positions[512]; float *x=falloc((int64_t)S*D); for(int s=0;sLc || !rows[s].kv->Rc || !rows[s].kv->kv_start || @@ -4151,7 +4213,11 @@ static void mtp_absorb(Model *m, const int *next_ids, const float *x, int S, int } static inline int argmax_v(const float *lo, int V){ - int b=0; float bv=lo[0]; for(int i=1;ibv){bv=lo[i];b=i;} return b; + /* skip NaN (x==x is false for NaN) so a poisoned logit can't pin the argmax + * to index 0 — pick the max finite/+Inf entry instead. */ + int b=-1; float bv=-INFINITY; + for(int i=0;ibv){ bv=x; b=i; } } + return b<0?0:b; } /* ---- METODO F: draft grammaticale (#48) ---- @@ -4255,9 +4321,28 @@ static void topp_siftdown(int *h, int n, int i){ * la coda troncata va AZZERATA in g_pbuf (dist_sample la legge direttamente per id). */ static void dist_build(const float *lo, int V){ if(!g_pbuf){ g_pbuf=falloc(V); g_pidx=malloc(V*sizeof(int)); } - float mx=lo[0]; for(int i=1;imx) mx=lo[i]; + /* Un solo logit NaN/+Inf avvelenava tutto (#369): +Inf diventava mx, NaN/Inf-mx + * -> expf NaN -> s NaN -> ogni prob NaN -> dist_sample cade sul fallback + * `g_pbuf[i]>0` (NaN>0 e' falso ovunque) e ritorna 0 PER SEMPRE, in silenzio. + * Difesa: mx solo sui finiti; un logit non finito contribuisce prob 0; + * se la distribuzione degenera (tutti non finiti / somma non valida) si + * ripiega sull'argmax dei finiti e si avvisa UNA volta, mai in silenzio. */ + int mxi=-1; float mx=0; + for(int i=0;imx)){ mx=lo[i]; mxi=i; } double s=0; float invt=1.f/(g_temp>1e-4f?g_temp:1e-4f); - for(int i=0;i=0){ + for(int i=0;i=0)?mxi:0; /* mxi = argmax dei logit FINITI (robusto + * anche se lo[0] e' NaN, dove argmax_v fallirebbe) */ + for(int i=0;i0 && g_nuc<1.f){ for(int i=0;it_ewait,m->t_emm,m->t_attn,m->t_kvb,m->t_head,elapsed-accounted); printf("ATTENTION: projection/RoPE %.3fs | score-softmax-value %.3fs | output projection %.3fs\n", m->t_aproj,m->t_acore,m->t_aout); + if(g_prof)printf("P0-EXEC: routed CPU %.3fs / %.2f GB/s (%llu row) | routed GPU critical %.3fs | router %.3fs | residual P2P %.3fs / %llu hop | orchestration %.3fs\n", + m->t_ecpu,m->t_ecpu>0?m->cpu_expert_bytes/1e9/m->t_ecpu:0.0, + (unsigned long long)m->cpu_expert_rows,m->t_egpu,m->t_route,m->t_p2p,(unsigned long long)m->n_p2p, + elapsed-m->t_ewait-m->t_emm-m->t_attn-m->t_head-m->t_route-m->t_p2p>0? + elapsed-m->t_ewait-m->t_emm-m->t_attn-m->t_head-m->t_route-m->t_p2p:0); #ifdef COLI_METAL if(g_metal_enabled){ uint64_t ok=0,fb=0,ex=0; double su=0,gp=0,sc=0; coli_metal_moe_counts(&ok,&fb,&ex); coli_metal_moe_times(&su,&gp,&sc); @@ -4571,6 +4661,8 @@ static void profile_print(Model *m, double elapsed){ static void profile_reset(Model *m){ m->t_ewait=m->t_emm=m->t_attn=m->t_kvb=m->t_head=0; + m->t_ecpu=m->t_egpu=m->t_route=m->t_p2p=0;m->n_p2p=0; + m->cpu_expert_bytes=0;m->cpu_expert_rows=0; m->t_aproj=m->t_acore=m->t_aout=0; atomic_store_explicit(&g_edisk_ns,0,memory_order_relaxed); } @@ -4607,19 +4699,31 @@ static void prof_report(Model *m, const ProfBase *b, double elapsed, int tokens, for(int i=0;i<=c->n_layers;i++){ if(m->npin)pinned+=m->npin[i]; if(m->ecn)lru+=m->ecn[i]; } double io_w=m->t_ewait-b->ewait; /* stall the compute thread felt */ double io_svc=edisk_s()-b->edisk; /* read service on the loading threads (overlaps compute) */ + uint64_t dhp=m->hit_pin-b->hit_pin, dhe=m->hit_ecache-b->hit_ecache; /* split #336 */ fprintf(f,"[PROF] expert I/O: %.3f GB fetched (%.1f MB/token, %.2f GB/s over the run%s) | " - "hit %.1f%% (%llu hit / %llu load) | %.1f loads/token | %.1fs read service / %.1fs felt wait\n", + "hit %.1f%% (%llu pin + %llu lru / %llu load) | %.1f loads/token | %.1fs read service / %.1fs felt wait\n", io/1e9, tokens>0?io/1e6/tokens:0.0, io/1e9/elapsed, g_mmap?"; COLI_MMAP=1: page cache may serve part":"", - hitp,(unsigned long long)dh,(unsigned long long)dm, tokens>0?(double)dq/tokens:0.0, + hitp,(unsigned long long)dhp,(unsigned long long)dhe,(unsigned long long)dm, tokens>0?(double)dq/tokens:0.0, io_svc,io_w); fprintf(f,"[PROF] resident experts: %d pinned (%.1f GB) + %d in LRU (%.1f GB, cap %d/layer)\n", pinned,pinned*eb/1e9,lru,lru*eb/1e9,m->ecap); - double emm=m->t_emm-b->emm, attn=m->t_attn-b->attn, head=m->t_head-b->head; - double other=elapsed-io_w-emm-attn-head; if(other<0) other=0; + double emm=m->t_emm-b->emm, ecpu=m->t_ecpu-b->ecpu, egpu=m->t_egpu-b->egpu; + double route=m->t_route-b->route,p2p=m->t_p2p-b->p2p; + uint64_t np2p=m->n_p2p-b->n_p2p; + int64_t cpu_bytes=m->cpu_expert_bytes-b->cpu_bytes; + uint64_t cpu_rows=m->cpu_expert_rows-b->cpu_rows; + double attn=m->t_attn-b->attn, head=m->t_head-b->head; + double other=elapsed-io_w-emm-attn-head-route-p2p; if(other<0) other=0; double f_io=io_w/elapsed, f_emm=emm/elapsed, f_attn=attn/elapsed; fprintf(f,"[PROF] time shares: expert-I/O %.0f%% | expert-matmul %.0f%% | attention %.0f%% | lm_head %.0f%% | other %.0f%%\n", 100*f_io,100*f_emm,100*f_attn,100*head/elapsed,100*other/elapsed); + double slow=ecpu>egpu?ecpu:egpu,fast=ecpu0?cpu_bytes/1e9/ecpu:0.0,(unsigned long long)cpu_rows, + egpu,fast>1e-9?slow/fast:0.0,route,p2p,(unsigned long long)np2p, + np2p?p2p*1e3/np2p:0.0,other); if(f_io>=0.30){ fprintf(f,"[PROF] verdict: I/O-bound — %.0f%% of the time waits on expert reads (hit %.0f%%).",100*f_io,hitp); if(hitp<90) fprintf(f," More cache is the lever: raise RAM_GB (or add RAM)."); @@ -4646,7 +4750,7 @@ static void run_replay(Model *m, const int *full, int nfull, int np){ if(np<2||nfull<=np){ fprintf(stderr,"REPLAY requires a non-empty prompt and continuation\n"); return; } kv_alloc(m,nfull+2); float *logit=step(m,full,np-1,0); free(logit); - m->hits=m->miss=m->ereq=m->gpu_expert_calls=0; + m->hits=m->miss=m->ereq=m->gpu_expert_calls=0; m->hit_pin=m->hit_ecache=0; profile_reset(m); ProfBase pb; prof_base(m,&pb); double t0=now_s(); int steps=0; @@ -4696,7 +4800,7 @@ static void run_text(Model *m, const char *snap, const char *prompt, int ngen){ } prefill_t=now_s()-prefill_t; printf("PROFILO PREFILL (%.2fs):\n",prefill_t); profile_print(m,prefill_t); - m->hits=m->miss=m->ereq=m->gpu_expert_calls=0; + m->hits=m->miss=m->ereq=m->gpu_expert_calls=0; m->hit_pin=m->hit_ecache=0; m->n_emit=m->n_fw=0; g_last_repin=0; profile_reset(m); @@ -4709,8 +4813,9 @@ static void run_text(Model *m, const char *snap, const char *prompt, int ngen){ double tot=m->hits+m->miss; int nsp=0; for(int i=0;in_layers;i++) if(m->L[i].sparse) nsp++; printf("\n---\nprefill %d tokens in %.2fs | decode %d tokens in %.2fs (%.2f tok/s) | " - "expert hit rate %.1f%% | RSS %.2f GB", - np,prefill_t,produced,dt,produced/dt,tot?100.0*m->hits/tot:0.0,rss_gb()); + "expert hit rate %.1f%% (pin %.1f%% + lru %.1f%%) | RSS %.2f GB", /* split #336: quale tier serve gli hit */ + np,prefill_t,produced,dt,produced/dt,tot?100.0*m->hits/tot:0.0, + tot?100.0*m->hit_pin/tot:0.0, tot?100.0*m->hit_ecache/tot:0.0, rss_gb()); if(g_cache_route && m->route_slots) printf(" | swap %.1f%% (%llu/%llu)", 100.0*m->route_swaps/m->route_slots, @@ -4819,7 +4924,68 @@ static int repin_pick(Model *m, RepinCand *out, int maxc){ } return nb; } +/* ---- RSS GUARD (#403) ----------------------------------------------------- + * La proiezione di cap_for_ram e' una STIMA: sul GB10 (#403) le generazioni + * lunghe l'hanno sforata di ~40 GB (proiettato 74.4, reale 115.6 -> 3 kill del + * kernel). La run D dell'issue prova che un cap piu' basso CONTIENE la crescita: + * questa guardia lo fa da sola, sull'RSS MISURATO invece che sul proiettato. + * Al safe point (stessa sede di repin: nessun moe in volo), ogni ~16 token + * emessi: se l'RSS supera il budget, svuota gli slot LRU meno usati e abbassa + * ecap perche' non ricrescano. Gli slab sono >128KB (mmap'd da glibc): la free + * restituisce le pagine al kernel subito, quindi l'RSS scende davvero. + * Lo slot NON viene compattato via: resta al suo posto con eid=-1/used=0 (primo + * candidato al riuso), perche' con PILOT_REAL il worker tiene puntatori dentro + * ecache[] durante i suoi pread e uno spostamento li invaliderebbe; per lo + * stesso motivo gli slot eid<0 (riservati/in caricamento) non si toccano e la + * selezione avviene sotto g_pilot_mx. resident_bytes resta invariato: gli slot + * LRU non sono mai contati li' (solo pin e densa). + * EN: evict = free the slab in place (eid=-1, used=0, never compact: PILOT_REAL + * EN: holds pointers into ecache[] across its preads), skip eid<0 reservations, + * EN: select under g_pilot_mx. RSS_GUARD_GB= forces an explicit ceiling. */ +static double g_ram_budget_gb=0; /* budget risolto, scritto da cap_for_ram */ +static uint64_t g_rssg_last=0; +static void rss_guard(Model *m){ + double lim = getenv("RSS_GUARD_GB") ? atof(getenv("RSS_GUARD_GB")) : g_ram_budget_gb; + if(lim<=0) return; + if(m->n_emit - g_rssg_last < 16) return; + g_rssg_last = m->n_emit; + double rss=rss_gb(); + if(rss <= lim*1.02+0.3) return; /* tolleranza: 2% + 300MB */ + Cfg *c=&m->c; + int64_t need=(int64_t)((rss-lim)*1e9), freed=0; int dropped=0; + for(int pass=0; pass<8 && freedn_layers && freedecache || !m->ecache[l]) continue; + pthread_mutex_lock(&g_pilot_mx); + int nn=m->ecn[l], lru=-1; + for(int z=0;zecache[l][z]; + if(cand->eid<0 || !cand->slab) continue; + if(lru<0 || cand->usedecache[l][lru].used) lru=z; + } + if(lru<0){ pthread_mutex_unlock(&g_pilot_mx); continue; } + ESlot *s=&m->ecache[l][lru]; + s->eid=-1; /* nascosto: nessun hit/evict altrui */ + pthread_mutex_unlock(&g_pilot_mx); + int64_t sb=s->slab_cap + s->fslab_cap*4; +#ifdef COLI_METAL + if(s->slab && g_metal_enabled) coli_metal_unregister(s->slab); +#endif + compat_aligned_free(s->slab); free(s->fslab); + s->slab=NULL; s->fslab=NULL; s->slab_cap=s->fslab_cap=0; + QT *q[3]={&s->g,&s->u,&s->d}; + for(int k=0;k<3;k++){ q[k]->qf=NULL; q[k]->q8=NULL; q[k]->q4=NULL; q[k]->s=NULL; } + s->used=0; /* primo candidato al riuso */ + freed += sb; dropped++; + } + if(m->ecap>2) m->ecap--; /* il tetto scende: niente ricrescita */ + } + if(dropped) + fprintf(stderr,"[RAM-GUARD] RSS %.1f GB over the %.1f GB budget (#403): " + "dropped %d cached experts, cap -> %d\n", rss, lim, dropped, m->ecap); +} static void repin_pass_limit(Model *m,int limit){ + rss_guard(m); /* #403: il budget si fa rispettare sull'RSS MISURATO */ if(g_repin<=0) return; if(m->n_emit - g_last_repin < (uint64_t)g_repin) return; g_last_repin = m->n_emit; @@ -5168,7 +5334,7 @@ static void run_serve_mux(Model *m, const char *snap){ g_draft=0; /* one scheduler owns every forward; MTP/speculation is not ragged-safe */ int maxctx=getenv("CTX")?atoi(getenv("CTX")):4096; int nctx=getenv("KV_SLOTS")?atoi(getenv("KV_SLOTS")):1; - if(nctx<1||nctx>16){fprintf(stderr,"KV_SLOTS deve essere tra 1 e 16\n");exit(2);} + if(nctx<1||nctx>512){fprintf(stderr,"KV_SLOTS must be between 1 and 512\n");exit(2);} g_kvsave=getenv("KVSAVE")?atoi(getenv("KVSAVE")):1; KVState *initial=m->kv; free(initial->kv_start); free(initial); ServeCtx *ctx=calloc(nctx,sizeof(*ctx)); ServeReq *req=calloc(nctx,sizeof(*req)); @@ -5224,7 +5390,7 @@ static void run_serve_mux(Model *m, const char *snap){ } active=0; for(int i=0;i1?atoi(argv[1]):64; int ebits= argc>2?atoi(argv[2]):8; int dbits= argc>3?atoi(argv[3]):ebits; - if(getenv("SERVE") && (kv_slot_count()<1 || kv_slot_count()>16)){ - fprintf(stderr,"KV_SLOTS must be between 1 and 16\n"); return 2; + int kv_limit=(getenv("SERVE_BATCH")&&atoi(getenv("SERVE_BATCH")))?512:16; + if(getenv("SERVE") && (kv_slot_count()<1 || kv_slot_count()>kv_limit)){ + fprintf(stderr,"KV_SLOTS must be between 1 and %d\n",kv_limit); return 2; } #ifdef COLI_CUDA if(getenv("COLI_CUDA") && atoi(getenv("COLI_CUDA"))){ @@ -6303,11 +6471,22 @@ int main(int argc, char **argv){ #endif } if(getenv("DSA_TOPK")) m.c.index_topk=atoi(getenv("DSA_TOPK")); /* override per test */ + /* Il path MUX (SERVE_BATCH=1, cioe' `coli serve`) forza g_draft=0 sotto — + * la speculazione non e' ragged-safe nel batch multi-slot. Segnalarlo QUI, + * altrimenti "MTP active (draft=8)" mentirebbe: il messaggio e' stampato + * prima della scelta del path (run_serve_mux, sotto), e con DRAFT=8 diceva + * "active" per poi disabilitarlo in silenzio (#358, LordMZTE). */ + int mux_will_disable_mtp = getenv("SERVE") && getenv("SERVE_BATCH") && atoi(getenv("SERVE_BATCH")); + int eff_draft = mux_will_disable_mtp ? 0 : g_draft; printf("loaded in %.2fs | resident dense: %.2f MB | layers=%d experts=%d | MTP %s (draft=%d)\n", now_s()-t0, m.resident_bytes/(1024.0*1024.0), m.c.n_layers, m.c.n_experts, - m.has_mtp?"ACTIVE":"absent", g_draft); + m.has_mtp?(mux_will_disable_mtp?"DISABLED (multiplexed serve)":"ACTIVE"):"absent", eff_draft); /* anche su stderr: e' il canale che le UI (coli) mostrano all'utente */ - fprintf(stderr,"[MTP] %s (draft=%d)\n", m.has_mtp?"active: native speculative decoding":"absent", g_draft); + if(mux_will_disable_mtp && m.has_mtp) + fprintf(stderr,"[MTP] disabled in multiplexed serve (SERVE_BATCH=1): speculation is not " + "ragged-safe across KV slots. Single-client interactive use (`coli chat`) keeps MTP.\n"); + else + fprintf(stderr,"[MTP] %s (draft=%d)\n", m.has_mtp?"active: native speculative decoding":"absent", eff_draft); #ifdef __linux__ { /* Only warn for a GENUINE 9p mount (WSL Windows drives, magic 0x01021997), where * fadvise is a no-op. The old check was `snap` starting with "/mnt/", which @@ -6457,8 +6636,9 @@ int main(int argc, char **argv){ double tot=m.hits+m.miss; printf("N-gram speculation (DRAFT=%d): %.2f tokens/forward (%llu forwards per %llu tokens)\n", g_draft, m.n_fw?(double)m.n_emit/m.n_fw:1.0, (unsigned long long)m.n_fw, (unsigned long long)m.n_emit); - printf("Expert cache hit rate: %.1f%% (hit=%llu miss=%llu) | RSS: %.2f GB | %.1f tok/s\n", - tot?100.0*m.hits/tot:0.0, (unsigned long long)m.hits, (unsigned long long)m.miss, rss_gb(), n_new/dt); + printf("Expert cache hit rate: %.1f%% (%llu pin + %llu lru / %llu miss) | RSS: %.2f GB | %.1f tok/s\n", + tot?100.0*m.hits/tot:0.0, (unsigned long long)m.hit_pin, (unsigned long long)m.hit_ecache, + (unsigned long long)m.miss, rss_gb(), n_new/dt); profile_print(&m,dt); if(g_prof) prof_report(&m,&pb,dt,n_new,stdout); #ifdef COLI_CUDA diff --git a/c/openai_server.py b/c/openai_server.py index d55fbc8..3900ba6 100644 --- a/c/openai_server.py +++ b/c/openai_server.py @@ -271,6 +271,14 @@ def parse_tool_calls(reply, tools=None): salvaged.append(name) calls.append({"id": "call_" + uuid.uuid4().hex[:24], "type": "function", "function": {"name": name, "arguments": json.dumps(args, ensure_ascii=False)}}) + if tools and not calls and re.search(r"||", reply): + # Diagnosi per la #401: il client ha dichiarato i tools e il modello ha PROVATO la + # sintassi, ma il parse rigoroso non ha agganciato nulla (tipico output int4 storpiato). + # EN: #401 field diagnosis: tools were declared and the model attempted the syntax, + # EN: but the strict parse matched nothing (typically quantization-mangled output). + sys.stderr.write("[api] tools declared and tool-call markers present, but no call " + "parsed -- output may be quantization-mangled; try COLI_TOOL_SALVAGE=1\n") + sys.stderr.flush() text = _BOX_RE.sub("", reply) if THINK_CLOSE in text: text = text.split(THINK_CLOSE, 1)[1] @@ -406,7 +414,10 @@ def generation_options(body, limit): maximum = body.get("max_tokens") maximum_param = "max_tokens" if maximum is None: - maximum = min(256, limit) + # Client omitted max_tokens: honor the operator's configured budget (--max-tokens / + # --ngen), not an arbitrary 256 — `coli serve --ngen 32768` must mean 32768 (#382). + # Generation still ends at EOS, so this is a cap, not a target. + maximum = limit temperature = body.get("temperature") top_p = body.get("top_p") temperature = 0.7 if temperature is None else temperature diff --git a/c/resource_plan.py b/c/resource_plan.py index 6b3f2fe..5f17a6a 100644 --- a/c/resource_plan.py +++ b/c/resource_plan.py @@ -203,6 +203,22 @@ def physical_cpu_count(): return os.cpu_count() or 1 +def cpu_socket_count(): + """Return the number of physical CPU sockets visible to this process.""" + if not sys.platform.startswith("linux"): + return 1 + try: + result = subprocess.run(["lscpu", "-p=socket"], text=True, + capture_output=True, check=True, timeout=5) + sockets = {int(line) for line in result.stdout.splitlines() + if line and not line.startswith("#")} + if sockets: + return len(sockets) + except (OSError, ValueError, subprocess.SubprocessError): + pass + return 1 + + POLICIES = { "quality": {"preserve_quantization": True, "preserve_router": True}, "balanced": {"preserve_quantization": True, "preserve_router": True}, @@ -212,11 +228,12 @@ POLICIES = { def build_plan(model, ram_gb=0, context=4096, gpu_indices=None, vram_gb=0, available_memory=None, available_disk=None, gpus=None, - policy="quality", physical_cpus=None): + policy="quality", physical_cpus=None, cpu_sockets=None): if policy not in POLICIES: raise ValueError(f"unknown policy: {policy}") info = analyze_model(model) physical_cpus = physical_cpu_count() if physical_cpus is None else physical_cpus + cpu_sockets = cpu_socket_count() if cpu_sockets is None else cpu_sockets cfg = info["config"] available_memory = memory_available() if available_memory is None else available_memory if available_disk is None: @@ -286,6 +303,7 @@ def build_plan(model, ram_gb=0, context=4096, gpu_indices=None, vram_gb=0, "quality_preserving": policy != "experimental-fast"}, "model": {key: value for key, value in info.items() if key != "config"}, "cpu": {"physical_cores": max(1, int(physical_cpus)), + "sockets": max(1, int(cpu_sockets)), "thread_policy": "physical-cores"}, "tiers": { "disk": {"role": "cold-backing", "model_bytes": info["model_bytes"], @@ -318,6 +336,10 @@ def environment_for_plan(plan, env=None, cuda_enabled=True): # ("Affinity not supported on this configuration"): non impostarle li'. result.setdefault("OMP_PROC_BIND", "spread") result.setdefault("OMP_PLACES", "cores") + if sys.platform.startswith("linux") and plan["cpu"].get("sockets", 1) > 1: + # Selectively interleave large expert/dense slabs across memory controllers. + # Unlike blanket numactl interleave, this leaves CUDA staging buffers local. + result.setdefault("COLI_NUMA", "1") if plan["policy"]["name"] == "balanced": result.setdefault("REPIN", "64") ram = plan["tiers"]["ram"] diff --git a/c/tests/bench_dsa_select.c b/c/tests/bench_dsa_select.c index 1fd3e6b..a1f2014 100644 --- a/c/tests/bench_dsa_select.c +++ b/c/tests/bench_dsa_select.c @@ -59,6 +59,12 @@ static double bench_ns(void (*fn)(const float*,int,int,int*,int*), return ts[N_REPEAT/2]; } +/* Sort an array of doubles ascending (median-of-medians aggregation below). */ +static void dsort(double *a, int n){ + for(int s=1;s=0 && a[b]>k){ a[b+1]=a[b]; b--; } a[b+1]=k; } +} + /* the NEW algorithm calls the real partial_select_desc + replicates the production * threshold derivation and position scans. */ static void keep_new(const float *isc, int nk, int keep, int *dst, int *nd_out){ @@ -74,6 +80,7 @@ static void keep_new(const float *isc, int nk, int keep, int *dst, int *nd_out){ /* deterministic score fill for three shapes */ static uint32_t brng = 0xA5A5A5A5u; +static void brng_seed(uint32_t s){ brng = s; } static double brand(void){ brng ^= brng << 13; brng ^= brng >> 17; brng ^= brng << 5; return (double)(brng >> 8) * (1.0 / 16777216.0); } static void fill_realistic(float *isc, int nk){ /* few hot, long distinct tail */ @@ -87,11 +94,24 @@ static void fill_plateau(float *isc, int nk){ /* tie blocks -> boundary path for(int i=0;i=u and fell + * through to return token 0 — so the engine silently emitted an unbroken run of + * token 0 with no error, on the DEFAULT serve path (TEMP>0, 0 +#include +#define main coli_glm_main_unused +#include "../glm.c" +#undef main + +static int approx1(double x){ return x > 0.999 && x < 1.001; } + +int main(void){ + /* --- argmax_v must skip NaN (greedy decode + speculative-verify paths) --- */ + { float lo[8]={NAN,1.f,5.f,2.f,NAN,-3.f,4.f,0.f}; + assert(argmax_v(lo,8)==2 && "pick max finite (idx2=5.0), not NaN-pinned idx0"); } + { float lo[8]={3.f,INFINITY,1.f,2.f,0.f,-1.f,2.5f,1.5f}; + assert(argmax_v(lo,8)==1 && "pick the +Inf position"); } + { float lo[8]; for(int i=0;i<8;i++) lo[i]=NAN; + assert(argmax_v(lo,8)==0 && "all-NaN: no crash, defined fallback"); } + + g_temp=0.7f; g_nuc=0.9f; /* the default serve/chat sampling path */ + + /* --- dist_build: a NaN logit must yield a finite one-hot, not all-NaN --- */ + { float lo[8]={0.5f,1.f,NAN,8.f,0.2f,-1.f,0.f,0.3f}; /* max finite = idx3 (8.0) */ + dist_build(lo,8); + double sum=0; int nan=0; + for(int i=0;i<8;i++){ if(!(g_pbuf[i]==g_pbuf[i])) nan=1; sum+=g_pbuf[i]; } + assert(!nan && "g_pbuf must be finite after a NaN logit"); + assert(approx1(sum) && "g_pbuf must normalize to 1"); + assert(approx1(g_pbuf[3]) && "mass must land on the max finite logit (idx3)"); + assert(dist_sample(8,-1)==3 && "sampler emits the finite argmax, not token 0"); } + + /* --- NaN at index 0: poisons the max scan itself (the old mx=lo[0] seed), + * the failure mode that starts before the sum (review note on #369) --- */ + { float lo[8]={NAN,1.f,0.5f,6.f,0.2f,-1.f,0.f,0.3f}; /* max finite = idx3 (6.0) */ + dist_build(lo,8); + double sum=0; int nan=0; + for(int i=0;i<8;i++){ if(!(g_pbuf[i]==g_pbuf[i])) nan=1; sum+=g_pbuf[i]; } + assert(!nan && "NaN at lo[0] must not poison via the mx seed"); + assert(approx1(sum) && "still normalizes to 1"); + assert(dist_sample(8,-1)==3 && "emits the max finite logit, not token 0"); } + + /* --- all-NaN logits: worst case — must stay finite, no crash --- */ + { float lo[8]; for(int i=0;i<8;i++) lo[i]=NAN; + dist_build(lo,8); + for(int i=0;i<8;i++) assert(g_pbuf[i]==g_pbuf[i] && "all-NaN: g_pbuf stays finite"); } + + /* --- regression: clean logits still produce a valid distribution --- */ + { float lo[8]={0.1f,0.2f,3.0f,0.4f,0.5f,0.6f,0.7f,0.8f}; /* peak = idx2 */ + dist_build(lo,8); + double sum=0; int nan=0; + for(int i=0;i<8;i++){ if(!(g_pbuf[i]==g_pbuf[i])) nan=1; sum+=g_pbuf[i]; } + assert(!nan && "clean softmax stays finite"); + assert(approx1(sum) && "clean softmax must sum to 1"); + assert(g_pbuf[2]>=g_pbuf[0] && "peak token keeps the most mass"); } + + printf("OK test_logit_nan: argmax_v NaN-skip + dist_build finite-collapse\n"); + return 0; +} diff --git a/c/tests/test_openai_tools_e2e.py b/c/tests/test_openai_tools_e2e.py new file mode 100644 index 0000000..d921843 --- /dev/null +++ b/c/tests/test_openai_tools_e2e.py @@ -0,0 +1,182 @@ +"""End-to-end tool-calling test for the OpenAI gateway (#401). + +Unlike the unit tests in test_openai_server.py (which call parse_tool_calls / +render_chat directly), this suite runs openai_server.py as a real subprocess +against a mock engine that speaks the actual SERVE wire protocol +(READY / SUBMIT / DATA / DONE), then talks to it over real HTTP. It pins down +the full path a coding client exercises: tool declaration rendering, marker +suppression in streamed deltas (across chunk boundaries), tool_calls in both +response shapes, and the <|observation|> round trip. +""" +import json +import os +import socket +import subprocess +import sys +import tempfile +import unittest +import urllib.request +from pathlib import Path + +SERVER = Path(__file__).resolve().parent.parent / "openai_server.py" +MODEL_ID = "glm-5.2-colibri" + +# Mock engine: replies are keyed on the prompt so one process covers every case. +# Prompts received are appended to MOCK_LOG for assertions on the rendering. +MOCK_ENGINE = r'''#!/usr/bin/env python3 +import sys, os +out, inp = sys.stdout.buffer, sys.stdin.buffer +out.write(b"\x01\x01READY\x01\x01\n" + b"STAT 0 0 0 0 0\n"); out.flush() + +def reply(rid, text, chunks=1): + data = text.encode("utf-8") + n = max(1, len(data) // chunks) + for i in range(0, len(data), n): + part = data[i:i+n] + out.write(("DATA %s %d\n" % (rid, len(part))).encode() + part + b"\n"); out.flush() + out.write(("DONE %s STAT %d 1.0 50.0 10.0 42 0\n" % (rid, len(text.split()))).encode()) + out.flush() + +while True: + line = inp.readline() + if not line: break + f = line.decode().strip().split() + if not f or f[0] != "SUBMIT": continue + rid, plen = f[1], int(f[3]) + prompt = inp.read(plen).decode("utf-8", "replace"); inp.read(1) + with open(os.environ["MOCK_LOG"], "a") as log: + log.write(prompt + "\n\x00\n") + if "" in prompt: + reply(rid, "25 degrees and sunny in Rome.") + elif "weather in Rome" in prompt: + reply(rid, "get_weathercity" + "Rome") + elif "weather in Milan" in prompt: + # split across many tiny DATA chunks: streamed marker suppression must + # hold even when a marker straddles a chunk boundary + reply(rid, "Checking. get_weathercity" + "Milan", chunks=20) + else: + reply(rid, "Hello from the mock engine.") +''' + +TOOLS = [{"type": "function", "function": { + "name": "get_weather", + "description": "Current weather for a city", + "parameters": {"type": "object", + "properties": {"city": {"type": "string"}}, + "required": ["city"]}}}] + + +@unittest.skipUnless(os.name == "posix", + "the mock engine is a shebang script the gateway execs directly; " + "Windows CreateProcess cannot run it. The gateway logic under test " + "is platform-independent and covered by the POSIX CI jobs.") +class ToolCallingE2E(unittest.TestCase): + @classmethod + def setUpClass(cls): + cls.tmp = tempfile.TemporaryDirectory() + mock = Path(cls.tmp.name) / "mock_engine.py" + mock.write_text(MOCK_ENGINE) + mock.chmod(0o755) + cls.mock_log = Path(cls.tmp.name) / "prompts.log" + cls.mock_log.touch() + with socket.socket() as probe: # free port, then hand it to the server + probe.bind(("127.0.0.1", 0)) + cls.port = probe.getsockname()[1] + env = dict(os.environ, MOCK_LOG=str(cls.mock_log)) + cls.server = subprocess.Popen( + [sys.executable, str(SERVER), "--model", cls.tmp.name, + "--engine", str(mock), "--port", str(cls.port)], + env=env, stderr=subprocess.DEVNULL) + cls.base = f"http://127.0.0.1:{cls.port}/v1" + for _ in range(100): + try: + urllib.request.urlopen(cls.base + "/models", timeout=2) + return + except OSError: + if cls.server.poll() is not None: + raise RuntimeError("gateway exited during startup") + import time + time.sleep(0.1) + raise RuntimeError("gateway did not come up") + + @classmethod + def tearDownClass(cls): + cls.server.terminate() + cls.server.wait(timeout=5) + cls.tmp.cleanup() + + def post(self, body, stream=False): + req = urllib.request.Request( + self.base + "/chat/completions", json.dumps(body).encode(), + {"Content-Type": "application/json"}) + resp = urllib.request.urlopen(req, timeout=30) + if not stream: + return json.loads(resp.read()) + events = [] + for raw in resp: + line = raw.decode().strip() + if line.startswith("data: ") and line != "data: [DONE]": + events.append(json.loads(line[6:])) + return events + + def test_tool_call_non_stream(self): + r = self.post({"model": MODEL_ID, "tools": TOOLS, + "messages": [{"role": "user", + "content": "What is the weather in Rome?"}]}) + choice = r["choices"][0] + self.assertEqual(choice["finish_reason"], "tool_calls") + calls = choice["message"]["tool_calls"] + self.assertEqual(len(calls), 1) + self.assertEqual(calls[0]["function"]["name"], "get_weather") + self.assertEqual(json.loads(calls[0]["function"]["arguments"]), {"city": "Rome"}) + self.assertNotIn("", choice["message"].get("content") or "") + + def test_tool_call_streamed_markers_suppressed(self): + events = self.post({"model": MODEL_ID, "tools": TOOLS, "stream": True, + "messages": [{"role": "user", + "content": "What is the weather in Milan?"}]}, + stream=True) + deltas = [e["choices"][0]["delta"] for e in events if e["choices"]] + text = "".join(d.get("content") or "" for d in deltas) + self.assertNotIn("", text) + self.assertNotIn("", text) + calls = [d["tool_calls"] for d in deltas if d.get("tool_calls")] + self.assertEqual(len(calls), 1) + self.assertEqual(calls[0][0]["function"]["name"], "get_weather") + self.assertEqual(json.loads(calls[0][0]["function"]["arguments"]), + {"city": "Milan"}) + finish = [e["choices"][0]["finish_reason"] for e in events + if e["choices"] and e["choices"][0].get("finish_reason")] + self.assertEqual(finish, ["tool_calls"]) + + def test_tool_result_round_trip(self): + r = self.post({"model": MODEL_ID, "tools": TOOLS, "messages": [ + {"role": "user", "content": "What is the weather in Rome?"}, + {"role": "assistant", "content": None, "tool_calls": [ + {"id": "call_x", "type": "function", + "function": {"name": "get_weather", + "arguments": "{\"city\": \"Rome\"}"}}]}, + {"role": "tool", "tool_call_id": "call_x", + "content": "25 degrees, sunny"}]}) + choice = r["choices"][0] + self.assertEqual(choice["finish_reason"], "stop") + self.assertFalse(choice["message"].get("tool_calls")) + self.assertIn("25 degrees", choice["message"]["content"]) + rendered = self.mock_log.read_text().split("\x00")[-2] + self.assertIn("<|observation|>25 degrees, sunny", + rendered) + self.assertIn("# Tools", rendered) + self.assertIn('"get_weather"', rendered) + + def test_no_tools_plain_text(self): + r = self.post({"model": MODEL_ID, + "messages": [{"role": "user", "content": "Hi!"}]}) + choice = r["choices"][0] + self.assertEqual(choice["finish_reason"], "stop") + self.assertIn("mock engine", choice["message"]["content"]) + + +if __name__ == "__main__": + unittest.main() diff --git a/c/tests/test_ragged_attention.cu b/c/tests/test_ragged_attention.cu new file mode 100644 index 0000000..dd1e2f8 --- /dev/null +++ b/c/tests/test_ragged_attention.cu @@ -0,0 +1,39 @@ +#include "../backend_cuda.h" + +#include +#include +#include + +int main(){ + int dev=0;if(!coli_cuda_init(&dev,1))return 77; + constexpr int S=3,H=2,Q=2,R=1,V=2,K=3,D=H*V,O=3,T=3; + std::vector w(H*(Q+V)*K),p(O*D),q(S*H*(Q+R)); + for(size_t i=0;i> l(S),r(S); + const float *lp[S],*rp[S]; + const void *keys[S]; + for(int s=0;s0` e' falso su NaN ovunque). + * Ora la distribuzione degenere ripiega sull'argmax dei logit FINITI e avvisa una volta. + * + * Il test verifica: (a) con logit sani il campionamento resta corretto; (b) con NaN/+Inf + * iniettato il token scelto e' l'argmax dei FINITI (mai 0 per default), su ogni posizione + * del NaN inclusa lo[0]; (c) nessun NaN sopravvive in g_pbuf. */ +#define main coli_glm_main_unused +#include "../glm.c" +#undef main +#include + +static uint32_t rs=0x1234abcd; static uint32_t xr(){rs^=rs<<13;rs^=rs>>17;rs^=rs<<5;return rs;} + +static int pbuf_has_nan(int V){ for(int i=0;i 0 e nessun NaN nel buffer */ + for(int i=0;ipv){pv=g_pbuf[i];picked=i;} + if(b<2 && picked!=amax){ /* NaN e +Inf: delta esatto su amax */ + printf(" FAIL: %s @ %d -> picked %d, atteso argmax finito %d\n",bn[b],at,picked,amax); fail=1; + } + } + } + if(!fail) printf(" NaN/+Inf iniettato: argmax dei finiti vince, mai 0/NaN ok\n"); + + /* (c) caso estremo: TUTTI non finiti -> non deve crashare, buffer valido */ + for(int i=0;i tuple[float, list[float]]: @@ -30,11 +36,20 @@ def parse_output(stdout: str, stderr: str = "") -> tuple[float, list[float]]: return float(speed.group(1)), [disk] + [float(value) for value in rest] -def execute(engine: str, env: dict[str, str]) -> tuple[float, list[float]]: +def parse_p0(stdout: str) -> list[float]: + """Extract the optional PROF=1 execution-layer breakdown.""" + row = P0_RE.search(stdout) + if not row: + raise RuntimeError("benchmark output missing P0-EXEC profile") + return [float(value) for value in row.groups()] + + +def execute(engine: str, env: dict[str, str]) -> tuple[float, list[float], list[float]]: run = subprocess.run( [engine, "4", "4", "4"], env=env, text=True, capture_output=True, check=True ) - return parse_output(run.stdout, run.stderr) + speed, profile = parse_output(run.stdout, run.stderr) + return speed, profile, parse_p0(run.stdout) def main() -> None: @@ -63,6 +78,8 @@ def main() -> None: OMP_NUM_THREADS=str(args.threads), OMP_PROC_BIND="spread", OMP_PLACES="cores", + DRAFT="0", + PROF="1", ) execute(args.engine, base | {"STATS": str(stats)}) @@ -86,13 +103,15 @@ def main() -> None: execute(args.engine, base | extra) # warm-up speeds = {name: [] for name in modes} profiles = {name: [] for name in modes} + p0_profiles = {name: [] for name in modes} names = list(modes) for run_index in range(args.runs): order = names[run_index % len(names):] + names[:run_index % len(names)] for name in order: - speed, profile = execute(args.engine, base | modes[name]) + speed, profile, p0 = execute(args.engine, base | modes[name]) speeds[name].append(speed) profiles[name].append(profile) + p0_profiles[name].append(p0) result = {} for name in names: @@ -103,6 +122,10 @@ def main() -> None: key: statistics.median(row[index] for row in profiles[name]) for index, key in enumerate(PROFILE_KEYS) }, + "median_p0": { + key: statistics.median(row[index] for row in p0_profiles[name]) + for index, key in enumerate(P0_KEYS) + }, } print(json.dumps(result, indent=2)) diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py index 382125e..eac51f9 100644 --- a/c/tools/convert_fp8_to_int4.py +++ b/c/tools/convert_fp8_to_int4.py @@ -286,6 +286,17 @@ def main(): # testa MTP a int4 = acceptance ~0-4% (misurato, issue #8): il draft sbaglia sempre # e la speculazione non parte mai. A int8: 39-59%, 2.2-2.8 token/forward. a.ebits = 8 if (a.mtp or a.indexer) else 4 + if a.mtp and a.ebits < 8 and a.group_size <= 0: + # Non solo lossy: eh_proj ha ~20-30x di asimmetria di scala fra le due meta' di + # colonna, quindi l'int4 per-riga (UNA scala per riga) arrotonda a ZERO l'intera + # meta' embedding -> il draft non vede il token -> acceptance ~0% (issue #8). + # EN: not merely lossy: eh_proj has ~20-30x column-scale asymmetry, so per-row + # EN: int4 rounds its ENTIRE embedding half to exact zeros -> the draft cannot + # EN: see the input token -> ~0% acceptance (issue #8). A container converted + # EN: this way is repairable in place with tools/repair_mtp_int8.py. + print(f"WARNING: --mtp with --ebits {a.ebits} and per-row scales ZEROES eh_proj's " + "embedding half -> MTP acceptance ~0% (issue #8). Use the default --ebits 8, " + "or add --group-size 128 for group-scaled int4.") if a.xbits is None: a.xbits = a.ebits # Build per-type bits map. If a type-specific arg is set, use it; otherwise the @@ -299,6 +310,16 @@ def main(): if bits_map: print(f"[MIXED] precision map: " + ", ".join(f"{k}={v}bit" for k,v in sorted(bits_map.items()))) + # Il PIANO risolto, PRIMA di toccare qualunque cosa (#383): --mtp/--indexer cambiano il + # default di ebits a 8 (testa int4 = acceptance ~0%, issue #8) e il ramo grouped e' + # gated su bits<=4 — combinazioni sorprendenti devono mostrarsi al secondo 1 di un job + # da ore, non nel size-check dopo. EN: print the RESOLVED plan before doing anything. + mode = "MTP head only" if a.mtp else "DSA indexer only" if a.indexer else "main model" + grp = f"grouped gs={a.group_size} (fmt=4)" if (a.group_size and a.ebits <= 4) else \ + (f"PER-ROW (grouped branch needs bits<=4; ebits={a.ebits} disables it)" if a.group_size else "per-row") + print(f"[PLAN] mode: {mode} | source: {'local ' + a.indir if a.indir else 'download ' + a.repo} | " + f"experts {a.ebits}-bit, embed/lm_head {a.io_bits}-bit, x {a.xbits}-bit | {grp}") + if a.selftest_nvfp4: import torch # 1) LUT e2m1: i 16 codici devono decodificare esattamente ai valori attesi. @@ -379,14 +400,99 @@ def main(): if a.indir: # conversione locale (test) shards = sorted(glob.glob(os.path.join(a.indir, "*.safetensors"))) from safetensors.numpy import save_file + # #383: se l'indice c'e', i passaggi --mtp/--indexer convertono SOLO gli shard + # che contengono i tensori richiesti (3 invece di scandire tutti i 141 — ogni + # scansione a vuoto apre comunque uno shard da 5 GB). Senza indice: scansione + # completa come prima. + # EN: #383: when the index is present, the --mtp/--indexer passes convert ONLY + # the shards that hold the requested tensors (3 instead of scanning all 141 — + # every empty scan still opens a 5 GB shard). Without the index: full scan as + # before. + if a.mtp or a.indexer: + idxp = os.path.join(a.indir, "model.safetensors.index.json") + if os.path.exists(idxp): + wmap = json.load(open(idxp))["weight_map"] + if a.mtp: + want = {v for k, v in wmap.items() if k.startswith(f"model.layers.{a.n_layers}.")} + else: + want = {v for k, v in wmap.items() if "indexer" in k and 0 <= layer_idx(k) < a.n_layers} + keep = [sp for sp in shards if os.path.basename(sp) in want] + print(f"[PLAN] index: {len(keep)}/{len(shards)} local shard(s) hold the requested tensors") + shards = keep + # BUG #355: questo ramo ignorava --mtp/--indexer. Con --mtp scriveva + # out-NNNNN (gli STESSI nomi di una conversione normale) in ebits=8 e + # keep_mtp=False -> il "secondo passaggio MTP" nella stessa outdir + # SOVRASCRIVEVA il container gia' finito con una riconversione int8 + # completa, in silenzio (137/141 shard distrutti prima di accorgersene). + # Ora il ramo locale rispecchia il download path: prefisso corretto, + # flag passate, shard vuoti saltati. + prefix = "out-mtp-" if a.mtp else "out-idx-" if a.indexer else "out-" + # RIPRESA (#383): i nomi out-NNNNN contano gli shard EMESSI, non l'indice di + # input (gli shard senza tensori rilevanti non producono file), quindi "il + # file esiste" non basta per saltare il lavoro gia' fatto. Un manifest + # sidecar ricorda input -> output (o "vuoto") e con quali parametri: la + # ripresa salta solo cio' che combacia, e parametri diversi sulla stessa + # outdir vengono rifiutati invece di mescolare container (il modo #355). + # EN: RESUME (#383): out-NNNNN names count EMITTED shards, not the input + # EN: index (shards with no relevant tensors emit no file), so "the file + # EN: exists" is not enough to skip completed work. A sidecar manifest + # EN: records input -> output (or "empty") plus the conversion parameters: + # EN: resume skips only what matches, and different parameters on the same + # EN: outdir are refused instead of mixing containers (the #355 failure mode). + params = {"ebits": a.ebits, "io_bits": a.io_bits, "xbits": a.xbits, + "group_size": a.group_size, "n_layers": a.n_layers, "bits_map": bits_map} + prog_path = os.path.join(a.outdir, f".{prefix}progress.json") + prog = {} + if os.path.exists(prog_path): + try: prog = json.loads(open(prog_path).read()) + except (OSError, ValueError): prog = {} + if prog and prog.get("params") != params: + print(f"ERROR: {prog_path} records a conversion with {prog.get('params')};\n" + f" this run uses {params}. Refusing to mix conversions in the same " + f"outdir — use a fresh --outdir (or delete the manifest and the " + f"{prefix}*.safetensors shards to redo).") + return + done = prog.setdefault("shards", {}); prog["params"] = params + n = 0; fresh = 0; skipped = 0 for i, sp in enumerate(shards): - out = {}; convert_shard(sp, out, a.n_layers, a.ebits, a.io_bits, a.xbits, group_size=a.group_size, bits_map=bits_map) - save_file(out, os.path.join(a.outdir, f"out-{i:05d}.safetensors")) - # copia config + tokenizer - for fn in ["config.json"]: - src = os.path.join(a.indir, fn) - if os.path.exists(src): shutil.copy(src, a.outdir) - print(f"converted {len(shards)} shards -> {a.outdir}") + key = os.path.basename(sp) + prev = done.get(key) # None = mai visto; "" = visto, vuoto; nome = emesso + if prev is not None and (prev == "" or os.path.exists(os.path.join(a.outdir, prev))): + if prev: n += 1 + skipped += 1 + continue + out = {} + convert_shard(sp, out, a.n_layers, a.ebits, a.io_bits, a.xbits, + keep_mtp=a.mtp, keep_idx=a.indexer, + group_size=a.group_size, bits_map=bits_map) + if not out: # shard senza MTP/idx: niente file (come il download path) + done[key] = "" + else: + name = f"{prefix}{n:05d}.safetensors" + save_file(out, os.path.join(a.outdir, name)) + done[key] = name; n += 1; fresh += 1 + tmp_prog = prog_path + ".tmp" # scrittura atomica: una ripresa non vede mai un manifest mezzo scritto + with open(tmp_prog, "w") as f: json.dump(prog, f, indent=1) # EN: atomic write: a resume never sees a half-written manifest + os.replace(tmp_prog, prog_path) + if skipped: print(f"[RESUME] {skipped} shard(s) already done in {a.outdir}, skipped") + # metadati per la conversione principale: gli stessi quattro file del download + # path — senza tokenizer.json chat/serve non partono. I passaggi mtp/idx vanno + # nella stessa outdir di un container gia' completo di metadati. + # EN: metadata for the main pass: the same four files as the download path — + # EN: chat/serve won't start without tokenizer.json. The mtp/idx passes target + # EN: an outdir whose container already has its metadata. + if not a.mtp and not a.indexer: + copied, missing = [], [] + for fn in ["config.json", "tokenizer.json", "tokenizer_config.json", "generation_config.json"]: + src = os.path.join(a.indir, fn) + if os.path.exists(src): shutil.copy(src, a.outdir); copied.append(fn) + else: missing.append(fn) + print(f"[META] copied from {a.indir}: {', '.join(copied) if copied else 'nothing'}") + if missing: + print(f"[META] WARNING: not found in {a.indir}: {', '.join(missing)}" + + (" — chat/serve need tokenizer.json" if "tokenizer.json" in missing else "")) + tag = "MTP" if a.mtp else "indexer" if a.indexer else "main" + print(f"converted {fresh} {tag} shard(s), {n} in container -> {a.outdir} ({prefix}NNNNN)") return # reale: scarica shard per shard, converte, cancella diff --git a/c/tools/repair_mtp_int8.py b/c/tools/repair_mtp_int8.py new file mode 100644 index 0000000..695e99f --- /dev/null +++ b/c/tools/repair_mtp_int8.py @@ -0,0 +1,229 @@ +"""Repair an existing colibri int4 container whose MTP head was quantized at int4. + +WHY THIS EXISTS + `model.layers..eh_proj.weight` [D, 2D] multiplies the MTP concat + [embedding_norm ; hidden_norm], and its two column halves differ in scale by + ~20-30x per row (embedding-half absmax ~0.05, hidden-half ~1.5 on GLM-5.2). + Per-row int4 uses ONE scale (= absmax/7) per row, so every embedding-half + weight lands below half a quantization step and np.rint rounds the ENTIRE + embedding half to exact zeros (packed bytes 0x88). The MTP head then drafts + garbage: acceptance ~0% (issue #8 measured 0-4% at int4; 39-59% at int8 — + which is why `convert_fp8_to_int4.py --mtp` defaults to --ebits 8). + + A container converted (or downloaded) with an int4 MTP head does not need a + full re-conversion: this script re-downloads ONLY the affected dense tensors + (~355 MB of HTTP range reads against the FP8 source repo), requantizes them + at int8 with the converter's exact math, and patches the local shards in + place. Originals are kept beside as *.bak-int4. + +WHAT IT TOUCHES + The MTP layer's dense tensors only (eh_proj, q_a/q_b/kv_a/kv_b/o_proj, + shared_experts.*): the ones that stream into RAM once and stay resident. + Routed experts (model.layers..mlp.experts.*) are NOT touched — they are + statistically like the main layers' experts and int4 is acceptable there. + The engine auto-detects int8 vs int4 by blob size (qt_from_disk), so no + engine or config change is needed. Cost: ~+133 MB on disk / resident RAM. + +USAGE + python3 tools/repair_mtp_int8.py --snap /path/to/glm52_i4 # repair + python3 tools/repair_mtp_int8.py --snap /path/to/glm52_i4 --dry-run # inspect only + + --source-repo defaults to zai-org/GLM-5.2-FP8 (the checkpoint the public + int4 containers were converted from). Requires numpy and network access; + no torch, no HF token (public repo, anonymous range reads). +""" +import argparse, glob, json, os, ssl, struct, sys, urllib.request +import numpy as np + +# macOS python.org builds ship no CA bundle: use certifi when available (Linux +# system Pythons generally have working system certs and skip this). +try: + import certifi + _SSL_CTX = ssl.create_default_context(cafile=certifi.where()) +except ImportError: + _SSL_CTX = ssl.create_default_context() + + +# ---------- HTTP range reads against the source repo ---------- +def http_range(url, start, length, tries=5): + req = urllib.request.Request(url, headers={"User-Agent": "colibri-mtp-repair", + "Range": f"bytes={start}-{start+length-1}"}) + for attempt in range(tries): + try: + with urllib.request.urlopen(req, timeout=30, context=_SSL_CTX) as r: + data = r.read() + if len(data) == length: + return data + except KeyboardInterrupt: + raise + except Exception as ex: + if attempt == tries - 1: + raise RuntimeError(f"range read failed for {url}: {ex}") + raise RuntimeError(f"short range read for {url}") + + +class SourceRepo: + def __init__(self, repo, revision="main"): + self.base = f"https://huggingface.co/{repo}/resolve/{revision}/" + with urllib.request.urlopen(self.base + "model.safetensors.index.json", timeout=30, context=_SSL_CTX) as r: + self.wmap = json.loads(r.read())["weight_map"] + self._hdr = {} + + def _shard_header(self, shard): + if shard not in self._hdr: + n = struct.unpack("> 3) & 0xF + mant = (b & 7).astype(np.float64) + v = (sign * np.where(e > 0, (1 + mant / 8) * np.exp2(e.astype(np.float64) - 7), + mant / 8 * np.exp2(-6.0))).reshape(m["shape"]) + sn = name + "_scale_inv" + sshard = self.wmap[sn] + shdr, sbase = self._shard_header(sshard) + sm = shdr[sn] + so0, so1 = sm["data_offsets"] + sc = np.frombuffer(http_range(self.base + sshard, sbase + so0, so1 - so0), + dtype=np.float32).reshape(sm["shape"]) + O, I = m["shape"] + scf = np.repeat(np.repeat(sc, 128, axis=0)[:O], 128, axis=1)[:, :I] + return (v * scf).astype(np.float32) + raise ValueError(f"{name}: unsupported source dtype {m['dtype']}") + + +# ---------- quantization: identical to convert_fp8_to_int4.quant_int8 ---------- +def quant_int8(w): + amax = np.abs(w).max(axis=1, keepdims=True) + s = np.maximum(amax / 127, 1e-8) + q = np.clip(np.rint(w / s), -128, 127).astype(np.int8) + return q.reshape(-1).view(np.uint8).copy(), s[:, 0].astype(np.float32) + + +# ---------- local safetensors IO (no deps; preserves byte-identity of untouched tensors) ---------- +def read_shard(path): + with open(path, "rb") as fh: + n = struct.unpack(" [tensor names to repair] + already_ok, skipped_experts = [], 0 + for f in sorted(glob.glob(os.path.join(a.snap, "*.safetensors"))): + with open(f, "rb") as fh: + n = struct.unpack(" --local-dir D:\glm52_i4 -``` - -Use the container recommended in the README (with **int8 MTP heads** — int4 heads silently give 0% draft acceptance). The download is resumable: if it stops, rerun the same command. Expect hours; everything below fits inside them. - -## 2. Build the engine (CPU) - -From a normal PowerShell, in the repo's `c\` directory: - -```powershell -make glm.exe ARCH=native # ARCH=native unlocks AVX-VNNI on Alder Lake+/Arrow Lake -make iobench.exe # disk benchmark, useful before committing to the download -``` - -Warnings about `#pragma comment` and unused variables are normal (MSVC-isms gcc ignores). The engine banner should print `idot: avx-vnni` on VNNI-capable CPUs — if it says avx2, you built without `ARCH=native`. - -### ⚠️ Smart App Control will block your fresh binary - -On Windows 11 machines with **Smart App Control** enforced (`VerifiedAndReputablePolicyState = 1`), running your self-compiled `glm.exe` fails with: - -``` -Program 'glm.exe' failed to run: An Application Control policy has blocked this file -``` - -This is not Defender and not Mark-of-the-Web — SAC blocks *all* unsigned, unknown binaries, which includes anything you compile yourself. **Fix:** Windows Security → App & browser control → Smart App Control settings → **Off**, then **reboot** (the policy only reloads on restart). Note SAC is one-way: re-enabling later requires resetting Windows. If the settings page is missing, the registry equivalent is setting `HKLM:\SYSTEM\CurrentControlSet\Control\CI\Policy\VerifiedAndReputablePolicyState` to `0` (admin PowerShell), then rebooting. Check your current state before touching anything: - -```powershell -(Get-ItemProperty "HKLM:\SYSTEM\CurrentControlSet\Control\CI\Policy").VerifiedAndReputablePolicyState -# 0 = off, 1 = enforced, 2 = evaluation -``` - -## 3. Build the CUDA DLL (GPU tier) - -nvcc needs MSVC as host compiler, so this one step must run from a shell with the MSVC environment: open **"x64 Native Tools Command Prompt for VS 2022"** from the Start menu (plain PowerShell will fail the `cl` check). Then: - -```cmd -make cuda-dll CUDA_ARCH=sm_120 # match your GPU: sm_120 Blackwell, sm_89 Ada, ... -make glm.exe CUDA_DLL=1 ARCH=native # relink host with the runtime loader -``` - -Two pitfalls, both fixed on current `dev` (#314) but worth knowing on older checkouts: - -- **Spaces in `CUDA_HOME`** (`C:\Program Files\...`) used to break the recipe → fixed; nvcc now comes from PATH and `"$(NVCC)"` is quoted. -- **`make glm.exe CUDA_DLL=1` after a CPU-only build** used to report `up to date` and silently keep the CPU-only binary (GPU tier never engages, no error). Current `dev` has a build-config stamp that forces the relink. On older trees: delete `glm.exe` first. - -Sanity check: first GPU run should print `[CUDA] device 0: , ... sm_XX` and `[CUDA] mode: routed experts + resident dense tensors`. - -## 4. First run - -```powershell -cd \c -$env:OMP_NUM_THREADS = "" -python coli run "Explain what a mixture-of-experts model is." --model D:\glm52_i4 --ngen 48 -``` - -The first run is cold — expect the profile to be dominated by `expert-disk` while the cache warms; hit rate climbs run over run. GPU tier on top: - -```powershell -$env:COLI_CUDA="1"; $env:COLI_GPU="0"; $env:CUDA_DENSE="1"; $env:CUDA_EXPERT_GB="4" -python coli run "..." --model D:\glm52_i4 --ngen 64 -``` - -Size `CUDA_EXPERT_GB` so dense (~10 GB) + experts + working set stays under your VRAM. Note MTP speculation is off by default under CUDA (#293, float-accumulation divergence between draft and verify) — `COLI_CUDA_MTP=1` opts back in. - -## 5. Reference numbers from this walkthrough's hardware - -285K / RTX 5080 / 128 GB / NVMe at 5.85 GB/s random-read (19 MB blocks, `iobench`): 0.26 tok/s cold CPU → 0.30 warm CPU (MTP 2.2–2.3 tok/forward) → 0.42 tok/s GPU tier + auto-pin, expert hit 66%, ~65% of wall time in expert-disk. Disk-bound is the expected shape at ~25% expert residency — a faster disk and more RAM move the floor, the GPU moves the compute. - -## Quick failure index - -| Symptom | Cause | Fix | -|---|---|---| -| `An Application Control policy has blocked this file` | Smart App Control | §2 — turn SAC off + **reboot** | -| `cuda-dll ... Error 1` immediately | old tree: spaced CUDA_HOME / MSVC rejects `-Wextra` | update to current `dev` (#314) | -| `glm.exe is up to date` but GPU never engages | old tree: stale CPU-only binary | update to `dev`, or delete `glm.exe` and rebuild | -| `cl.exe (MSVC) not in PATH` | built from plain PowerShell | use the x64 Native Tools prompt | -| `nvcc fatal: unsupported gpu architecture 'sm_120'` | CUDA < 12.8 | install CUDA 12.8+ | -| MTP `0% (0/0)` on CPU path | int4 MTP heads in the container | use the int8-MTP container | -| MTP `draft=0` under CUDA | intended default since #293 | `COLI_CUDA_MTP=1` to opt in | diff --git a/docs/api.md b/docs/api.md index d3b09f1..62010bd 100644 --- a/docs/api.md +++ b/docs/api.md @@ -49,6 +49,82 @@ errors before streaming headers are sent. `GET /health` exposes active/queued/completed/rejected counters, and successful generation responses include `x-colibri-queue-wait-ms`. +## Connect a coding CLI or editor + +The API is OpenAI-compatible, so most coding CLIs and editor extensions work by +pointing them at Colibri as an *OpenAI-compatible* provider. Three settings: + +- **Base URL** — `http://localhost:8000/v1` +- **Model** — `glm-5.2-colibri` (or whatever you pass to `--model-id`) +- **API key** — any non-empty string, e.g. `local` + +Colibri needs **no** API key by default, but many clients refuse to start without +one — give them any dummy value. The key is only enforced if you set `COLI_API_KEY`. + +Smoke-test the endpoint first (no key needed unless you set one): + +```bash +curl http://127.0.0.1:8000/v1/chat/completions \ + -H 'Content-Type: application/json' \ + -d '{"model":"glm-5.2-colibri","messages":[{"role":"user","content":"hi"}]}' +``` + +**aider** + +```bash +export OPENAI_API_BASE=http://localhost:8000/v1 +export OPENAI_API_KEY=local +aider --model openai/glm-5.2-colibri # the openai/ prefix routes to OPENAI_API_BASE +``` + +**crush** — add a provider to `crush.json` (`~/.config/crush/crush.json`, or +`%USERPROFILE%\AppData\Local\crush\crush.json` on Windows): + +```json +{ + "$schema": "https://charm.land/crush.json", + "providers": { + "colibri": { + "name": "Colibri", + "type": "openai-compat", + "base_url": "http://localhost:8000/v1/", + "api_key": "local", + "models": [ + { "name": "GLM-5.2 (Colibri)", "id": "glm-5.2-colibri", + "context_window": 131072, "default_max_tokens": 1024 } + ] + } + } +} +``` + +The `"api_key": "local"` dummy is what satisfies clients that demand a key. +`context_window` is only the client's budget display — set it to whatever your +KV configuration actually allows. + +**Continue, Cline / Roo, `llm`, the OpenAI SDKs, …** — set the provider's base +URL to `http://localhost:8000/v1`, the model to `glm-5.2-colibri`, and any dummy +key (`OPENAI_API_KEY` / `OPENAI_BASE_URL` for env-based tools). + +> **Set your expectations before connecting an agentic CLI.** Two costs dominate, +> and the first one is invisible until you know it's there: +> +> 1. **Prefill.** Coding agents (crush, aider in repo-map mode, Cline, …) send a +> large system prompt plus tool definitions — often 10–20k tokens — *before +> your first word*. Prefill on the CPU-streaming path runs at a few tokens per +> second (it is attention-bound, see #153), so a 15k-token agent preamble is +> **an hour of silent "thinking" before the first output token**. The client +> looks hung; it isn't. Smoke-test with the tiny `curl` above first — if that +> answers in about a minute, the pipeline works and what you're paying for is +> prompt size. +> 2. **Decode.** Roughly 1 tok/s for a large model, so multi-turn agent loops +> (which re-pay the growing context every turn) compound the cost. +> +> Practical guidance: single surgical asks with a short context work; iterative +> agent sessions against a disk-streaming 744B model do not resemble a hosted +> API and mostly won't be worth the wait. If your client lets you trim or disable +> its system preamble and tool catalog, do it. + ## Isolated KV contexts `coli serve --kv-slots N` allocates up to 16 independent sequence contexts. diff --git a/docs/benchmarks.md b/docs/benchmarks.md index f9c83f2..09d3752 100644 --- a/docs/benchmarks.md +++ b/docs/benchmarks.md @@ -116,8 +116,12 @@ can match an RTX 5090 on expert matmul ([#101](https://github.com/JustVugg/colib so **the GPU tier earns its VRAM only when the CPU is the weak link**. On multi-socket hosts, NUMA placement is a further lever: interleaving the resident weights across nodes measured **+13% (2-socket) and +40% (4-socket CPU-only)** -([#82](https://github.com/JustVugg/colibri/issues/82)) — but never blanket-interleave -a GPU host (measured 10× regression via the DMA staging pages). +([#82](https://github.com/JustVugg/colibri/issues/82)). On a 2-socket Xeon Silver +4510 host with 6× RTX 5090, selective `COLI_NUMA=1` raised effective CPU-expert +bandwidth from **42.42 to 58.26/65.89 GB/s** and greedy decode from **7.66 to +9.02/9.17 tok/s** (64 tokens, `TEMP=0 DRAFT=0`, byte-identical output). Do not +blanket-interleave a GPU host: it also spreads DMA staging pages and has measured +up to a 10× regression; generated plans enable only the selective slab policy. ## Quality benchmark diff --git a/docs/windows.md b/docs/windows.md index 94a4f69..ecb81e7 100644 --- a/docs/windows.md +++ b/docs/windows.md @@ -1,25 +1,107 @@ -# Windows 11 (native, no WSL) +# Windows 11 native install — a complete walkthrough (no WSL) -colibrì builds and runs natively on Windows 11 x86-64 with MinGW-w64. The port -adds a `_WIN32` compatibility layer in `c/compat.h` that maps POSIX I/O to the -Windows API (pread → ReadFile+OVERLAPPED, posix_fadvise no-op, aligned -allocation, MoveFileEx rename, GlobalMemoryStatusEx RAM detection). All platform -differences stay in `compat.h`; the engine source is unchanged. +A start-to-finish, reproducible path from a fresh Windows 11 machine to GLM-5.2 generating tokens, with the GPU tier. Every step and every failure mode below was hit and verified on real hardware: Core Ultra 9 285K (AVX-VNNI) / RTX 5080 (sm_120) / 128 GB RAM / Windows 11 24H2 (issue #306). Steps are ordered so the long downloads run while you build. -**Toolchain:** GCC via [winlibs](https://winlibs.com/) or MSYS2 MinGW-w64. -Tested with GCC 16.1.0 (x86_64-ucrt-posix-seh). +## 0. What you need + +| Piece | Why | Get it | +|---|---|---| +| git, Python 3 | clone + `coli` launcher | winget / python.org | +| MinGW-w64 gcc + make | builds the engine (MSVC can't) | `scoop install mingw-winlibs`, MSYS2, or portable **w64devkit** (no admin, unzip and go) | +| CUDA Toolkit ≥ 12.8 | GPU tier; ≥12.8 required for Blackwell/sm_120 | `winget install Nvidia.CUDA` | +| MSVC Build Tools (C++ workload) | nvcc's host compiler for the CUDA DLL | `winget install Microsoft.VisualStudio.2022.BuildTools` + "Desktop development with C++" | +| ~400 GB free on a local NVMe | the int4 model (~370–384 GB) | NTFS is fine; **never** a network mount | + +RAM: 16 GB minimum, more = bigger expert cache = faster. The build itself needs none of the CUDA/MSVC pieces — do the CPU build first, add the GPU tier later. + +## 1. Start the model download first (it's the long pole) ```powershell -# One-time toolchain install (pick one): -scoop install mingw-winlibs # portable, no shell needed -# or: pacman -S mingw-w64-x86_64-gcc make # via MSYS2 +python -m pip install -U "huggingface_hub[hf_transfer]" +$env:HF_HUB_ENABLE_HF_TRANSFER = "1" +hf download --local-dir D:\glm52_i4 +``` -# Build (from c/ directory): -make glm.exe # GLM-5.2 engine (static, no DLL dependencies) -make olmoe.exe # OLMoE engine (same shims) -make iobench.exe # disk I/O benchmark -make test-c # run C tests -make test-python # run Python tests (requires python) +Use the container recommended in the README (with **int8 MTP heads** — int4 heads silently give 0% draft acceptance). The download is resumable: if it stops, rerun the same command. Expect hours; everything below fits inside them. + +## 2. Build the engine (CPU) + +From a normal PowerShell, in the repo's `c\` directory: + +```powershell +make glm.exe ARCH=native # ARCH=native unlocks AVX-VNNI on Alder Lake+/Arrow Lake +make iobench.exe # disk benchmark, useful before committing to the download +``` + +Warnings about `#pragma comment` and unused variables are normal (MSVC-isms gcc ignores). The engine banner should print `idot: avx-vnni` on VNNI-capable CPUs — if it says avx2, you built without `ARCH=native`. + +### ⚠️ Smart App Control will block your fresh binary + +On Windows 11 machines with **Smart App Control** enforced (`VerifiedAndReputablePolicyState = 1`), running your self-compiled `glm.exe` fails with: + +``` +Program 'glm.exe' failed to run: An Application Control policy has blocked this file +``` + +This is not Defender and not Mark-of-the-Web — SAC blocks *all* unsigned, unknown binaries, which includes anything you compile yourself. **Fix:** Windows Security → App & browser control → Smart App Control settings → **Off**, then **reboot** (the policy only reloads on restart). Note SAC is one-way: re-enabling later requires resetting Windows. If the settings page is missing, the registry equivalent is setting `HKLM:\SYSTEM\CurrentControlSet\Control\CI\Policy\VerifiedAndReputablePolicyState` to `0` (admin PowerShell), then rebooting. Check your current state before touching anything: + +```powershell +(Get-ItemProperty "HKLM:\SYSTEM\CurrentControlSet\Control\CI\Policy").VerifiedAndReputablePolicyState +# 0 = off, 1 = enforced, 2 = evaluation +``` + +## 3. Build the CUDA DLL (GPU tier) + +nvcc needs MSVC as host compiler, so this one step must run from a shell with the MSVC environment: open **"x64 Native Tools Command Prompt for VS 2022"** from the Start menu (plain PowerShell will fail the `cl` check). Then: + +```cmd +make cuda-dll CUDA_ARCH=sm_120 # match your GPU: sm_120 Blackwell, sm_89 Ada, ... +make glm.exe CUDA_DLL=1 ARCH=native # relink host with the runtime loader +``` + +Two pitfalls, both fixed on current `dev` (#314) but worth knowing on older checkouts: + +- **Spaces in `CUDA_HOME`** (`C:\Program Files\...`) used to break the recipe → fixed; nvcc now comes from PATH and `"$(NVCC)"` is quoted. +- **`make glm.exe CUDA_DLL=1` after a CPU-only build** used to report `up to date` and silently keep the CPU-only binary (GPU tier never engages, no error). Current `dev` has a build-config stamp that forces the relink. On older trees: delete `glm.exe` first. + +Sanity check: first GPU run should print `[CUDA] device 0: , ... sm_XX` and `[CUDA] mode: routed experts + resident dense tensors`. + +## 4. First run + +```powershell +cd \c +$env:OMP_NUM_THREADS = "" +python coli run "Explain what a mixture-of-experts model is." --model D:\glm52_i4 --ngen 48 +``` + +The first run is cold — expect the profile to be dominated by `expert-disk` while the cache warms; hit rate climbs run over run. GPU tier on top: + +```powershell +$env:COLI_CUDA="1"; $env:COLI_GPU="0"; $env:CUDA_DENSE="1"; $env:CUDA_EXPERT_GB="4" +python coli run "..." --model D:\glm52_i4 --ngen 64 +``` + +Size `CUDA_EXPERT_GB` so dense (~10 GB) + experts + working set stays under your VRAM. Note MTP speculation is off by default under CUDA (#293, float-accumulation divergence between draft and verify) — `COLI_CUDA_MTP=1` opts back in. + +## 5. Reference numbers from this walkthrough's hardware + +285K / RTX 5080 / 128 GB / NVMe at 5.85 GB/s random-read (19 MB blocks, `iobench`): 0.26 tok/s cold CPU → 0.30 warm CPU (MTP 2.2–2.3 tok/forward) → 0.42 tok/s GPU tier + auto-pin, expert hit 66%, ~65% of wall time in expert-disk. Disk-bound is the expected shape at ~25% expert residency — a faster disk and more RAM move the floor, the GPU moves the compute. + +## Quick failure index + +| Symptom | Cause | Fix | +|---|---|---| +| `An Application Control policy has blocked this file` | Smart App Control | §2 — turn SAC off + **reboot** | +| `cuda-dll ... Error 1` immediately | old tree: spaced CUDA_HOME / MSVC rejects `-Wextra` | update to current `dev` (#314) | +| `glm.exe is up to date` but GPU never engages | old tree: stale CPU-only binary | update to `dev`, or delete `glm.exe` and rebuild | +| `cl.exe (MSVC) not in PATH` | built from plain PowerShell | use the x64 Native Tools prompt | +| `nvcc fatal: unsupported gpu architecture 'sm_120'` | CUDA < 12.8 | install CUDA 12.8+ | +| MTP `0% (0/0)` on CPU path | int4 MTP heads in the container | use the int8-MTP container | +| MTP `draft=0` under CUDA | intended default since #293 | `COLI_CUDA_MTP=1` to opt in | + +--- + +## Reference: build flags & warmup # AVX-VNNI: Intel Alder Lake+ (and Meteor Lake+) CPUs have a 128-bit int8 # dot-product instruction (VPDPBUSD) the engine can use for ~1.3x faster