From 72d3d37231e922a6fa9afca16e08fa45842d5eb4 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Thu, 16 Jul 2026 19:39:40 +0200
Subject: [PATCH 01/31] ci: run the checks on main too (#144)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

@ZacharyZcR's workflow, already green on dev (78c77bf): engine + C suite,
web build + vitest, python suite, and a real CUDA syntax check.

Only the workflow file is cherry-picked here — main's engine code stays at
54cfe56 and nothing else from dev comes along. main is where the checks are
worth the most and where they were missing entirely.

Includes the two fixes the first run exposed: the pinned action's version
table stops at 12.6.2 (12.6.3 failed the install), and `nvcc ... | head -40`
took its exit status from head, so the job could never fail.

Co-Authored-By: ZacharyZcR <#144>

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 .github/workflows/ci.yml | 72 ++++++++++++++++++++++++++++++++++++++++
 1 file changed, 72 insertions(+)
 create mode 100644 .github/workflows/ci.yml

diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
new file mode 100644
index 0000000..e18f1b3
--- /dev/null
+++ b/.github/workflows/ci.yml
@@ -0,0 +1,72 @@
+name: CI
+
+on:
+  pull_request:
+    branches: [dev, main]
+  push:
+    branches: [dev, main]
+
+jobs:
+  engine:
+    name: Engine (Linux, CPU)
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
+      - name: Build glm
+        run: cd c && make glm
+      - name: C test suite
+        run: cd c && make test-c
+
+  engine-cuda-syntax:
+    name: CUDA syntax check
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
+      - name: Install CUDA toolkit (compiler only)
+        uses: Jimver/cuda-toolkit@v0.2.19
+        with:
+          # v0.2.19's version table stops at 12.6.2 — 12.6.3 fails the install
+          # step with "Version not available" before nvcc is ever reached.
+          cuda: '12.6.2'
+          method: network
+          sub-packages: '["nvcc"]'
+      - name: Compile backend_cuda.cu (syntax only, no GPU)
+        run: |
+          cd c
+          # NOT `nvcc ... | head -40`: a pipeline exits with the status of its
+          # LAST command, so head's 0 masked every nvcc error and the job could
+          # only ever be green. A check that cannot fail is worse than no check —
+          # it buys false confidence in the one file nothing else compiles.
+          nvcc -O2 -std=c++17 -arch=sm_80 -c backend_cuda.cu -o /dev/null \
+            -Xcompiler=-Wall,-Wextra
+          echo "CUDA syntax check passed"
+
+  web:
+    name: Web UI
+    runs-on: ubuntu-latest
+    defaults:
+      run:
+        working-directory: web
+    steps:
+      - uses: actions/checkout@v4
+      - uses: actions/setup-node@v4
+        with:
+          node-version: '22'
+          cache: npm
+          cache-dependency-path: web/package-lock.json
+      - run: npm ci
+      - name: Build
+        run: npm run build
+      - name: Test
+        run: npx vitest run
+
+  python:
+    name: Python tests
+    runs-on: ubuntu-latest
+    steps:
+      - uses: actions/checkout@v4
+      - uses: actions/setup-python@v5
+        with:
+          python-version: '3.12'
+      - name: Python test suite
+        run: cd c && python3 -m unittest discover -s tests -p 'test_*.py'

From 5e2be61a2c49ba9b10f6710853a85a545e1ea448 Mon Sep 17 00:00:00 2001
From: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
Date: Fri, 17 Jul 2026 08:27:51 -0400
Subject: [PATCH 02/31] fix(test): make test_stops build on Windows (mkdtemp
 compat shim)

test_stops.c uses POSIX mkdtemp() to make a scratch dir in the CWD, but
MinGW-w64 does not declare mkdtemp, so the test failed to compile on the
Windows job - and only there:

  tests/test_stops.c:73:9: error: implicit declaration of function 'mkdtemp';
     did you mean 'mktemp'? [-Wimplicit-function-declaration]
  make: *** [Makefile:318: tests/test_stops.exe] Error 1

That halted `make test` at the C-test stage on Windows (test_uring is
correctly Linux-only, so test_stops was the only blocker).

Added a compat_mkdtemp shim to compat.h, following the file's existing
convention (every platform difference lives there; the .c stays clean):
_mktemp fills the trailing X's in place (same contract as mkdtemp), then
_mkdir creates the directory. Also added <direct.h> for _mkdir. On Linux
compat.h is a complete no-op, so POSIX mkdtemp is untouched there.

Verified: test_stops builds clean on MinGW and all 5 sub-cases pass:
  tokenizer special-flag parsing, config/generation_config eos union,
  no-generation_config fallback, both-configs-mutilated tokenizer sweep,
  and the T=NULL validation path.
---
 c/compat.h | 14 ++++++++++++++
 1 file changed, 14 insertions(+)

diff --git a/c/compat.h b/c/compat.h
index 82de8b6..c477170 100644
--- a/c/compat.h
+++ b/c/compat.h
@@ -80,6 +80,7 @@ static inline int compat_open_direct(const char *path){
 #endif
 #include <windows.h>
 #include <io.h>
+#include <direct.h>   /* _mkdir (for the mkdtemp shim below) */
 #include <process.h>
 #include <malloc.h>
 #include <fcntl.h>
@@ -304,6 +305,19 @@ static inline int compat_setenv(const char *name, const char *value, int overwri
 }
 #define setenv(name,value,overwrite) compat_setenv(name,value,overwrite)
 
+/* --- 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 */
 
 /* --- compat_aligned_free su piattaforme diverse da Windows ---

From 1223f9ca2e12a5bfeaccf1f1eca90462fbd0ad97 Mon Sep 17 00:00:00 2001
From: ZacharyZcR <zacharyzcr1984@gmail.com>
Date: Sat, 18 Jul 2026 01:54:31 +0800
Subject: [PATCH 03/31] cuda: batch ragged attention across independent streams

---
 c/backend_cuda.cu                | 71 ++++++++++++++++++++++++++++++++
 c/backend_cuda.h                 |  6 ++-
 c/glm.c                          | 26 +++++++++---
 c/tests/test_ragged_attention.cu | 35 ++++++++++++++++
 4 files changed, 132 insertions(+), 6 deletions(-)
 create mode 100644 c/tests/test_ragged_attention.cu

diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu
index 9ce142f..188330d 100644
--- a/c/backend_cuda.cu
+++ b/c/backend_cuda.cu
@@ -338,6 +338,40 @@ __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 KV sequence per row. latent/rope are packed as [S,T,*], while
+ * lengths selects the valid prefix for each row. */
+__global__ static void attention_absorb_ragged_kernel(float *ctx,const float *q,
+        const float *latent,const float *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+(size_t)s*T*K,*rs=rope+(size_t)s*T*R;
+    for(int k=tid;k<K;k+=blockDim.x){float a=0;for(int d=0;d<Q;d++)
+        a+=qs[d]*weight_at(weights,fmt,(size_t)(rbase+d)*row_bytes(fmt,K),k)*
+          (fmt?wscale[rbase+d]:1.f);qa[k]=a;}
+    __syncthreads();
+    for(int t=tid;t<nt;t+=blockDim.x){float a=0;const float *lt=ls+(size_t)t*K;
+        const float *rt=rs+(size_t)t*R;for(int k=0;k<K;k++)a+=qa[k]*lt[k];
+        for(int d=0;d<R;d++)a+=qs[Q+d]*rt[d];scores[t]=a*scale;}
+    __syncthreads();
+    float local=-3.402823466e+38F;for(int t=tid;t<nt;t+=blockDim.x)local=fmaxf(local,scores[t]);
+    red[tid]=local;__syncthreads();
+    for(int n=blockDim.x>>1;n;n>>=1){if(tid<n)red[tid]=fmaxf(red[tid],red[tid+n]);__syncthreads();}
+    float mx=red[0];local=0;for(int t=tid;t<nt;t+=blockDim.x){float e=expf(scores[t]-mx);scores[t]=e;local+=e;}
+    red[tid]=local;__syncthreads();
+    for(int n=blockDim.x>>1;n;n>>=1){if(tid<n)red[tid]+=red[tid+n];__syncthreads();}
+    float inv=1.f/red[0];for(int t=tid;t<nt;t+=blockDim.x)scores[t]*=inv;
+    __syncthreads();
+    for(int k=tid;k<K;k+=blockDim.x){float a=0;for(int t=0;t<nt;t++)a+=scores[t]*ls[(size_t)t*K+k];cl[k]=a;}
+    __syncthreads();
+    for(int v=tid;v<V;v+=blockDim.x){int row=rbase+Q+v;float a=0;size_t rb=row_bytes(fmt,K);
+        for(int k=0;k<K;k++)a+=cl[k]*weight_at(weights,fmt,(size_t)row*rb,k);
+        ctx[((size_t)s*H+h)*V+v]=a*(fmt?wscale[row]:1.f);}
+}
+
 static int reserve(float **ptr, size_t *cap, size_t bytes) {
     if (*cap >= bytes) return 1;
     if (*ptr) cudaFree(*ptr);
@@ -770,6 +804,43 @@ 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 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||!latent||!rope||!lengths||S<1||S>512||T<1||T>512||
+       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;
+    size_t ln=(size_t)S*T*K,rn=(size_t)S*T*R;
+    float *lh=(float*)std::calloc(ln,sizeof(float)),*rh=(float*)std::calloc(rn,sizeof(float));
+    if(!lh||!rh){std::free(lh);std::free(rh);return 0;}
+    for(int s=0;s<S;s++){
+        if(lengths[s]<1||lengths[s]>T){std::free(lh);std::free(rh);return 0;}
+        std::memcpy(lh+(size_t)s*T*K,latent[s],(size_t)lengths[s]*K*sizeof(float));
+        std::memcpy(rh+(size_t)s*T*R,rope[s],(size_t)lengths[s]*R*sizeof(float));
+    }
+    DeviceContext *dc=find_ctx(w->device);
+    if(!select_ctx(dc)){std::free(lh);std::free(rh);return 0;}
+    size_t qb=(size_t)S*H*(Q+R)*sizeof(float),lb=ln*sizeof(float),rb=rn*sizeof(float);
+    size_t cb=(size_t)S*H*V*sizeof(float),ob=(size_t)S*proj->O*sizeof(float);
+    int ok=reserve(&dc->aq,&dc->aq_cap,qb)&&reserve(&dc->al,&dc->al_cap,lb)&&
+           reserve(&dc->ar,&dc->ar_cap,rb)&&reserve(&dc->ac,&dc->ac_cap,cb)&&
+           reserve(&dc->y,&dc->y_cap,ob)&&
+           reserve_bytes(&dc->group_desc,&dc->group_desc_cap,(size_t)S*sizeof(int));
+    if(ok)ok=cuda_ok(cudaMemcpyAsync(dc->aq,q,qb,cudaMemcpyHostToDevice,dc->stream),"ragged q upload")&&
+             cuda_ok(cudaMemcpyAsync(dc->al,lh,lb,cudaMemcpyHostToDevice,dc->stream),"ragged latent upload")&&
+             cuda_ok(cudaMemcpyAsync(dc->ar,rh,rb,cudaMemcpyHostToDevice,dc->stream),"ragged rope upload")&&
+             cuda_ok(cudaMemcpyAsync(dc->group_desc,lengths,(size_t)S*sizeof(int),cudaMemcpyHostToDevice,dc->stream),"ragged lengths upload");
+    std::free(lh);std::free(rh);if(!ok)return 0;
+    size_t shared=(size_t)(2*K+T+256)*sizeof(float);
+    attention_absorb_ragged_kernel<<<dim3(H,S),256,shared,dc->stream>>>(dc->ac,dc->aq,dc->al,dc->ar,
+        (const int*)dc->group_desc,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
+    quant_matmul<<<dim3(proj->O,S),256,0,dc->stream>>>(dc->y,dc->ac,proj->weights,
+        proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I));
+    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);
diff --git a/c/backend_cuda.h b/c/backend_cuda.h
index acbe4bc..edddbfe 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,10 @@ 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 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 +148,3 @@ COLI_CUDA_DLLEXPORT int coli_cuda_pipe_sync(int device);
 #endif
 
 #endif
-
diff --git a/c/glm.c b/c/glm.c
index 003a42d..298065b 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -1644,7 +1644,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);
@@ -2640,7 +2640,23 @@ 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));
+            int *rn=malloc((size_t)S*sizeof(*rn)); int mt=0;
+            if(rl&&rr&&rn){
+                for(int s=0;s<S;s++){
+                    int pos=positions[s],st0=kvs[s]->kv_start[layer]; rn[s]=pos+1-st0;
+                    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,rl,rr,rn,S,H,c->qk_nope,c->qk_rope,vh,kvl,mt,c->attn_scale);
+            }
+            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;d<n;d++)for(int s=0;s<S;s++)memcpy(
@@ -3975,7 +3991,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;s<S;s++) embed_row(m, ids[s], x+(int64_t)s*D);
     layers_forward(m,x,S,pos_base);
-    if(m->h_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;s<S;s++){ rmsnorm(row, x+(int64_t)s*D, m->final_norm, D, c->eps);
@@ -3988,8 +4004,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;s<S;s++){
         if(!rows[s].kv || !rows[s].kv->Lc || !rows[s].kv->Rc || !rows[s].kv->kv_start ||
diff --git a/c/tests/test_ragged_attention.cu b/c/tests/test_ragged_attention.cu
new file mode 100644
index 0000000..c00a8eb
--- /dev/null
+++ b/c/tests/test_ragged_attention.cu
@@ -0,0 +1,35 @@
+#include "../backend_cuda.h"
+
+#include <cmath>
+#include <cstdio>
+#include <vector>
+
+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<float> w(H*(Q+V)*K),p(O*D),q(S*H*(Q+R));
+    for(size_t i=0;i<w.size();i++)w[i]=((int)(i%11)-5)*.07f;
+    for(size_t i=0;i<p.size();i++)p[i]=((int)(i%7)-3)*.09f;
+    for(size_t i=0;i<q.size();i++)q[i]=((int)(i%13)-6)*.05f;
+    ColiCudaTensor *tw=nullptr,*tp=nullptr;
+    if(!coli_cuda_tensor_upload(&tw,w.data(),nullptr,0,K,H*(Q+V),dev)||
+       !coli_cuda_tensor_upload(&tp,p.data(),nullptr,0,D,O,dev))return 1;
+    int n[S]={1,2,3};std::vector<std::vector<float>> l(S),r(S);
+    const float *lp[S],*rp[S];
+    for(int s=0;s<S;s++){
+        l[s].resize(n[s]*K);r[s].resize(n[s]*R);
+        for(size_t i=0;i<l[s].size();i++)l[s][i]=((int)((i+s*3)%9)-4)*.08f;
+        for(size_t i=0;i<r[s].size();i++)r[s][i]=((int)((i+s)%5)-2)*.06f;
+        lp[s]=l[s].data();rp[s]=r[s].data();
+    }
+    float got[S*O],ref[S*O];
+    if(!coli_cuda_attention_project_ragged(tw,tp,got,q.data(),lp,rp,n,S,H,Q,R,V,K,T,.2f))return 2;
+    for(int s=0;s<S;s++)if(!coli_cuda_attention_project_batch(tw,tp,ref+s*O,
+        q.data()+s*H*(Q+R),lp[s],rp[s],1,H,Q,R,V,K,n[s],.2f))return 3;
+    double e=0,z=0;for(int i=0;i<S*O;i++){
+        double d=got[i]-ref[i];e+=d*d;z+=(double)ref[i]*ref[i];
+    }
+    double rms=std::sqrt(e/(z+1e-30));std::printf("ragged_relative_rms=%.9g\n",rms);
+    coli_cuda_tensor_free(tw);coli_cuda_tensor_free(tp);coli_cuda_shutdown();
+    return rms<1e-6?0:4;
+}

From ef97f3cf2b0a6f411de0ff02ae9bd0fba2d11923 Mon Sep 17 00:00:00 2001
From: volognamor <romangolov2001@gmail.com>
Date: Fri, 17 Jul 2026 21:49:28 +0300
Subject: [PATCH 04/31] Windows: fix non-ASCII chat prompt corruption (ANSI
 codepage vs UTF-8)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

coli passes the chat prompt to glm.exe through the PROMPT/COLI_PROMPT
environment variable. On Windows, plain getenv() is populated by the CRT
from the ANSI-codepage view of the environment block, not UTF-8 — so any
non-ASCII prompt text (Cyrillic, CJK, ...) is silently mangled before the
byte-level tokenizer ever sees it, even though the parent process (coli's
Python subprocess call) sets the value correctly via the wide env block.

Add compat_getenv_utf8() in compat.h: reads the variable through
GetEnvironmentVariableW and converts straight to UTF-8, bypassing the ANSI
codepage entirely. No-op passthrough to getenv() on non-Windows platforms.
coli_user_prompt() in glm.c now uses it for both COLI_PROMPT and PROMPT.

Verified: tiny-oracle self-test still 32/32 after rebuild, and a real run
against the full GLM-5.2-int4 model with a Cyrillic prompt now round-trips
correctly end to end (input echoed intact, coherent Cyrillic output),
where it previously produced replacement-character garbage on both input
and output.
---
 c/compat.h | 31 +++++++++++++++++++++++++++++++
 c/glm.c    |  4 ++--
 2 files changed, 33 insertions(+), 2 deletions(-)

diff --git a/c/compat.h b/c/compat.h
index 82de8b6..06779e0 100644
--- a/c/compat.h
+++ b/c/compat.h
@@ -304,8 +304,39 @@ 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)
+
 #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/glm.c b/c/glm.c
index 00c7076..b0e419c 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -5701,9 +5701,9 @@ static void cap_for_ram(Model *m, double ram_gb, int ebits, int max_ctx){
  * self-test, and would "generate" from "$P$G". So on Windows a PROMPT carrying
  * cmd's $-metacodes is ignored; set COLI_PROMPT to pass a real prompt from cmd. */
 static const char *coli_user_prompt(void){
-    const char *p = getenv("COLI_PROMPT");
+    const char *p = getenv_utf8("COLI_PROMPT");
     if(p) return p;
-    p = getenv("PROMPT");
+    p = getenv_utf8("PROMPT");
 #ifdef _WIN32
     if(p) for(const char *q=p; q[0]; q++)
         if(q[0]=='$' && q[1] && strchr("ABCDEFGHLNPQSTV_+|$", q[1]&~0x20)){ p=NULL; break; }

From 7f70a8db5d92516c267f867d41e0a12bcd0021c8 Mon Sep 17 00:00:00 2001
From: KingIcyCreamProjects <safereesekapellen@gmail.com>
Date: Fri, 17 Jul 2026 14:08:50 -0500
Subject: [PATCH 05/31] sampling: guard against non-finite logits (was silent
 token-0 spew)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

On the default serve path (TEMP>0, 0<NUCLEUS<1) a single NaN or +Inf in
the logits — a bad streamed expert tile, or an fp overflow in the matmul
at a low-RAM eviction boundary — poisoned softmax: g_pbuf became all-NaN,
dist_sample never satisfied cum>=u, and the fallback returned token 0. The
engine then emitted an unbroken run of token 0 with NO error. The greedy
path was equally blind: argmax_v started bv=lo[0] and `lo[i]>NaN` is always
false, so a NaN at index 0 pinned the argmax to 0.

- argmax_v: skip NaN (x==x) and seed from -inf, so it returns the max
  finite/+Inf entry instead of being NaN-pinned to 0. Covers greedy decode
  and the speculative-verify argmax path.
- dist_build: after the softmax sum, if s is non-finite or <=0, collapse
  g_pbuf to a one-hot over the finite argmax and warn once, instead of
  dividing every entry into NaN. Covers the nucleus and verify paths.

Both are O(1)/free on the happy path (one branch after the existing loop;
one extra comparison inside the existing argmax loop). Degrade + diagnose,
never silently corrupt.

test_logit_nan (wired into TEST_BINS): asserts argmax_v skips NaN/picks
+Inf, dist_build yields a finite normalized one-hot on the max finite
logit, dist_sample emits that token (not 0), and clean logits still give a
valid distribution. Fails on stock dev, passes with this change.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 c/Makefile               |  5 +++-
 c/glm.c                  | 15 ++++++++++-
 c/tests/test_logit_nan.c | 55 ++++++++++++++++++++++++++++++++++++++++
 3 files changed, 73 insertions(+), 2 deletions(-)
 create mode 100644 c/tests/test_logit_nan.c

diff --git a/c/Makefile b/c/Makefile
index feb3045..062ef75 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_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
@@ -331,6 +331,9 @@ tests/bench_topp$(EXE): tests/bench_topp.c glm.c st.h uring.h json.h tok.h tok_u
 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/glm.c b/c/glm.c
index 92a4d78..9125587 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -4151,7 +4151,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;i<V;i++) if(lo[i]>bv){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;i<V;i++){ float x=lo[i]; if(x==x && x>bv){ bv=x; b=i; } }
+    return b<0?0:b;
 }
 
 /* ---- METODO F: draft grammaticale (#48) ----
@@ -4258,6 +4262,15 @@ static void dist_build(const float *lo, int V){
     float mx=lo[0]; for(int i=1;i<V;i++) if(lo[i]>mx) mx=lo[i];
     double s=0; float invt=1.f/(g_temp>1e-4f?g_temp:1e-4f);
     for(int i=0;i<V;i++){ g_pbuf[i]=expf((lo[i]-mx)*invt); s+=g_pbuf[i]; }
+    /* A single NaN/+Inf logit (bad streamed expert tile, matmul overflow at an
+     * eviction boundary) makes s non-finite; the old code then divided every entry
+     * into NaN and dist_sample silently returned token 0 forever. Collapse to a
+     * one-hot over the max finite logit and warn once — degrade, don't corrupt. */
+    if(!(s==s) || s<=0.0 || s>1e300){
+        static int warned=0; if(!warned){ warned=1; fprintf(stderr,"[sample] non-finite logits; emitting argmax of finite entries\n"); }
+        int b=argmax_v(lo,V); for(int i=0;i<V;i++) g_pbuf[i]=(i==b)?1.f:0.f;
+        return;
+    }
     for(int i=0;i<V;i++) g_pbuf[i]/=(float)s;
     if(g_nuc>0 && g_nuc<1.f){
         for(int i=0;i<V;i++) g_pidx[i]=i;
diff --git a/c/tests/test_logit_nan.c b/c/tests/test_logit_nan.c
new file mode 100644
index 0000000..036c3ea
--- /dev/null
+++ b/c/tests/test_logit_nan.c
@@ -0,0 +1,55 @@
+/* Regression for non-finite-logit poisoning of sampling.
+ *
+ * A single NaN or +Inf in the logits (a bad streamed expert tile, or an fp
+ * overflow in the matmul at a low-RAM eviction boundary) used to make softmax
+ * produce an all-NaN g_pbuf; dist_sample then never satisfied cum>=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<NUCLEUS<1).
+ *
+ * Fix under test: argmax_v() skips NaN (picks the max finite/+Inf entry instead
+ * of being pinned to index 0), and dist_build() detects a non-finite softmax sum
+ * and collapses to a one-hot over the finite argmax (warning once) rather than
+ * dividing every entry into NaN. Degrade + diagnose, never silently corrupt.
+ *
+ * No model file needed: exercises argmax_v / dist_build / dist_sample directly. */
+#include <assert.h>
+#include <math.h>
+#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"); }
+
+    /* --- 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;
+}

From e2d39abd2d4951606ca5c65fe025dc5cf8ea0fec Mon Sep 17 00:00:00 2001
From: KingIcyCreamProjects <safereesekapellen@gmail.com>
Date: Fri, 17 Jul 2026 14:39:14 -0500
Subject: [PATCH 06/31] sampling: NaN-skip the mx scan too (review follow-up on
 #369)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Per review: the collapse starts one line before the sum — seeding
mx=lo[0] means a NaN at index 0 makes mx NaN, every (lo[i]-mx) NaN, and
the softmax is doomed at the max-finding, not the normalize. Seed mx
from -INFINITY and skip NaNs (x==x), mirroring the argmax_v change; if
nothing finite survives, fall back to mx=0 and let the post-sum guard
decide. The isfinite(s) guard is now the second line of defense rather
than the only one.

Clean logits take a byte-identical path (the extra x==x compare is
noise next to V expf calls). test_logit_nan gains the NaN-at-index-0
and all-NaN dist_build cases; test_topp's 123-case sweep still passes
on this tree, confirming no interaction with the #354 heap select.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 c/glm.c                  |  8 +++++++-
 c/tests/test_logit_nan.c | 15 +++++++++++++++
 2 files changed, 22 insertions(+), 1 deletion(-)

diff --git a/c/glm.c b/c/glm.c
index 9125587..54dd7fa 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -4259,7 +4259,13 @@ 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;i<V;i++) if(lo[i]>mx) mx=lo[i];
+    /* NaN-skip the max scan (x==x rules out NaN): seeding mx=lo[0] let a NaN at
+     * index 0 poison every (lo[i]-mx) before the sum — the softmax was doomed at
+     * the max-finding, not the normalize. The post-sum guard below stays as the
+     * second line of defense. */
+    float mx=-INFINITY;
+    for(int i=0;i<V;i++) if(lo[i]==lo[i] && lo[i]>mx) mx=lo[i];
+    if(!isfinite(mx)) mx=0.f;   /* all-NaN (or +Inf) logits: let the guard below decide */
     double s=0; float invt=1.f/(g_temp>1e-4f?g_temp:1e-4f);
     for(int i=0;i<V;i++){ g_pbuf[i]=expf((lo[i]-mx)*invt); s+=g_pbuf[i]; }
     /* A single NaN/+Inf logit (bad streamed expert tile, matmul overflow at an
diff --git a/c/tests/test_logit_nan.c b/c/tests/test_logit_nan.c
index 036c3ea..981d74d 100644
--- a/c/tests/test_logit_nan.c
+++ b/c/tests/test_logit_nan.c
@@ -41,6 +41,21 @@ int main(void){
       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);

From 786f98d4719b80dd353d9c5e3a09ebaafe977bba Mon Sep 17 00:00:00 2001
From: ZacharyZcR <zacharyzcr1984@gmail.com>
Date: Sat, 18 Jul 2026 03:50:48 +0800
Subject: [PATCH 07/31] cuda: load ragged attention entry point on Windows

---
 c/backend_loader.c | 13 +++++++++++++
 1 file changed, 13 insertions(+)

diff --git a/c/backend_loader.c b/c/backend_loader.c
index b743d1b..bbcb8ca 100644
--- a/c/backend_loader.c
+++ b/c/backend_loader.c
@@ -61,6 +61,9 @@ 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 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 +110,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 +204,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 +347,14 @@ 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 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,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);

From c3a90eca36fa3387e78a8f48a31e5b34a5ff1bdf Mon Sep 17 00:00:00 2001
From: ZacharyZcR <zacharyzcr1984@gmail.com>
Date: Sat, 18 Jul 2026 04:54:59 +0800
Subject: [PATCH 08/31] profile CPU GPU tier execution costs

---
 c/glm.c                                | 49 ++++++++++++++++++++------
 c/tests/test_benchmark_cuda_fixture.py |  6 +++-
 c/tools/benchmark_cuda_fixture.py      | 28 +++++++++++++--
 3 files changed, 69 insertions(+), 14 deletions(-)

diff --git a/c/glm.c b/c/glm.c
index 92a4d78..8ec64a1 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -201,7 +201,9 @@ typedef struct {
     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 */
+                                                 /* 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 +309,16 @@ 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; uint64_t hits,miss,ereq,n_fw,n_emit,nlat,n_p2p;
 } 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->n_fw=m->n_fw; b->n_emit=m->n_emit; b->nlat=g_prof_nlat; b->n_p2p=m->n_p2p;
 }
 
 static float *falloc(int64_t n){
@@ -2848,6 +2851,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
         if(!rank_buf||!rank_w){ free(rank_buf); free(rank_w); rank_buf=NULL; rank_w=NULL; do_cache_route=0; }
     }
     /* ---- FASE A: routing di tutte le S posizioni ---- */
+    double route_t0=g_prof?now_s():0;
     int *idxs=malloc((size_t)S*K*sizeof(int)); float *ws=malloc((size_t)S*K*sizeof(float));
     int *keff=malloc(S*sizeof(int));
     /* router in UN matmul batch: stessa matematica, via le S chiamate S=1 */
@@ -2997,6 +3001,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
         for(int d=0;d<D;d++) out[(int64_t)s*D+d]=0;
     }
     free(rank_buf); free(rank_w);
+    if(g_prof)m->t_route+=now_s()-route_t0;
     if(g_route_fp) g_route_call++;
     if(g_couple && cp_pred && S<=8)
         for(int s2=0;s2<S;s2++) couple_prefetch(m,layer,idxs+(int64_t)s2*K,keff[s2]);
@@ -3263,7 +3268,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;r<nr;r++){ float *os=out+(int64_t)rows[r]*D,wgt=rw[r],*hr=hh+(int64_t)r*D;
                     for(int d=0;d<D;d++) os[d]+=wgt*hr[d]; }
-                m->t_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 +3277,7 @@ 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;r<nr;r++){ float *os=out+(int64_t)rows[r]*D, wgt=rw[r], *hr=hh+(int64_t)r*D;
                 for(int d=0;d<D;d++) os[d]+=wgt*hr[d]; }
-            m->t_emm += now_s()-t0;
+            double dt=now_s()-t0;m->t_emm+=dt;if(g_prof)m->t_ecpu+=dt;
         }
 #ifdef COLI_CUDA
         ColiCudaTensor *dev_g[COLI_CUDA_MAX_DEVICES][64],*dev_u[COLI_CUDA_MAX_DEVICES][64];
@@ -3280,6 +3285,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;di<g_cuda_ndev;di++) for(int q=0;q<ngroup;q++)
             if(group_e[q]->g.cuda_device==g_cuda_devices[di]) dev_total[di]+=group_n[q];
         for(int di=1;di<g_cuda_ndev;di++) dev_off[di]=dev_off[di-1]+dev_total[di-1];
@@ -3296,21 +3302,26 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
         }
         double tg=now_s();
         #pragma omp parallel for if(g_cuda_ndev>1) schedule(static)
-        for(int di=0;di<g_cuda_ndev;di++) if(dev_nc[di])
+        for(int di=0;di<g_cuda_ndev;di++) if(dev_nc[di]){
+            double td=g_prof?now_s():0;
             dev_ok[di]=coli_cuda_expert_group(dev_g[di],dev_u[di],dev_d[di],dev_rows[di],dev_nc[di],
                 group_y+(int64_t)dev_off[di]*D,group_x+(int64_t)dev_off[di]*D);
+            if(g_prof)dev_time[di]=now_s()-td;
+        }
         for(int di=0;di<g_cuda_ndev;di++){
             int off=dev_off[di];
             for(int q=0;q<dev_nc[di];q++){
                 int gi=dev_which[di][q],nr=group_n[gi]; ESlot *e=group_e[gi];
                 if(!dev_ok[di]){
                     for(int r=0;r<nr;r++) memcpy(xg+(int64_t)r*D,x+(int64_t)group_row[(int64_t)gi*S+r]*D,D*sizeof(float));
+                    double tc=g_prof?now_s():0;
                     if(!coli_cuda_expert_mlp(e->g.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->t_ecpu+=now_s()-tc;
                 }
                 float *src=dev_ok[di]?group_y+(int64_t)off*D:hh;
                 for(int r=0;r<nr;r++){ float *os=out+(int64_t)group_row[(int64_t)gi*S+r]*D,wgt=group_weight[(int64_t)gi*S+r];
@@ -3318,6 +3329,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
                 off+=nr;
             }
         }
+        if(g_prof){double mx=0;for(int di=0;di<g_cuda_ndev;di++)if(dev_time[di]>mx)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 +3937,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;
@@ -4558,6 +4574,10 @@ static void profile_print(Model *m, double elapsed){
         edisk_s(),m->t_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 | routed GPU critical %.3fs | router %.3fs | residual P2P %.3fs / %llu hop | orchestration %.3fs\n",
+        m->t_ecpu,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 +4591,7 @@ 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->t_aproj=m->t_acore=m->t_aout=0;
     atomic_store_explicit(&g_edisk_ns,0,memory_order_relaxed);
 }
@@ -4615,11 +4636,19 @@ static void prof_report(Model *m, const ProfBase *b, double elapsed, int tokens,
         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;
+    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=ecpu<egpu?ecpu:egpu;
+    fprintf(f,"[PROF] P0 execution: routed CPU %.3fs | routed GPU critical %.3fs | tier straggler %.2fx | "
+              "router %.3fs | residual P2P %.3fs (%llu hop, %.3f ms/hop) | orchestration %.3fs\n",
+        ecpu,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).");
diff --git a/c/tests/test_benchmark_cuda_fixture.py b/c/tests/test_benchmark_cuda_fixture.py
index a704052..ca317c9 100644
--- a/c/tests/test_benchmark_cuda_fixture.py
+++ b/c/tests/test_benchmark_cuda_fixture.py
@@ -1,11 +1,12 @@
 import unittest
 
-from tools.benchmark_cuda_fixture import parse_output
+from tools.benchmark_cuda_fixture import parse_output, parse_p0
 
 
 SAMPLE = """
 REPLAY decode: 4 tokens | 12.34 tok/s
 PROFILE: expert-disk 1.25s | expert-matmul 2.50s | attention 0.75s | lm_head 0.10s | other -0.05s
+P0-EXEC: routed CPU 1.200s | routed GPU critical 0.150s | router 0.200s | residual P2P 0.030s / 75 hop | orchestration 0.100s
 """
 
 
@@ -19,6 +20,9 @@ class ParseOutputTest(unittest.TestCase):
         with self.assertRaisesRegex(RuntimeError, "benchmark output missing"):
             parse_output("REPLAY decode: 4 tokens | 12.34 tok/s", "engine failed")
 
+    def test_extracts_p0_profile(self):
+        self.assertEqual(parse_p0(SAMPLE), [1.2, 0.15, 0.2, 0.03, 75.0, 0.1])
+
 
 if __name__ == "__main__":
     unittest.main()
diff --git a/c/tools/benchmark_cuda_fixture.py b/c/tools/benchmark_cuda_fixture.py
index bced18e..849bb55 100644
--- a/c/tools/benchmark_cuda_fixture.py
+++ b/c/tools/benchmark_cuda_fixture.py
@@ -17,6 +17,11 @@ PROFILE_RE = re.compile(
     r"\| attention ([0-9.]+)s .* lm_head ([0-9.]+)s \| other ([0-9.-]+)s"
 )
 PROFILE_KEYS = ("disk", "expert_matmul", "attention", "lm_head", "other")
+P0_RE = re.compile(
+    r"P0-EXEC: routed CPU ([0-9.]+)s \| routed GPU critical ([0-9.]+)s \| "
+    r"router ([0-9.]+)s \| residual P2P ([0-9.]+)s / ([0-9]+) hop \| orchestration ([0-9.]+)s"
+)
+P0_KEYS = ("routed_cpu", "routed_gpu_critical", "router", "p2p", "p2p_hops", "orchestration")
 
 
 def parse_output(stdout: str, stderr: str = "") -> tuple[float, list[float]]:
@@ -30,11 +35,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 +77,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 +102,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 +121,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))
 

From 946fcd4f9f9d7c6558f8704f2393eb97910abba7 Mon Sep 17 00:00:00 2001
From: bokiko <bokiko@gmail.com>
Date: Sat, 18 Jul 2026 00:19:39 +0300
Subject: [PATCH 09/31] =?UTF-8?q?glm:=20fmt=3D4=20support=20in=20qt=5Faddr?=
 =?UTF-8?q?ow/qt=5Fmatvec=5Frows=20=E2=80=94=20CPU=20absorb=20path=20decod?=
 =?UTF-8?q?ed=20grouped=20int4=20as=20int2?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

An all-grouped container (kv_b_proj at fmt=4) generates one correct token
and then EOS: qt_addrow and qt_matvec_rows handle fmt 0/1/2 and fall
through to the int2 decoder, so grouped-int4 kv_b was unpacked as 2-bit
pairs under a per-row scale that does not exist in the [O,ng] layout.
Prefill (S>4, reconstruction) is unaffected, which made the failure look
like an EOS bug rather than an attention bug.

Same class as #298 (CUDA absorb kernels missing fmt=4), CPU side. Existing
containers escape it because the recommended mixed-precision recipe keeps
kv_b at int8 (#237).

Adds per-group branches mirroring matmul_i4_grouped semantics. fmt 0/1/2/3
paths are untouched.

Co-Authored-By: Claude <noreply@anthropic.com>
---
 c/glm.c | 19 +++++++++++++++++++
 1 file changed, 19 insertions(+)

diff --git a/c/glm.c b/c/glm.c
index 92a4d78..c869b9c 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -2284,6 +2284,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;i<I;i++) acc[i]+=coef*w[i]; return; }
+    /* fmt=4 PRIMA del calcolo di c: s[] e' [O,ng] per-gruppo, s[row] sarebbe la scala
+     * sbagliata. Senza questo ramo il fall-through int2 decodificava i nibble int4 come
+     * coppie di valori a 2 bit — lo stesso bug di #298 sui kernel absorb CUDA, lato CPU.
+     * EN: fmt=4 BEFORE computing c: s[] is [O,ng] per-group, s[row] would be the wrong
+     * scale. Without this branch the int2 fall-through decoded int4 nibbles as pairs of
+     * 2-bit values — the same bug #298 fixed in the CUDA absorb kernels, CPU side. */
+    if(t->fmt==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<I;i+=2){ uint8_t b=w[i>>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;i<I;i++) acc[i]+=c*(float)w[i]; return; }
     if(t->fmt==2){ const uint8_t *w=t->q4+(int64_t)row*((I+1)/2);
@@ -2302,6 +2314,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<I;i+=2){ uint8_t b=w[i>>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*gs<I; g++){ int base=g*gs, end=base+gs>I?I:base+gs; float acc=0;
+                for(int i=base;i<end;i++){ uint8_t b=w[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<I;i++){ uint8_t b=w[i>>2]; acc+=((int)((b>>((i&3)*2))&3)-2)*x[i]; } a=acc*s; }
         y[j]=(float)a;

From 73e0e5ec82e98ef40bcfab5abd42c6776e6a381e Mon Sep 17 00:00:00 2001
From: KingIcyCreamProjects <safereesekapellen@gmail.com>
Date: Fri, 17 Jul 2026 16:24:33 -0500
Subject: [PATCH 10/31] docs(api): add "Connect a coding CLI or editor" recipe
 (#373)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Documents connecting OpenAI-compatible coding CLIs to `coli serve`.
Covers the common snag reported in #373 — clients like crush refuse to
start without an API key even though the local endpoint needs none — by
showing that any dummy key works (Colibri only enforces COLI_API_KEY if
set). Concrete recipes for aider and crush, a curl smoke test, the
generic base-URL/model/key pattern for other tools, and an honest
tok/s-latency caveat for the streaming path.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 docs/api.md | 61 +++++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 61 insertions(+)

diff --git a/docs/api.md b/docs/api.md
index d3b09f1..cb29eb2 100644
--- a/docs/api.md
+++ b/docs/api.md
@@ -49,6 +49,67 @@ 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).
+
+> On the CPU-streaming path a large model decodes at roughly 1 tok/s, so
+> interactive agent loops will feel slow — it connects and works, but the latency
+> is very different from a hosted model.
+
 ## Isolated KV contexts
 
 `coli serve --kv-slots N` allocates up to 16 independent sequence contexts.

From 570d738ed5c7ae8a43aed5a4eb23b8948c48c41f Mon Sep 17 00:00:00 2001
From: ZacharyZcR <zacharyzcr1984@gmail.com>
Date: Sat, 18 Jul 2026 05:34:41 +0800
Subject: [PATCH 11/31] profile effective CPU expert bandwidth

---
 c/glm.c                                | 23 +++++++++++++++++------
 c/tests/test_benchmark_cuda_fixture.py |  4 ++--
 c/tools/benchmark_cuda_fixture.py      |  5 +++--
 3 files changed, 22 insertions(+), 10 deletions(-)

diff --git a/c/glm.c b/c/glm.c
index 8ec64a1..f5dd33b 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -203,6 +203,7 @@ typedef struct {
     double route_kl_sum; uint64_t route_kl_n;     /* mean KL(true||chosen) on gate mass */
     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) */
@@ -310,7 +311,7 @@ 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,ecpu,egpu,route,p2p,attn,head;
-    int64_t io; uint64_t hits,miss,ereq,n_fw,n_emit,nlat,n_p2p;
+    int64_t io,cpu_bytes; uint64_t hits,miss,ereq,n_fw,n_emit,nlat,n_p2p,cpu_rows;
 } ProfBase;
 static void prof_base(Model *m, ProfBase *b){
     b->edisk=edisk_s(); b->ewait=m->t_ewait; b->emm=m->t_emm;
@@ -319,6 +320,7 @@ static void prof_base(Model *m, ProfBase *b){
     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->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){
@@ -3277,7 +3279,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;r<nr;r++){ float *os=out+(int64_t)rows[r]*D, wgt=rw[r], *hr=hh+(int64_t)r*D;
                 for(int d=0;d<D;d++) os[d]+=wgt*hr[d]; }
-            double dt=now_s()-t0;m->t_emm+=dt;if(g_prof)m->t_ecpu+=dt;
+            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];
@@ -3320,6 +3324,8 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
                         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;
                 }
@@ -4574,8 +4580,9 @@ static void profile_print(Model *m, double elapsed){
         edisk_s(),m->t_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 | routed GPU critical %.3fs | router %.3fs | residual P2P %.3fs / %llu hop | orchestration %.3fs\n",
-        m->t_ecpu,m->t_egpu,m->t_route,m->t_p2p,(unsigned long long)m->n_p2p,
+    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
@@ -4592,6 +4599,7 @@ 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);
 }
@@ -4639,15 +4647,18 @@ static void prof_report(Model *m, const ProfBase *b, double elapsed, int tokens,
     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=ecpu<egpu?ecpu:egpu;
-    fprintf(f,"[PROF] P0 execution: routed CPU %.3fs | routed GPU critical %.3fs | tier straggler %.2fx | "
+    fprintf(f,"[PROF] P0 execution: routed CPU %.3fs / %.2f GB/s (%llu row) | routed GPU critical %.3fs | tier straggler %.2fx | "
               "router %.3fs | residual P2P %.3fs (%llu hop, %.3f ms/hop) | orchestration %.3fs\n",
-        ecpu,egpu,fast>1e-9?slow/fast:0.0,route,p2p,(unsigned long long)np2p,
+        ecpu,ecpu>0?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);
diff --git a/c/tests/test_benchmark_cuda_fixture.py b/c/tests/test_benchmark_cuda_fixture.py
index ca317c9..3feaf11 100644
--- a/c/tests/test_benchmark_cuda_fixture.py
+++ b/c/tests/test_benchmark_cuda_fixture.py
@@ -6,7 +6,7 @@ from tools.benchmark_cuda_fixture import parse_output, parse_p0
 SAMPLE = """
 REPLAY decode: 4 tokens | 12.34 tok/s
 PROFILE: expert-disk 1.25s | expert-matmul 2.50s | attention 0.75s | lm_head 0.10s | other -0.05s
-P0-EXEC: routed CPU 1.200s | routed GPU critical 0.150s | router 0.200s | residual P2P 0.030s / 75 hop | orchestration 0.100s
+P0-EXEC: routed CPU 1.200s / 123.40 GB/s (456 row) | routed GPU critical 0.150s | router 0.200s | residual P2P 0.030s / 75 hop | orchestration 0.100s
 """
 
 
@@ -21,7 +21,7 @@ class ParseOutputTest(unittest.TestCase):
             parse_output("REPLAY decode: 4 tokens | 12.34 tok/s", "engine failed")
 
     def test_extracts_p0_profile(self):
-        self.assertEqual(parse_p0(SAMPLE), [1.2, 0.15, 0.2, 0.03, 75.0, 0.1])
+        self.assertEqual(parse_p0(SAMPLE), [1.2, 123.4, 456.0, 0.15, 0.2, 0.03, 75.0, 0.1])
 
 
 if __name__ == "__main__":
diff --git a/c/tools/benchmark_cuda_fixture.py b/c/tools/benchmark_cuda_fixture.py
index 849bb55..26512dd 100644
--- a/c/tools/benchmark_cuda_fixture.py
+++ b/c/tools/benchmark_cuda_fixture.py
@@ -18,10 +18,11 @@ PROFILE_RE = re.compile(
 )
 PROFILE_KEYS = ("disk", "expert_matmul", "attention", "lm_head", "other")
 P0_RE = re.compile(
-    r"P0-EXEC: routed CPU ([0-9.]+)s \| routed GPU critical ([0-9.]+)s \| "
+    r"P0-EXEC: routed CPU ([0-9.]+)s / ([0-9.]+) GB/s \(([0-9]+) row\) \| routed GPU critical ([0-9.]+)s \| "
     r"router ([0-9.]+)s \| residual P2P ([0-9.]+)s / ([0-9]+) hop \| orchestration ([0-9.]+)s"
 )
-P0_KEYS = ("routed_cpu", "routed_gpu_critical", "router", "p2p", "p2p_hops", "orchestration")
+P0_KEYS = ("routed_cpu", "routed_cpu_gb_s", "routed_cpu_rows", "routed_gpu_critical",
+           "router", "p2p", "p2p_hops", "orchestration")
 
 
 def parse_output(stdout: str, stderr: str = "") -> tuple[float, list[float]]:

From 153ee16f61b50633c6ce30008ce56c13aea49069 Mon Sep 17 00:00:00 2001
From: woolcoxm <13604288+woolcoxm@users.noreply.github.com>
Date: Fri, 17 Jul 2026 18:17:29 -0400
Subject: [PATCH 12/31] =?UTF-8?q?bench:=20multi-seed=20bench=5Fdsa=5Fselec?=
 =?UTF-8?q?t=20=E2=80=94=20kill=20single-input=20pivot-luck=20spike=20(#35?=
 =?UTF-8?q?7)?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

@KingIcyCreamProjects caught a real artifact on the 9950X3D (#357 thread): the
nk=8192 cell reported ~75x, far above the real ~13-40x algorithm. Root cause was
two compounding bench bugs, both verified against the code:

  1. ONE frozen input per cell. partial_select_desc's median-of-three pivot is
     fully deterministic (no RNG), and bench_dsa_select froze the input then ran
     2000 reps on that identical array -- so all 2000 reps hit the exact same
     pivot sequence. A single lucky input spiked one nk row (stable across re-runs
     because it's deterministic, not because it's real).
  2. brng was a static global, initialized once and NEVER reset between cells.
     So each (shape,nk) cell's input depended on every prior cell's brand() draws
     -- reordering nks[] or adding a shape silently shifted all later inputs.

Fix:
  - brng_seed() reseeds per (shape, nk, seed) so every cell is reproducible and
    independent of cell ordering.
  - Report the MEDIAN of N_SEEDS=11 independent inputs, each itself a median over
    N_REPEAT=2000 timing reps. A lucky pivot now moves one of 11 samples, not the
    reported number.

Measured on Intel Core Ultra 9 185H (this fix, median of 11 seeds):
  realistic@8192  10.97x   (was 13.57x single-seed on this box)
  uniform@8192     9.39x   (was 7.48x)
  plateau@8192    16.11x   (plateau ignores the RNG, unchanged as expected)
  band: 6-26x across all cells, no anomalous spikes.

Correctness is unaffected (test_dsa_select, 129 cases, unchanged). This is a bench
fidelity fix only -- no engine code touched.
---
 c/tests/bench_dsa_select.c | 61 +++++++++++++++++++++++++++++---------
 1 file changed, 47 insertions(+), 14 deletions(-)

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<n;s++){ double k=a[s]; int b=s-1;
+        while(b>=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<nk;i++) isc[i]=(float)(-(double)(i/7));
 }
 
+/* Multi-seed aggregation. Per @KingIcyCreamProjects (#357 thread): with a single frozen
+ * input per cell, quickselect's deterministic median-of-three pivot means one lucky input
+ * can spike a single nk row (an observed ~75x at nk=8192 on a 9950X3D that was really a
+ * ~13-40x algorithm). Two bugs compounded it: (1) the old bench drew ONE input per cell;
+ * (2) brng was never reset, so each cell's input depended on every prior cell's draws --
+ * reordering nks[] silently shifted all later inputs. Both fixed here: brng is reseeded
+ * per (shape, nk, seed), and we take the MEDIAN of N_SEEDS independent inputs, each itself
+ * a median over N_REPEAT timing reps. A lucky pivot now moves one of the N_SEEDS samples,
+ * not the reported number. */
+
 int main(void){
     int keep = 2048;   /* GLM-5.2 index_topk */
     int nks[] = {2049, 4096, 8192, 16384, 32768, 65536};
+    const int N_SEEDS = 11;   /* odd so the median is a real sample, not interpolated */
     float *isc = malloc((size_t)65536*sizeof(float));
     int *dst   = malloc((size_t)65536*sizeof(int)); int nd;
+    double *old_seeds = malloc((size_t)N_SEEDS*sizeof(double));
+    double *new_seeds = malloc((size_t)N_SEEDS*sizeof(double));
 
     struct { const char *name; void (*fill)(float*,int); } shapes[] = {
         { "realistic", fill_realistic },
@@ -99,27 +119,40 @@ int main(void){
         { "plateau",   fill_plateau   },
     };
 
-    printf("bench_dsa_select: DSA top-keep, old (qsort) vs new (partial-select)  keep=%d\n", keep);
+    printf("bench_dsa_select: DSA top-keep, old (qsort) vs new (partial-select)  keep=%d  (median of %d seeds x %d reps)\n",
+           keep, N_SEEDS, N_REPEAT);
     printf("%-12s %7s %14s %14s %9s\n", "shape", "nk", "old ns/call", "new ns/call", "speedup");
     printf("------------------------------------------------------------------------\n");
 
     for(size_t sh=0; sh<sizeof(shapes)/sizeof(shapes[0]); sh++){
         for(size_t ni=0; ni<sizeof(nks)/sizeof(nks[0]); ni++){
             int nk=nks[ni];
-            shapes[sh].fill(isc,nk);
-            /* warmup both paths so caches/branch predictors are primed */
-            for(int w=0; w<50; w++){ keep_old(isc,nk,keep,dst,&nd); keep_new(isc,nk,keep,dst,&nd); }
-            /* sanity: both must keep exactly `keep` (correctness is test_dsa_select's
-             * job, but a count divergence here would make the timing meaningless) */
-            keep_old(isc,nk,keep,dst,&nd); int na=nd;
-            keep_new(isc,nk,keep,dst,&nd); int nb=nd;
-            if(na!=keep || nb!=keep){
-                printf("%-12s %7d   (BAD COUNTS: old=%d new=%d, skipped)\n",
-                       shapes[sh].name, nk, na, nb);
+            int bad = 0;
+            for(int sd=0; sd<N_SEEDS; sd++){
+                /* reseed per (shape,nk,seed) so each cell's input is reproducible and
+                 * independent of cell ordering, and so lucky pivots are sampled, not fixed. */
+                brng_seed(0xA5A5A5A5u + (uint32_t)(sd*0x9E3779B9u));
+                shapes[sh].fill(isc,nk);
+                /* warmup both paths so caches/branch predictors are primed */
+                for(int w=0; w<50; w++){ keep_old(isc,nk,keep,dst,&nd); keep_new(isc,nk,keep,dst,&nd); }
+                /* sanity: both must keep exactly `keep` (correctness is test_dsa_select's
+                 * job, but a count divergence here would make the timing meaningless) */
+                keep_old(isc,nk,keep,dst,&nd); int na=nd;
+                keep_new(isc,nk,keep,dst,&nd); int nb=nd;
+                if(na!=keep || nb!=keep){ bad++; continue; }
+                old_seeds[sd] = bench_ns(keep_old,isc,nk,keep);
+                new_seeds[sd] = bench_ns(keep_new,isc,nk,keep);
+            }
+            if(bad == N_SEEDS){
+                printf("%-12s %7d   (BAD COUNTS on all %d seeds, skipped)\n",
+                       shapes[sh].name, nk, bad);
                 continue;
             }
-            double t_old=bench_ns(keep_old,isc,nk,keep);
-            double t_new=bench_ns(keep_new,isc,nk,keep);
+            /* report median-of-seed-medians: robust to a single lucky/unlucky pivot */
+            dsort(old_seeds, N_SEEDS);
+            dsort(new_seeds, N_SEEDS);
+            double t_old = old_seeds[N_SEEDS/2];
+            double t_new = new_seeds[N_SEEDS/2];
             printf("%-12s %7d %14.0f %14.0f %8.2fx\n",
                    shapes[sh].name, nk, t_old, t_new, t_old/t_new);
         }
@@ -127,6 +160,6 @@ int main(void){
     }
 
     printf("bench_dsa_select: done (lower ns is better; speedup = old/new)\n");
-    free(isc); free(dst);
+    free(isc); free(dst); free(old_seeds); free(new_seeds);
     return 0;
 }

From 4fb5b4e9754653fd978b59d4d10f73c08291b885 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sat, 18 Jul 2026 01:10:33 +0200
Subject: [PATCH 13/31] sampling: survive non-finite logits instead of emitting
 token 0 forever (#369)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

A single NaN or +Inf logit silently broke the default sampling path. +Inf
became `mx`, then `expf((Inf-mx))`/`expf((NaN-mx))` is NaN, the softmax sum went
NaN, every probability went NaN — and dist_sample's fallback loop
`if(g_pbuf[i]>0)` is false for NaN at every index, so it returned 0. Every
subsequent token: 0. No error, no warning. @KingIcyCreamProjects found it.

dist_build now takes `mx` over finite logits only, gives a non-finite logit
probability 0, and when the distribution is unusable (no finite logit, or a
non-finite/zero sum) collapses to a delta on the finite argmax and warns ONCE
on stderr — degraded, but a valid token and a visible cause, never a silent
stream of zeros. The finite argmax uses the index found during the mx pass
(robust even when lo[0] itself is NaN, where argmax_v would wrongly return 0).

tests/test_sample_nan.c: healthy logits still sample correctly; NaN/+Inf
injected at lo[0], the middle, and the last position all pick the finite
argmax; an all-non-finite vocab leaves no NaN in the buffer and doesn't crash.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 c/Makefile                |  5 +++-
 c/glm.c                   | 23 +++++++++++++--
 c/tests/test_sample_nan.c | 62 +++++++++++++++++++++++++++++++++++++++
 3 files changed, 87 insertions(+), 3 deletions(-)
 create mode 100644 c/tests/test_sample_nan.c

diff --git a/c/Makefile b/c/Makefile
index feb3045..7210cbb 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)
 ifneq (,$(LINUX))
 TEST_BINS += tests/test_uring$(EXE)
 endif
@@ -328,6 +328,9 @@ 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)
 
diff --git a/c/glm.c b/c/glm.c
index bf8acb1..2ee12ed 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -4296,9 +4296,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;i<V;i++) if(lo[i]>mx) 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;i<V;i++) if(isfinite(lo[i]) && (mxi<0 || lo[i]>mx)){ 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<V;i++){ g_pbuf[i]=expf((lo[i]-mx)*invt); s+=g_pbuf[i]; }
+    if(mxi>=0){
+        for(int i=0;i<V;i++){ g_pbuf[i]=isfinite(lo[i])?expf((lo[i]-mx)*invt):0.f; s+=g_pbuf[i]; }
+    }
+    if(mxi<0 || !isfinite(s) || s<=0.0){          /* distribuzione inutilizzabile */
+        static int warned=0;
+        if(!warned){ warned=1; fprintf(stderr,
+            "[SAMPLE] warning: non-finite logits (NaN/Inf) — falling back to argmax; "
+            "output may be degraded. This usually means a numerical blow-up upstream.\n"); }
+        int a=(mxi>=0)?mxi:0;                       /* mxi = argmax dei logit FINITI (robusto
+                                                     * anche se lo[0] e' NaN, dove argmax_v fallirebbe) */
+        for(int i=0;i<V;i++) g_pbuf[i]=0.f; g_pbuf[a]=1.f;
+        return;                                    /* delta su un token valido, niente top-p su NaN */
+    }
     for(int i=0;i<V;i++) g_pbuf[i]/=(float)s;
     if(g_nuc>0 && g_nuc<1.f){
         for(int i=0;i<V;i++) g_pidx[i]=i;
diff --git a/c/tests/test_sample_nan.c b/c/tests/test_sample_nan.c
new file mode 100644
index 0000000..fc33268
--- /dev/null
+++ b/c/tests/test_sample_nan.c
@@ -0,0 +1,62 @@
+/* Regressione #369: un solo logit NaN o +Inf faceva emettere a dist_build/dist_sample
+ * il token 0 PER SEMPRE, in silenzio (il fallback `g_pbuf[i]>0` 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 <stdio.h>
+
+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<V;i++) if(!isfinite(g_pbuf[i])) return 1; return 0; }
+
+int main(void){
+    int fail=0, V=2000;
+    float *lo=malloc(V*sizeof(float));
+    g_temp=0.7f; g_nuc=0.90f;                     /* il path serve di default */
+
+    /* (a) logit sani: il campionato deve avere prob > 0 e nessun NaN nel buffer */
+    for(int i=0;i<V;i++) lo[i]=(float)((int)(xr()%2000)-1000)/100.f;
+    int known=1337; lo[known]=50.f;               /* picco netto */
+    dist_build(lo,V);
+    if(pbuf_has_nan(V)){ printf("  FAIL: NaN in g_pbuf con logit sani\n"); fail=1; }
+    if(g_pbuf[known]<=0.f){ printf("  FAIL: il picco ha prob 0\n"); fail=1; }
+    if(!fail) printf("  logit sani: distribuzione valida, picco vivo            ok\n");
+
+    /* (b) NaN/+Inf iniettato in varie posizioni; il finito-argmax deve vincere */
+    float bad[3]; bad[0]=NAN; bad[1]=INFINITY; bad[2]=-INFINITY;
+    const char *bn[3]={"NaN","+Inf","-Inf"};
+    for(int b=0;b<3;b++){
+        for(int pos=0;pos<3;pos++){               /* lo[0], meta', ultimo */
+            for(int i=0;i<V;i++) lo[i]=(float)((int)(xr()%400)-200)/100.f;
+            int amax=777; lo[amax]=9.0f;           /* massimo FINITO atteso */
+            int at = pos==0?0 : pos==1?V/2 : V-1;
+            if(at==amax) amax=amax+1;              /* non sovrapporre */
+            lo[amax]=9.0f;
+            lo[at]=bad[b];                         /* veleno */
+            dist_build(lo,V);
+            if(pbuf_has_nan(V)){ printf("  FAIL: NaN sopravvive (%s @ %d)\n",bn[b],at); fail=1; continue; }
+            /* con -Inf il fallback puo' non scattare (max finito resta), ma il buffer
+             * deve restare valido e sommare ~1: con +Inf/NaN scatta il delta su amax */
+            int picked=-1; float pv=-1;
+            for(int i=0;i<V;i++) if(g_pbuf[i]>pv){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<V;i++) lo[i]=NAN;
+    dist_build(lo,V);
+    if(pbuf_has_nan(V)){ printf("  FAIL: tutti-NaN lascia NaN nel buffer\n"); fail=1; }
+    else printf("  tutti non-finiti: nessun crash, buffer valido               ok\n");
+
+    printf(fail?"test_sample_nan: FAIL\n":"test_sample_nan: ok\n");
+    return fail;
+}

From eb6d4325dc90b51aa9c14f6499d22f90bc18dcd9 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sat, 18 Jul 2026 01:15:48 +0200
Subject: [PATCH 14/31] docs: collapse WINDOWS.md/windows.md case collision
 into one file (#371)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

docs/WINDOWS.md (the #329 install walkthrough) and docs/windows.md (the concise
reference from the README restructure) differ only in case. On case-insensitive
filesystems — every default Windows and macOS checkout — the two paths collide:
the working tree can never be clean, and `git rebase`/`git status` refuse to
run. @KingIcyCreamProjects reported it.

Everything links to the lowercase docs/windows.md (README.md, README.zh-TW.md,
docs/cuda.md), so that is the canonical name. This keeps it and folds in both
contents: the step-by-step walkthrough (download-first, Smart App Control, CUDA
DLL, first run, reference numbers, failure index) followed by the build-flags
reference (AVX-VNNI, warmup). docs/WINDOWS.md is removed. No links change.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 docs/WINDOWS.md | 100 -----------------------------------------
 docs/windows.md | 116 +++++++++++++++++++++++++++++++++++++++++-------
 2 files changed, 99 insertions(+), 117 deletions(-)
 delete mode 100644 docs/WINDOWS.md

diff --git a/docs/WINDOWS.md b/docs/WINDOWS.md
deleted file mode 100644
index 52d83cd..0000000
--- a/docs/WINDOWS.md
+++ /dev/null
@@ -1,100 +0,0 @@
-# Windows 11 native install — a complete walkthrough (no WSL)
-
-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.
-
-## 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
-python -m pip install -U "huggingface_hub[hf_transfer]"
-$env:HF_HUB_ENABLE_HF_TRANSFER = "1"
-hf download <model-repo> --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: <your GPU>, ... sm_XX` and `[CUDA] mode: routed experts + resident dense tensors`.
-
-## 4. First run
-
-```powershell
-cd <repo>\c
-$env:OMP_NUM_THREADS = "<physical cores>"
-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/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 <model-repo> --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: <your GPU>, ... sm_XX` and `[CUDA] mode: routed experts + resident dense tensors`.
+
+## 4. First run
+
+```powershell
+cd <repo>\c
+$env:OMP_NUM_THREADS = "<physical cores>"
+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

From db09466d3dc95015c77c8953d083d8d568f0657d Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sat, 18 Jul 2026 01:24:43 +0200
Subject: [PATCH 15/31] convert(fp8->int4): --mtp/--indexer respected on the
 --indir path, no more container overwrite (#355)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

The local (`--indir`) branch of main() ignored --mtp and --indexer entirely:
it always called convert_shard() without keep_mtp/keep_idx and always wrote
`out-NNNNN.safetensors`. So the documented two-pass workflow — convert the main
model into an outdir, then run `--mtp` into the SAME outdir for the head — did
the opposite on the local path: with --mtp defaulting ebits to 8 and keep_mtp
staying False, it silently re-converted the whole model to per-row int8 and
overwrote the finished fmt=4 container shard by shard, printing nothing wrong.
@mohamedmastouri2000-boop lost 137 of 141 freshly-converted g64 shards to this
while building the public container for #298/#326.

The --indir branch now mirrors the download path: keep_mtp=a.mtp /
keep_idx=a.indexer passed through, output named out-mtp-/out-idx-/out- by mode,
empty shards skipped (an MTP pass emits only shards containing layer n_layers),
and config/tokenizer copied only on the main pass (the head/idx passes land in
an already-complete outdir).

Verified with a synthetic 2-layer + MTP-shard model: after the main pass, a
second --mtp pass into the same outdir leaves out-00000's md5 BYTE-IDENTICAL
and writes out-mtp-00000 alongside it. The bug would have changed that md5.

Reported-by: mohamedmastouri2000-boop <#355>

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 c/tools/convert_fp8_to_int4.py | 32 +++++++++++++++++++++++++-------
 1 file changed, 25 insertions(+), 7 deletions(-)

diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py
index 382125e..fb2fb73 100644
--- a/c/tools/convert_fp8_to_int4.py
+++ b/c/tools/convert_fp8_to_int4.py
@@ -379,14 +379,32 @@ def main():
     if a.indir:    # conversione locale (test)
         shards = sorted(glob.glob(os.path.join(a.indir, "*.safetensors")))
         from safetensors.numpy import save_file
+        # 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-"
+        n = 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}")
+            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)
+                continue
+            save_file(out, os.path.join(a.outdir, f"{prefix}{n:05d}.safetensors"))
+            n += 1
+        # config/tokenizer solo per la conversione principale — i passaggi mtp/idx
+        # vanno nella stessa outdir di un container gia' completo di metadati.
+        if not a.mtp and not a.indexer:
+            for fn in ["config.json"]:
+                src = os.path.join(a.indir, fn)
+                if os.path.exists(src): shutil.copy(src, a.outdir)
+        tag = "MTP" if a.mtp else "indexer" if a.indexer else "main"
+        print(f"converted {n} {tag} shard(s) -> {a.outdir} ({prefix}NNNNN)")
         return
 
     # reale: scarica shard per shard, converte, cancella

From 679c0742b2ec09a85d9987f85230125658e86d80 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sat, 18 Jul 2026 01:43:50 +0200
Subject: [PATCH 16/31] telemetry: split expert hits into pin-tier vs LRU
 ecache (#336)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

The expert lookup counted a hit identically whether the pinned hot-store or the
LRU ecache served it — both Phase C branches bumped the single `m->hits`. With a
warm pin profile the pin tier absorbs most hot experts, so the ecache could be
serving anywhere from ~0% to most hits and the logs couldn't tell which. That
made every cache-policy question unanswerable, including #223's "at what cap
does the eviction policy start to win?" (a flat A/B can mean "pin absorbed
everything" or "genuine floor" and the lumped counter can't distinguish them).

Two counters `hit_pin`/`hit_ecache` bumped at the two lookup branches (the
existing `m->hits++` stays, so all existing math is unchanged), snapshotted in
ProfBase and reset alongside `hits`. Surfaced in the human-readable summaries:
  decode:  expert hit rate 31.3% (pin 22.1% + lru 9.2%)
  [PROF]:  hit 31.3% (X pin + Y lru / Z load)
  tiny:    hit rate 88.1% (0 pin + 74 lru / 10 miss)
The serve-mux STAT protocol line is untouched (openai_server.py parses it
positionally).

Invariant hit_pin + hit_ecache == hits holds by construction — exactly two
sites bump hits and each also bumps one split counter, nothing else mutates
hits. Verified numerically on the tiny oracle: 0 pin + 74 lru = 74 hits, 88.1%.
Zero cost: two increments on paths that already increment.

Reported-by: KingIcyCreamProjects <#336>

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 c/glm.c | 26 ++++++++++++++++----------
 1 file changed, 16 insertions(+), 10 deletions(-)

diff --git a/c/glm.c b/c/glm.c
index 2ee12ed..1446606 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -196,6 +196,7 @@ 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 */
@@ -312,6 +313,7 @@ static void prof_lat(double s){ g_prof_lat[g_prof_nlat++ % PROF_LAT_CAP]=s; }
 typedef struct {
     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;
@@ -319,6 +321,7 @@ static void prof_base(Model *m, ProfBase *b){
     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->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;
 }
@@ -3132,9 +3135,9 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
         ESlot *use[64]; int missk[64]; int qof[64]; int nmiss=0;
         for(int j=0;j<nb;j++){ int eid=uniq[base+j]; use[j]=NULL; qof[j]=-1;
             ESlot *P=m->pin[layer];
-            for(int z=0;z<m->npin[layer];z++) if(P[z].eid==eid){ m->hits++; use[j]=&P[z]; break; }
+            for(int z=0;z<m->npin[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;z<nn;z++) if(Sl[z].eid==eid){ m->hits++; Sl[z].used=(uint64_t)__atomic_add_fetch(&m->eclock,1,__ATOMIC_RELAXED); use[j]=&Sl[z]; break; } }
+                for(int z=0;z<nn;z++) if(Sl[z].eid==eid){ m->hits++; 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++; } }
         }
@@ -4674,11 +4677,12 @@ 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);
@@ -4724,7 +4728,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;
@@ -4774,7 +4778,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);
@@ -4787,8 +4791,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;i<c->n_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,
@@ -6535,8 +6540,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

From b09273f73d33bc2b18ac1c7a8328a2ad69451ea7 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sat, 18 Jul 2026 01:57:17 +0200
Subject: [PATCH 17/31] serve: MTP status message tells the truth in
 multiplexed mode (#358)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

`coli serve` runs the engine with SERVE_BATCH=1 (openai_server.py), which
selects run_serve_mux, which sets g_draft=0 — speculation isn't ragged-safe
across the multi-slot batch. But the "[MTP] active: native speculative decoding
(draft=N)" line is printed in main() BEFORE the serve path is chosen, so
`coli serve DRAFT=8` announced draft=8 and then silently disabled it.
@LordMZTE reported the misleading message.

The load line and the [MTP] stderr line now detect the mux case (SERVE +
SERVE_BATCH) and report it truthfully: "MTP DISABLED (multiplexed serve)" with
draft=0, plus a one-line explanation that single-client interactive use
(`coli chat`, which spawns run_serve without SERVE_BATCH) keeps MTP. No more
draft=8 claim on a path that runs draft=0.

This is the honesty half of #358. The feature half — MTP inside the HTTP
server for a single client — is a real enhancement but needs engine work: the
mux decode kernel would have to run the speculative path when exactly one KV
slot is active (S=1 is not actually ragged), and it needs a server round-trip
test. Tracked separately; the message no longer lies in the meantime.

Reported-by: LordMZTE <#358>

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 c/glm.c | 15 +++++++++++++--
 1 file changed, 13 insertions(+), 2 deletions(-)

diff --git a/c/glm.c b/c/glm.c
index 1446606..4c4feef 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -6386,11 +6386,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

From 8a1dccae3b1a4c82ea9173e3d1ea8c002f14e785 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sat, 18 Jul 2026 02:33:23 +0200
Subject: [PATCH 18/31] docs(api): tell coding-CLI users about the prefill cost
 before they hit it (#373)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

yurivict connected crush per the walkthrough and got 'thought for 1h8m and did
nothing' — which is not a hang: agent CLIs send a 10-20k-token system preamble,
and prefill on the CPU-streaming path runs at a few tok/s (attention-bound,
#153). An hour of silent prefill looks exactly like a dead server. The note now
spells out the arithmetic, the curl smoke-test that separates slow-but-working
from broken, and honest guidance on what agent workloads are (not) viable.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 docs/api.md | 21 ++++++++++++++++++---
 1 file changed, 18 insertions(+), 3 deletions(-)

diff --git a/docs/api.md b/docs/api.md
index cb29eb2..62010bd 100644
--- a/docs/api.md
+++ b/docs/api.md
@@ -106,9 +106,24 @@ KV configuration actually allows.
 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).
 
-> On the CPU-streaming path a large model decodes at roughly 1 tok/s, so
-> interactive agent loops will feel slow — it connects and works, but the latency
-> is very different from a hosted model.
+> **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
 

From f8e4dc95b5add6cd55105b25731293d2c17fae44 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sat, 18 Jul 2026 12:28:27 +0200
Subject: [PATCH 19/31] serve: honor --ngen when the client omits max_tokens;
 convert: print the resolved plan (#382, #383)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Two operator-surprise fixes:

openai_server.py (#382, LordMZTE): a request without max_tokens defaulted to
min(256, limit) — so `coli serve --ngen 32768` still cut every answer at 256
tokens for clients that omit the field. The operator's configured budget IS the
default now; generation still ends at EOS, so it's a cap, not a target.

convert_fp8_to_int4.py (#383, bokiko): the resolved conversion plan (mode,
source, ebits/io/x bits, grouped-vs-per-row) prints as a [PLAN] line before any
work. The two traps in #383 — --mtp defaulting ebits to 8 with the grouped
branch silently gated off, and --mtp appearing to ignore --indir — are already
defused by the #355 fix (the --mtp pass emits ONLY head tensors on the local
path), but a 3.5-hour job must show its plan at second 1, not in a post-hoc
size sanity check. bokiko's exact invocation now prints:
  [PLAN] mode: MTP head only | source: local ./fp8 | experts 8-bit, ... |
         PER-ROW (grouped branch needs bits<=4; ebits=8 disables it)

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
---
 c/openai_server.py             |  5 ++++-
 c/tools/convert_fp8_to_int4.py | 10 ++++++++++
 2 files changed, 14 insertions(+), 1 deletion(-)

diff --git a/c/openai_server.py b/c/openai_server.py
index d55fbc8..f56e20a 100644
--- a/c/openai_server.py
+++ b/c/openai_server.py
@@ -406,7 +406,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/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py
index fb2fb73..9cd6937 100644
--- a/c/tools/convert_fp8_to_int4.py
+++ b/c/tools/convert_fp8_to_int4.py
@@ -299,6 +299,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.

From 6d3a6168e40ef2d490f708e20baacd0908b67d92 Mon Sep 17 00:00:00 2001
From: ZacharyZcR <zacharyzcr1984@gmail.com>
Date: Sat, 18 Jul 2026 04:14:46 +0800
Subject: [PATCH 20/31] cuda: keep ragged KV resident across decode steps

---
 c/backend_cuda.cu                | 127 ++++++++++++++++++++++++-------
 c/backend_cuda.h                 |   3 +-
 c/backend_loader.c               |   8 +-
 c/glm.c                          |  17 +++--
 c/tests/test_ragged_attention.cu |  10 ++-
 5 files changed, 122 insertions(+), 43 deletions(-)

diff --git a/c/backend_cuda.cu b/c/backend_cuda.cu
index 188330d..413e53b 100644
--- a/c/backend_cuda.cu
+++ b/c/backend_cuda.cu
@@ -8,12 +8,21 @@
 #include <cstring>
 #include <mutex>
 
+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;
@@ -338,17 +347,17 @@ __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 KV sequence per row. latent/rope are packed as [S,T,*], while
- * lengths selects the valid prefix for each row. */
+/* 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 *latent,const float *rope,const int *lengths,
+        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+(size_t)s*T*K,*rs=rope+(size_t)s*T*R;
+    const float *ls=latent[s],*rs=rope[s];
     for(int k=tid;k<K;k+=blockDim.x){float a=0;for(int d=0;d<Q;d++)
         a+=qs[d]*weight_at(weights,fmt,(size_t)(rbase+d)*row_bytes(fmt,K),k)*
           (fmt?wscale[rbase+d]:1.f);qa[k]=a;}
@@ -372,6 +381,15 @@ __global__ static void attention_absorb_ragged_kernel(float *ctx,const float *q,
         ctx[((size_t)s*H+h)*V+v]=a*(fmt?wscale[row]:1.f);}
 }
 
+__global__ static void ragged_kv_append(float *const *latent,float *const *rope,
+        const float *packed,const int *old_len,const int *add,const int *offset,int K,int R){
+    int s=blockIdx.x,n=add[s],base=offset[s];
+    for(int i=threadIdx.x;i<n*(K+R);i+=blockDim.x){
+        if(i<n*K)latent[s][(size_t)old_len[s]*K+i]=packed[base+i];
+        else rope[s][(size_t)old_len[s]*R+i-n*K]=packed[base+i];
+    }
+}
+
 static int reserve(float **ptr, size_t *cap, size_t bytes) {
     if (*cap >= bytes) return 1;
     if (*ptr) cudaFree(*ptr);
@@ -440,16 +458,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;
@@ -805,35 +824,81 @@ extern "C" int coli_cuda_attention_project_batch(ColiCudaTensor *w,ColiCudaTenso
 }
 
 extern "C" int coli_cuda_attention_project_ragged(ColiCudaTensor *w,ColiCudaTensor *proj,
-        float *out,const float *q,const float *const *latent,const float *const *rope,
+        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||!latent||!rope||!lengths||S<1||S>512||T<1||T>512||
+    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;
-    size_t ln=(size_t)S*T*K,rn=(size_t)S*T*R;
-    float *lh=(float*)std::calloc(ln,sizeof(float)),*rh=(float*)std::calloc(rn,sizeof(float));
-    if(!lh||!rh){std::free(lh);std::free(rh);return 0;}
-    for(int s=0;s<S;s++){
-        if(lengths[s]<1||lengths[s]>T){std::free(lh);std::free(rh);return 0;}
-        std::memcpy(lh+(size_t)s*T*K,latent[s],(size_t)lengths[s]*K*sizeof(float));
-        std::memcpy(rh+(size_t)s*T*R,rope[s],(size_t)lengths[s]*R*sizeof(float));
-    }
     DeviceContext *dc=find_ctx(w->device);
-    if(!select_ctx(dc)){std::free(lh);std::free(rh);return 0;}
-    size_t qb=(size_t)S*H*(Q+R)*sizeof(float),lb=ln*sizeof(float),rb=rn*sizeof(float);
+    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;s<S;s++){
+        if(!keys[s]||lengths[s]<1||lengths[s]>T){std::free(dl);std::free(dr);std::free(old);std::free(add);std::free(off);return 0;}
+        RaggedKVEntry *e=nullptr;
+        for(int i=0;i<w->ragged_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]<e->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);
-    int ok=reserve(&dc->aq,&dc->aq_cap,qb)&&reserve(&dc->al,&dc->al_cap,lb)&&
-           reserve(&dc->ar,&dc->ar_cap,rb)&&reserve(&dc->ac,&dc->ac_cap,cb)&&
-           reserve(&dc->y,&dc->y_cap,ob)&&
-           reserve_bytes(&dc->group_desc,&dc->group_desc_cap,(size_t)S*sizeof(int));
+    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;s<S;s++)if(add[s]){
+            float *p=dc->host_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(dc->al,lh,lb,cudaMemcpyHostToDevice,dc->stream),"ragged latent upload")&&
-             cuda_ok(cudaMemcpyAsync(dc->ar,rh,rb,cudaMemcpyHostToDevice,dc->stream),"ragged rope upload")&&
-             cuda_ok(cudaMemcpyAsync(dc->group_desc,lengths,(size_t)S*sizeof(int),cudaMemcpyHostToDevice,dc->stream),"ragged lengths upload");
-    std::free(lh);std::free(rh);if(!ok)return 0;
+             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<<<S,256,0,dc->stream>>>(ddl,ddr,dc->al,dold,dadd,doff,K,R);
+    if(ok)for(int s=0;s<S;s++){
+        for(int i=0;i<w->ragged_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<<<dim3(H,S),256,shared,dc->stream>>>(dc->ac,dc->aq,dc->al,dc->ar,
-        (const int*)dc->group_desc,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
+    attention_absorb_ragged_kernel<<<dim3(H,S),256,shared,dc->stream>>>(dc->ac,dc->aq,ddl,ddr,
+        dn,w->weights,w->scales,w->fmt,S,H,Q,R,V,K,T,scale);
     quant_matmul<<<dim3(proj->O,S),256,0,dc->stream>>>(dc->y,dc->ac,proj->weights,
         proj->scales,proj->fmt,S,proj->I,proj->O,row_bytes(proj->fmt,proj->I));
     return cuda_ok(cudaGetLastError(),"ragged attention launch")&&
@@ -852,6 +917,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;i<tensor->ragged_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 edddbfe..63c1f11 100644
--- a/c/backend_cuda.h
+++ b/c/backend_cuda.h
@@ -94,7 +94,8 @@ COLI_CUDA_DLLEXPORT int coli_cuda_attention_project_batch(ColiCudaTensor *kv_b,C
                                       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 float *const *latent,const float *const *rope,
+        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);
diff --git a/c/backend_loader.c b/c/backend_loader.c
index bbcb8ca..eedbd50 100644
--- a/c/backend_loader.c
+++ b/c/backend_loader.c
@@ -62,7 +62,8 @@ typedef int (*fn_attention_absorb_batch_dev)(ColiCudaTensor *kv_b_shard,float *c
 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 float *const *latent,const float *const *rope,
+        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);
@@ -348,10 +349,11 @@ int coli_cuda_attention_project_batch(ColiCudaTensor *kv_b,ColiCudaTensor *o_pro
 }
 
 int coli_cuda_attention_project_ragged(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj,
-        float *out,const float *q,const float *const *latent,const float *const *rope,
+        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,latent,rope,lengths,
+    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);
 }
 
diff --git a/c/glm.c b/c/glm.c
index 70da320..a45fbc9 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -2717,18 +2717,20 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p
            !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(rl&&rr&&rn){
+            if(rk&&rl&&rr&&rn){
                 for(int s=0;s<S;s++){
                     int pos=positions[s],st0=kvs[s]->kv_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,rl,rr,rn,S,H,c->qk_nope,c->qk_rope,vh,kvl,mt,c->attn_scale);
+                    out,Q,rk,rl,rr,rn,S,H,c->qk_nope,c->qk_rope,vh,kvl,mt,c->attn_scale);
             }
-            free(rl);free(rr);free(rn);
+            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);
@@ -5267,7 +5269,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));
@@ -5323,7 +5325,7 @@ static void run_serve_mux(Model *m, const char *snap){
         }
         active=0; for(int i=0;i<nctx;i++) active+=req[i].active;
         if(!active){ if(eof) break; continue; }
-        DecodeRow rows[16]; int slots[16], S=0;
+        DecodeRow rows[512]; int slots[512], S=0;
         for(int i=0;i<nctx;i++) if(req[i].active){
             rows[S]=(DecodeRow){&ctx[i].kv,req[i].pending,ctx[i].len}; slots[S++]=i;
         }
@@ -6330,8 +6332,9 @@ int main(int argc, char **argv){
     int cap  = argc>1?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"))){
diff --git a/c/tests/test_ragged_attention.cu b/c/tests/test_ragged_attention.cu
index c00a8eb..dd1e2f8 100644
--- a/c/tests/test_ragged_attention.cu
+++ b/c/tests/test_ragged_attention.cu
@@ -16,14 +16,18 @@ int main(){
        !coli_cuda_tensor_upload(&tp,p.data(),nullptr,0,D,O,dev))return 1;
     int n[S]={1,2,3};std::vector<std::vector<float>> l(S),r(S);
     const float *lp[S],*rp[S];
+    const void *keys[S];
     for(int s=0;s<S;s++){
         l[s].resize(n[s]*K);r[s].resize(n[s]*R);
         for(size_t i=0;i<l[s].size();i++)l[s][i]=((int)((i+s*3)%9)-4)*.08f;
         for(size_t i=0;i<r[s].size();i++)r[s][i]=((int)((i+s)%5)-2)*.06f;
-        lp[s]=l[s].data();rp[s]=r[s].data();
+        lp[s]=l[s].data();rp[s]=r[s].data();keys[s]=&l[s];
     }
-    float got[S*O],ref[S*O];
-    if(!coli_cuda_attention_project_ragged(tw,tp,got,q.data(),lp,rp,n,S,H,Q,R,V,K,T,.2f))return 2;
+    float got[S*O],ref[S*O],warm[S*O];
+    int first[S]={1,1,1};
+    if(!coli_cuda_attention_project_ragged(tw,tp,warm,q.data(),keys,lp,rp,first,
+        S,H,Q,R,V,K,1,.2f))return 2;
+    if(!coli_cuda_attention_project_ragged(tw,tp,got,q.data(),keys,lp,rp,n,S,H,Q,R,V,K,T,.2f))return 2;
     for(int s=0;s<S;s++)if(!coli_cuda_attention_project_batch(tw,tp,ref+s*O,
         q.data()+s*H*(Q+R),lp[s],rp[s],1,H,Q,R,V,K,n[s],.2f))return 3;
     double e=0,z=0;for(int i=0;i<S*O;i++){

From 9c17dc6d292095cf56ca2c4cc92712e65a0e1e3a Mon Sep 17 00:00:00 2001
From: ZacharyZcR <zacharyzcr1984@gmail.com>
Date: Sat, 18 Jul 2026 18:38:40 +0800
Subject: [PATCH 21/31] plan NUMA interleave on multi-socket Linux

---
 c/resource_plan.py            | 24 +++++++++++++++++++++++-
 c/tests/test_resource_plan.py | 19 +++++++++++++++++--
 docs/ENVIRONMENT.md           |  2 +-
 docs/benchmarks.md            |  8 ++++++--
 4 files changed, 47 insertions(+), 6 deletions(-)

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/test_resource_plan.py b/c/tests/test_resource_plan.py
index 023cb93..b184c80 100644
--- a/c/tests/test_resource_plan.py
+++ b/c/tests/test_resource_plan.py
@@ -10,6 +10,7 @@ from resource_plan import (
     GB,
     analyze_model,
     build_plan,
+    cpu_socket_count,
     environment_for_plan,
     format_plan,
     memory_available,
@@ -66,15 +67,19 @@ class ResourcePlanTest(unittest.TestCase):
         # 0 slots/layer. The value must be a sane positive number of bytes.
         self.assertGreater(memory_available(), 0)
 
+    def test_cpu_socket_count_is_positive(self):
+        self.assertGreaterEqual(cpu_socket_count(), 1)
+
     def test_builds_bounded_three_tier_plan(self):
         gpus = [{"index": 0, "name": "test-gpu", "total_bytes": 12 * GB,
                  "free_bytes": 10 * GB}]
         plan = build_plan(self.model, ram_gb=16, context=32, vram_gb=20,
                           available_memory=32 * GB, available_disk=100 * GB, gpus=gpus,
-                          physical_cpus=24)
+                          physical_cpus=24, cpu_sockets=2)
         self.assertEqual(plan["version"], 2)
         self.assertEqual(plan["policy"]["name"], "quality")
         self.assertEqual(plan["cpu"]["physical_cores"], 24)
+        self.assertEqual(plan["cpu"]["sockets"], 2)
         self.assertTrue(plan["policy"]["preserve_quantization"])
         self.assertFalse(plan["tiers"]["vram"]["requires_host_backing"])
         self.assertEqual(plan["tiers"]["ram"]["budget_bytes"], 16 * GB)
@@ -105,7 +110,7 @@ class ResourcePlanTest(unittest.TestCase):
             {"index": 1, "name": "b", "total_bytes": 12 * GB, "free_bytes": 10 * GB},
         ]
         plan = build_plan(self.model, ram_gb=16, available_memory=32 * GB,
-                          available_disk=1, gpus=gpus)
+                          available_disk=1, gpus=gpus, cpu_sockets=2)
         env = environment_for_plan(plan, {"RAM_GB": "12", "PIN": "stats.txt",
                                                "COLI_GPUS": "1"})
         self.assertEqual(env["RAM_GB"], "12")
@@ -125,6 +130,16 @@ class ResourcePlanTest(unittest.TestCase):
                                                         "OMP_PROC_BIND": "close"})
         self.assertEqual(explicit_threads["OMP_NUM_THREADS"], "7")
         self.assertEqual(explicit_threads["OMP_PROC_BIND"], "close")
+
+        if sys.platform.startswith("linux"):
+            self.assertEqual(env["COLI_NUMA"], "1")
+            explicit_numa = environment_for_plan(plan, {"COLI_NUMA": "0"})
+            self.assertEqual(explicit_numa["COLI_NUMA"], "0")
+
+    def test_single_socket_plan_does_not_enable_numa(self):
+        plan = build_plan(self.model, available_memory=16 * GB, available_disk=1,
+                          gpus=[], physical_cpus=8, cpu_sockets=1)
+        self.assertNotIn("COLI_NUMA", environment_for_plan(plan))
     def test_cpu_binary_does_not_apply_gpu_tier(self):
         plan = build_plan(self.model, available_memory=16 * GB, available_disk=1,
                           gpus=[{"index": 0, "name": "a", "total_bytes": 8 * GB,
diff --git a/docs/ENVIRONMENT.md b/docs/ENVIRONMENT.md
index 535288a..9f0d749 100644
--- a/docs/ENVIRONMENT.md
+++ b/docs/ENVIRONMENT.md
@@ -43,7 +43,7 @@ Format: `VAR` — default — effect.
 | `URING` | `0` (off) | Linux-only queued expert I/O. `URING=1` implies `PIPE=1`, forces cold reads through io-wq (`IOSQE_ASYNC`), replaces blocking loader pthreads and spin waits with batched SQEs/CQEs, and batches `PILOT_REAL` loads on a separate ring. Use `DIRECT=1` for cold NVMe to avoid page-cache copy/readahead limits. Fails clearly if the kernel denies io_uring; incompatible with `COLI_MMAP=1`. |
 | `DIRECT` | `0` (off) | Use `O_DIRECT`/unbuffered reads for expert slabs. Helps sustained NVMe; keeps the zero-copy GPU path. |
 | `COLI_NO_OMP_TUNE` | off | **Kill-switch** for the OpenMP hot-thread tuning (`OMP_WAIT_POLICY=active` spin + proc-bind). Set `=1` when the CPU is mostly waiting on the GPU (Metal) so spin doesn't steal the shared power budget. |
-| `COLI_NUMA` | off (Linux only) | `COLI_NUMA=1` interleaves expert slabs across NUMA nodes via `mbind` (raw syscall, no libnuma). Helps multi-socket hosts (+7–40% expert matmul); silent no-op on single-node or non-Linux. |
+| `COLI_NUMA` | auto in generated plans on multi-socket Linux; otherwise off | `COLI_NUMA=1` selectively interleaves large expert and dense slabs across NUMA nodes via `mbind` (raw syscall, no libnuma). Helps multi-socket hosts (+7–40% expert matmul); silent no-op on single-node or non-Linux. Explicit `COLI_NUMA=0` overrides the generated plan. |
 | `MLOCK` | `-1` (auto: on for macOS) | Wire the streamed expert cache into physical RAM (`mlock`) to dodge the memory compressor. `0` off, `1` force. |
 | `CAP_RAISE` | `1` (on) | Let the engine raise the expert-cache cap above `topk` when RAM allows (bigger batches). `0` fixes the cap. |
 | `PREFETCH` | `0` | Prefetch depth for streamed experts. |
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
 

From 6a2ab47f3b644a07e7a0841741bf08509dd57075 Mon Sep 17 00:00:00 2001
From: bokiko <bokiko@gmail.com>
Date: Sat, 18 Jul 2026 15:15:26 +0300
Subject: [PATCH 22/31] convert: resolve --mtp/--indexer shards from the local
 index instead of scanning all of them (#383)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

When --indir contains model.safetensors.index.json, the head/indexer
passes convert only the shards that actually hold the requested tensors
(3 instead of scanning 141 for --mtp — every empty scan still opens a
5 GB shard). Prints the selection in the [PLAN] block. Without the index
the full scan is unchanged.

Output is byte-identical to the unfiltered path (sha256-verified on the
GLM-5.2 FP8 head shards, with a decoy shard proving the filter excludes
rather than reorders).

Co-Authored-By: Claude <noreply@anthropic.com>
---
 c/tools/convert_fp8_to_int4.py | 19 +++++++++++++++++++
 1 file changed, 19 insertions(+)

diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py
index 9cd6937..f851c37 100644
--- a/c/tools/convert_fp8_to_int4.py
+++ b/c/tools/convert_fp8_to_int4.py
@@ -389,6 +389,25 @@ 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

From 799fe41f88a963d039af3dd9c56f85c3b26f3e82 Mon Sep 17 00:00:00 2001
From: andres-gb10 <a@nextars.co>
Date: Sat, 18 Jul 2026 16:45:19 +0100
Subject: [PATCH 23/31] coli: honor THINK=1 in run mode

coli run hardcoded the nothink template (<|assistant|><think></think>),
so THINK=1 had no effect in one-shot runs while serve mode honors it
(glm.c serve template picks <think> vs <think></think> from the same
env var). Pick the template the same way here.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---
 c/coli | 7 +++++--
 1 file changed, 5 insertions(+), 2 deletions(-)

diff --git a/c/coli b/c/coli
index 0dd9de3..c7fdaa2 100755
--- a/c/coli
+++ b/c/coli
@@ -480,8 +480,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; <think></think> = risposta diretta (nothink)
-    e=env_for(a); e["PROMPT"]=f"[gMASK]<sop><|user|>{prompt}<|assistant|><think></think>"
+    # template ufficiale GLM-5.2: niente \n dopo i ruoli; <think></think> = risposta diretta (nothink).
+    # THINK=1 lascia <think> aperto, stessa convenzione del serve mode (glm.c). EN: THINK=1 leaves
+    # <think> open so the engine emits its reasoning block; the default stays nothink.
+    tk="<think>" if os.environ.get("THINK","0")=="1" else "<think></think>"
+    e=env_for(a); e["PROMPT"]=f"[gMASK]<sop><|user|>{prompt}<|assistant|>{tk}"
     sys.exit(subprocess.call([GLM, str(a.cap)], env=e))
 
 def server_probe(base, api_key=None, timeout=1.5):

From 308e269f468797211a32884f2a942083ac773aac Mon Sep 17 00:00:00 2001
From: andres-gb10 <a@nextars.co>
Date: Sat, 18 Jul 2026 21:11:40 +0100
Subject: [PATCH 24/31] convert(fp8->int4): --indir copies the full metadata
 set and resumes interrupted passes (#383)

Two gaps in the local path, both raised in #383:

- The main --indir pass copied only config.json while the download path
  copies four files; without tokenizer.json the converted container can't
  run chat/serve. Copy the same four, print what was copied, and warn
  when the source is missing any (tokenizer.json called out explicitly).

- Local passes restarted from shard 0 on every rerun. out-NNNNN names
  count EMITTED shards, not input indexes (shards with no relevant
  tensors emit nothing), so the download path's exists-check can't be
  mirrored directly: a sidecar manifest per prefix records
  input -> output (or empty) plus the conversion parameters, written
  atomically after each shard. Rerun skips what matches and refuses a
  parameter mismatch on the same outdir instead of mixing containers
  (the #355 failure mode).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---
 c/tools/convert_fp8_to_int4.py | 66 +++++++++++++++++++++++++++++-----
 1 file changed, 57 insertions(+), 9 deletions(-)

diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py
index 9cd6937..bef3f74 100644
--- a/c/tools/convert_fp8_to_int4.py
+++ b/c/tools/convert_fp8_to_int4.py
@@ -397,24 +397,72 @@ def main():
         # 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-"
-        n = 0
+        # 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):
+            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)
-                continue
-            save_file(out, os.path.join(a.outdir, f"{prefix}{n:05d}.safetensors"))
-            n += 1
-        # config/tokenizer solo per la conversione principale — i passaggi mtp/idx
-        # vanno nella stessa outdir di un container gia' completo di metadati.
+                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:
-            for fn in ["config.json"]:
+            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)
+                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 {n} {tag} shard(s) -> {a.outdir} ({prefix}NNNNN)")
+        print(f"converted {fresh} {tag} shard(s), {n} in container -> {a.outdir} ({prefix}NNNNN)")
         return
 
     # reale: scarica shard per shard, converte, cancella

From d6676c10e9d78b0f59ced6e51a4709bc2a22b5c5 Mon Sep 17 00:00:00 2001
From: recviking <recreational.viking@gmail.com>
Date: Sat, 18 Jul 2026 16:27:29 -0400
Subject: [PATCH 25/31] =?UTF-8?q?tools:=20repair=5Fmtp=5Fint8.py=20?=
 =?UTF-8?q?=E2=80=94=20fix=20int4-converted=20MTP=20heads=20in=20place;=20?=
 =?UTF-8?q?warn=20on=20--mtp=20--ebits=20<8?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Per-row int4 quantization of the MTP head's eh_proj [D,2D] zeroes its ENTIRE
embedding half: the tensor's two column halves differ in scale by ~20-30x per
row (embedding-half absmax ~0.05, hidden-half ~1.5 on GLM-5.2), so a single
per-row scale (absmax/7) puts every embedding-half weight below half a
quantization step and np.rint rounds it to exact zero (packed bytes 0x88).
The draft head then cannot see the input token and MTP acceptance collapses
to ~0% — the mechanism behind issue #8's measurement (int4: 0-4%; int8:
39-59%), and the reason --mtp already defaults to --ebits 8.

This showed up in the wild: a published pre-converted container had its MTP
shards made with an explicit --ebits 4 and drafts were pure garbage on every
backend (deterministic, data-side; verified byte-for-byte by reproducing the
published bytes with quant_int4 on the official BF16 rows).

- tools/repair_mtp_int8.py: repairs such a container IN PLACE — finds the MTP
  layer's per-row-int4 dense tensors (eh_proj, q/kv/o projections, shared
  experts; routed experts untouched), re-downloads only those from the FP8
  source repo (~355 MB of HTTP range reads, no torch, no token), requantizes
  at int8 with quant_int8's exact math, and rewrites the shards atomically
  with *.bak-int4 backups. --dry-run to inspect. The engine picks up int8
  automatically (qt_from_disk detects format by blob size).
  Validated on GLM-5.2 744B: MTP acceptance 0% -> 100% (greedy), 1.11 -> 3.20
  tokens/forward, decode 0.18 -> 0.30 tok/s (Metal, 4-bit KV).

- convert_fp8_to_int4.py: print a warning when --mtp is combined with
  --ebits <8 and per-row scales (the exact footgun above); --group-size 128
  remains a valid int4 alternative since group scales give the embedding
  half its own scale.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---
 c/tools/convert_fp8_to_int4.py |  11 ++
 c/tools/repair_mtp_int8.py     | 229 +++++++++++++++++++++++++++++++++
 2 files changed, 240 insertions(+)
 create mode 100644 c/tools/repair_mtp_int8.py

diff --git a/c/tools/convert_fp8_to_int4.py b/c/tools/convert_fp8_to_int4.py
index 382125e..e398d6f 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
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.<N>.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.<N>.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("<Q", http_range(self.base + shard, 0, 8))[0]
+            self._hdr[shard] = (json.loads(http_range(self.base + shard, 8, n)), 8 + n)
+        return self._hdr[shard]
+
+    def meta(self, name):
+        shard = self.wmap.get(name)
+        if not shard:
+            return None
+        hdr, _ = self._shard_header(shard)
+        return hdr[name]
+
+    def fetch_f32(self, name):
+        """Download one tensor and dequantize to f32 [O, I] (BF16 or FP8+block scales)."""
+        shard = self.wmap[name]
+        hdr, base = self._shard_header(shard)
+        m = hdr[name]
+        o0, o1 = m["data_offsets"]
+        raw = http_range(self.base + shard, base + o0, o1 - o0)
+        if m["dtype"] == "BF16":
+            u = np.frombuffer(raw, dtype=np.uint16).astype(np.uint32) << 16
+            return u.view(np.float32).reshape(m["shape"]).astype(np.float32)
+        if m["dtype"] == "F8_E4M3":
+            b = np.frombuffer(raw, dtype=np.uint8).astype(np.uint16)
+            sign = np.where(b & 0x80, -1.0, 1.0)
+            e = (b >> 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("<Q", fh.read(8))[0]
+        hdr = json.loads(fh.read(n))
+        base = 8 + n
+        order = sorted(((k, v) for k, v in hdr.items() if k != "__metadata__"),
+                       key=lambda kv: kv[1]["data_offsets"][0])
+        out = {}
+        for k, v in order:
+            fh.seek(base + v["data_offsets"][0])
+            out[k] = (v["dtype"], v["shape"], fh.read(v["data_offsets"][1] - v["data_offsets"][0]))
+    return out, hdr.get("__metadata__")
+
+
+def write_shard(path, tensors, meta):
+    hdr = {}
+    off = 0
+    for k, (dt, shape, raw) in tensors.items():
+        hdr[k] = {"dtype": dt, "shape": shape, "data_offsets": [off, off + len(raw)]}
+        off += len(raw)
+    if meta:
+        hdr["__metadata__"] = meta
+    hj = json.dumps(hdr).encode()
+    hj += b" " * ((8 - len(hj) % 8) % 8)
+    with open(path, "wb") as fh:
+        fh.write(struct.pack("<Q", len(hj)))
+        fh.write(hj)
+        for _, (_, _, raw) in tensors.items():
+            fh.write(raw)
+
+
+def main():
+    ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
+    ap.add_argument("--snap", required=True, help="local colibri int4 container directory")
+    ap.add_argument("--source-repo", default="zai-org/GLM-5.2-FP8")
+    ap.add_argument("--revision", default="main")
+    ap.add_argument("--dry-run", action="store_true", help="report what would change, touch nothing")
+    a = ap.parse_args()
+
+    cfg = json.load(open(os.path.join(a.snap, "config.json")))
+    L = cfg["num_hidden_layers"]
+    pref = f"model.layers.{L}."
+    src = SourceRepo(a.source_repo, a.revision)
+
+    # find the MTP layer's dense int4 tensors across the local shards
+    plan = {}   # shard path -> [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("<Q", fh.read(8))[0]
+            hdr = json.loads(fh.read(n))
+        for name, v in hdr.items():
+            if not name.startswith(pref) or name.endswith(".qs") or v.get("dtype") != "U8":
+                continue
+            if ".mlp.experts." in name:
+                skipped_experts += 1
+                continue
+            m = src.meta(name)
+            if m is None:
+                print(f"  ?? {name}: not in source repo, skipping")
+                continue
+            O, I = m["shape"]
+            nb = v["data_offsets"][1] - v["data_offsets"][0]
+            if nb == O * I:
+                already_ok.append(name)
+            elif nb == O * ((I + 1) // 2):
+                plan.setdefault(f, []).append((name, O, I))
+            else:
+                print(f"  ?? {name}: unexpected blob size {nb} for shape {O}x{I}, skipping")
+
+    n_fix = sum(len(v) for v in plan.values())
+    print(f"MTP layer {L}: {n_fix} dense tensor(s) at per-row int4 to repair, "
+          f"{len(already_ok)} already int8, {skipped_experts} routed-expert tensors left as-is")
+    if not n_fix:
+        print("nothing to do."); return
+    if a.dry_run:
+        for f, names in plan.items():
+            for nm, O, I in names:
+                print(f"  would repair {nm} [{O},{I}] in {os.path.basename(f)}")
+        return
+
+    for f, names in plan.items():
+        print(f"patching {os.path.basename(f)} ({len(names)} tensors)")
+        tensors, meta = read_shard(f)
+        for nm, O, I in names:
+            print(f"  {nm}: fetching source + requantizing at int8...", flush=True)
+            w = src.fetch_f32(nm)
+            assert w.shape == (O, I), f"{nm}: source shape {w.shape} != container {O}x{I}"
+            q, s = quant_int8(w)
+            tensors[nm] = ("U8", [len(q)], q.tobytes())
+            tensors[nm + ".qs"] = ("F32", [O], s.tobytes())
+            # verify first row against the source
+            loc = q[:I].view(np.int8).astype(np.float64) * s[0]
+            ref = w[0].astype(np.float64)
+            cos = float(ref @ loc / (np.linalg.norm(ref) * np.linalg.norm(loc) + 1e-30))
+            print(f"    row-0 cosine vs source: {cos:.5f}")
+        bak = f + ".bak-int4"
+        write_shard(f + ".new", tensors, meta)
+        os.replace(f, bak)
+        os.replace(f + ".new", f)
+        print(f"  saved; original kept as {os.path.basename(bak)}")
+    print("done. Re-run with --dry-run to confirm (should report 'already int8').")
+
+
+if __name__ == "__main__":
+    main()

From 05bba7994c7641c91fe47c0b28d623f1baeb5828 Mon Sep 17 00:00:00 2001
From: ZacharyZcR <zacharyzcr1984@gmail.com>
Date: Sun, 19 Jul 2026 05:01:41 +0800
Subject: [PATCH 26/31] release: version infrastructure + GitHub Release
 workflow

- c/version.py: single source of truth (__version__ = "1.0.0")
- coli: reads version.py, banner shows dynamic version, --version flag
- .github/workflows/release.yml: tag push triggers cross-platform build
  (Linux x86_64, macOS ARM64, Windows x86_64) and creates a GitHub
  Release with packaged binaries + changelog notes
- CHANGELOG.md: v1.0.0 baseline documenting all shipped features

To cut a release:
  1. bump c/version.py
  2. add a CHANGELOG section
  3. git tag v1.0.0 && git push --tags
---
 .github/workflows/release.yml | 99 +++++++++++++++++++++++++++++++++++
 CHANGELOG.md                  | 53 +++++++++++++++++++
 c/coli                        |  5 +-
 c/version.py                  |  3 ++
 4 files changed, 159 insertions(+), 1 deletion(-)
 create mode 100644 .github/workflows/release.yml
 create mode 100644 CHANGELOG.md
 create mode 100644 c/version.py

diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml
new file mode 100644
index 0000000..245d24f
--- /dev/null
+++ b/.github/workflows/release.yml
@@ -0,0 +1,99 @@
+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: ""
+            shell: msys2
+    runs-on: ${{ matrix.os }}
+    defaults:
+      run:
+        shell: ${{ matrix.shell || 'bash' }}
+    steps:
+      - uses: actions/checkout@v4
+
+      - if: matrix.shell == 'msys2'
+        uses: msys2/setup-msys2@v2
+        with:
+          msystem: UCRT64
+          update: false
+          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/coli b/c/coli
index 4c27415..81bc787 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 "",
@@ -712,6 +714,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/version.py b/c/version.py
new file mode 100644
index 0000000..9951ca1
--- /dev/null
+++ b/c/version.py
@@ -0,0 +1,3 @@
+"""Single source of truth for the colibrì version number."""
+
+__version__ = "1.0.0"

From d2d3a7b5598f0feb49ccc8e4665c2b7375d4afb2 Mon Sep 17 00:00:00 2001
From: Stonki13 <204933532+Stonki13@users.noreply.github.com>
Date: Sat, 18 Jul 2026 23:18:16 +0200
Subject: [PATCH 27/31] Fix CUDA detection on Windows:
 cuda_binary()/cuda_linkage() always returned False
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

Both coli's cuda_binary() and doctor.py's cuda_linkage() detect CUDA support
by running `ldd` on the engine binary and looking for a linked libcudart —
but that's Linux-only (cuda_binary() checks sys.platform != "linux", and
cuda_linkage() checks os.name != "posix", both short-circuiting to False on
win32). Windows CUDA_DLL=1 builds never link libcudart at all: glm.exe loads
coli_cuda.dll dynamically via LoadLibrary at startup (backend_loader.c), so
there's no import-table entry for ldd/dumpbin to find in the first place.

The practical effect: `coli doctor` always reported "NVIDIA GPU detected but
the engine is CPU-only" on Windows, and `coli run/chat/serve --gpu ...`
always hard-exited with "--gpu needs the CUDA build" — even on a correctly
built CUDA_DLL=1 binary with coli_cuda.dll sitting right next to glm.exe.

Fix: on win32, detect a COLI_CUDA build by scanning glm.exe for the marker
string "[CUDA] mode: routed experts", which only exists in code compiled
under #ifdef COLI_CUDA (see glm.c's cuda init block), then confirm
coli_cuda.dll actually sits next to the binary — mirroring the Linux
"linked but missing" distinction. Linux/macOS detection is unchanged.

Verified on Windows 11 with mingw-w64 GCC 16.1.0 + CUDA 13.2 + RTX 5080:
`coli doctor` now reports "CUDA engine and devices are available", and
`coli run --gpu 0 --vram 8` populates VRAM (confirmed via the engine's own
"[CUDA] resident set: N tensors, X GB VRAM" runtime log) instead of exiting.
---
 c/coli      | 24 ++++++++++++++++++------
 c/doctor.py | 35 ++++++++++++++++++++++++++---------
 2 files changed, 44 insertions(+), 15 deletions(-)

diff --git a/c/coli b/c/coli
index 4c27415..4566ea6 100755
--- a/c/coli
+++ b/c/coli
@@ -141,12 +141,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))
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, *,

From 2122c004d9bdbaa387c64ae4d6c00f24ace128b1 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sun, 19 Jul 2026 10:56:00 +0200
Subject: [PATCH 28/31] serve: end-to-end tool-calling regression test +
 unparsed-marker diagnosis (#401)

The gateway's tool-calling path had unit coverage (parse_tool_calls,
render_chat) but nothing exercised the real subprocess wire protocol or the
HTTP surface a coding client actually hits. #401 reports plain-text replies
where tool_calls were expected; every documented path checks out, so pin the
whole path down with a mock engine speaking SUBMIT/DATA/DONE and assert:

- non-stream: tool_calls populated, finish_reason tool_calls, no raw markers
- stream: markers suppressed across 20-way chunk splits, tool_calls delta
- tool-result round trip: <|observation|><tool_response> rendering, text reply
- no tools: plain text untouched

Also emit a stderr diagnosis when tools are declared and tool-call markers
are present in the reply but the strict parse matches nothing (typically
quantization-mangled output) pointing at COLI_TOOL_SALVAGE=1 -- the likely
field condition behind #401.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---
 c/openai_server.py               |   8 ++
 c/tests/test_openai_tools_e2e.py | 178 +++++++++++++++++++++++++++++++
 2 files changed, 186 insertions(+)
 create mode 100644 c/tests/test_openai_tools_e2e.py

diff --git a/c/openai_server.py b/c/openai_server.py
index f56e20a..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"</?tool_call>|</?arg_key>|</?arg_value>", 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]
diff --git a/c/tests/test_openai_tools_e2e.py b/c/tests/test_openai_tools_e2e.py
new file mode 100644
index 0000000..3eabee1
--- /dev/null
+++ b/c/tests/test_openai_tools_e2e.py
@@ -0,0 +1,178 @@
+"""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|><tool_response> 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 "<tool_response>" in prompt:
+        reply(rid, "25 degrees and sunny in Rome.")
+    elif "weather in Rome" in prompt:
+        reply(rid, "<tool_call>get_weather<arg_key>city</arg_key>"
+                   "<arg_value>Rome</arg_value></tool_call>")
+    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. <tool_call>get_weather<arg_key>city</arg_key>"
+                   "<arg_value>Milan</arg_value></tool_call>", 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"]}}}]
+
+
+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("<tool_call>", 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("<tool_call>", text)
+        self.assertNotIn("<arg_key>", 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|><tool_response>25 degrees, sunny</tool_response>",
+                      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()

From 0f606bc4465df0473f39e809f9510b318708631e Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sun, 19 Jul 2026 11:01:26 +0200
Subject: [PATCH 29/31] test: skip the tools e2e suite on Windows (shebang mock
 engine)

CreateProcess cannot exec a shebang script, so the gateway exits during
setUpClass on the windows CI job. The gateway logic under test is
platform-independent and stays covered by the POSIX jobs.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---
 c/tests/test_openai_tools_e2e.py | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/c/tests/test_openai_tools_e2e.py b/c/tests/test_openai_tools_e2e.py
index 3eabee1..d921843 100644
--- a/c/tests/test_openai_tools_e2e.py
+++ b/c/tests/test_openai_tools_e2e.py
@@ -68,6 +68,10 @@ TOOLS = [{"type": "function", "function": {
                    "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):

From c90e2cc4382e4a7f7f48914a03a9b0c181c38694 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sun, 19 Jul 2026 11:02:31 +0200
Subject: [PATCH 30/31] =?UTF-8?q?glm:=20measured-RSS=20guard=20=E2=80=94?=
 =?UTF-8?q?=20the=20RAM=20budget=20enforces=20itself=20at=20the=20safe=20p?=
 =?UTF-8?q?oint=20(#403)?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

cap_for_ram's projection is an estimate: on the GB10 (#403) long generations
overshot it by ~40 GB (projected 74.4, real 115.6) and the kernel killed the
engine three times. Run D of the issue proves a low cap CONTAINS the growth;
this guard does that automatically, keyed on MEASURED RSS instead of the
projection.

At the repin safe point (no moe in flight), every ~16 emitted tokens: if RSS
exceeds the resolved budget (RAM_GB/auto, or an explicit RSS_GUARD_GB
ceiling), free the least-used LRU expert slabs in place and lower ecap so the
cache cannot regrow. Slabs are >128 KB so glibc returns the pages to the
kernel immediately -- RSS actually drops.

Eviction never compacts the array: with PILOT_REAL the pilot worker holds
pointers into ecache[] across its preads, so the slot stays in place with
eid=-1/used=0 (first candidate for reuse); reserved slots (eid<0) are never
touched and victim selection happens under g_pilot_mx. resident_bytes is left
alone: LRU slots are never accounted there (only pin + dense).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---
 c/glm.c | 62 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 62 insertions(+)

diff --git a/c/glm.c b/c/glm.c
index 11c09a6..a3e7ab7 100644
--- a/c/glm.c
+++ b/c/glm.c
@@ -4924,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=<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 && freed<need; pass++){
+        for(int l=0; l<=c->n_layers && freed<need; l++){
+            if(!m->ecache || !m->ecache[l]) continue;
+            pthread_mutex_lock(&g_pilot_mx);
+            int nn=m->ecn[l], lru=-1;
+            for(int z=0;z<nn;z++){                        /* solo slot pubblicati e con slab */
+                ESlot *cand=&m->ecache[l][z];
+                if(cand->eid<0 || !cand->slab) continue;
+                if(lru<0 || cand->used<m->ecache[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;
@@ -6024,6 +6085,7 @@ static void cap_for_ram(Model *m, double ram_gb, int ebits, int max_ctx){
     if(auto_b){ ram_gb = g_mem_avail_boot*0.88;   /* misurata PRIMA del load: il residente gia'
                                                    * allocato viene sottratto sotto, non due volte */
         if(ram_gb<4){ fprintf(stderr,"[RAM] MemAvailable is unreadable or too low; assuming 8 GB\n"); ram_gb=8; } }
+    g_ram_budget_gb = ram_gb;                    /* #403: la RSS-guard usa il budget RISOLTO */
     /* slack ONESTO, non forfettario (l'OOM del 2026-07-04 veniva da qui):
      *  ws[64] slab del working-set (si materializzano TUTTI nel prefill batch-union),
      *  KV cache a max_ctx, kvb_all della ricostruzione k/v in attention,

From 3cd2674f68e2015f9b4168841c796a9fcbf32ed6 Mon Sep 17 00:00:00 2001
From: JustVugg <JustVugg@users.noreply.github.com>
Date: Sun, 19 Jul 2026 12:29:33 +0200
Subject: [PATCH 31/31] =?UTF-8?q?release:=20fix=20the=20Windows=20job=20sh?=
 =?UTF-8?q?ell=20=E2=80=94=20'msys2=20{0}',=20not=20'msys2'=20(+=20inherit?=
 =?UTF-8?q?=20PATH=20for=207z)?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

GitHub Actions shells are format strings; the bare 'msys2' string made the
v1.0.0 tag build fail before its first step ('Invalid shell option'). Same
invocation ci.yml already uses. path-type: inherit so the Package step can
reach the runner's 7z.exe.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
---
 .github/workflows/release.yml | 9 +++++++--
 1 file changed, 7 insertions(+), 2 deletions(-)

diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml
index 245d24f..e471fb9 100644
--- a/.github/workflows/release.yml
+++ b/.github/workflows/release.yml
@@ -24,7 +24,9 @@ jobs:
             name: windows-x86_64
             ext: ".exe"
             make_args: ""
-            shell: msys2
+            # 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:
@@ -32,11 +34,14 @@ jobs:
     steps:
       - uses: actions/checkout@v4
 
-      - if: matrix.shell == 'msys2'
+      - 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