Merge remote-tracking branch 'upstream/main' into feat/gpu-backend-hardening
This commit is contained in:
noobdev-ph
@@ -0,0 +1,104 @@
|
||||
name: Release
|
||||
|
||||
on:
|
||||
push:
|
||||
tags: ['v*']
|
||||
|
||||
permissions:
|
||||
contents: write
|
||||
|
||||
jobs:
|
||||
build:
|
||||
strategy:
|
||||
matrix:
|
||||
include:
|
||||
- os: ubuntu-latest
|
||||
name: linux-x86_64
|
||||
ext: ""
|
||||
make_args: "ARCH=x86-64-v3"
|
||||
- os: macos-latest
|
||||
name: macos-arm64
|
||||
ext: ""
|
||||
make_args: ""
|
||||
- os: windows-latest
|
||||
name: windows-x86_64
|
||||
ext: ".exe"
|
||||
make_args: ""
|
||||
# GitHub shells are format strings: 'msys2 {0}', never bare 'msys2'
|
||||
# (the v1.0.0 tag build failed on exactly this). Same as ci.yml.
|
||||
shell: "msys2 {0}"
|
||||
runs-on: ${{ matrix.os }}
|
||||
defaults:
|
||||
run:
|
||||
shell: ${{ matrix.shell || 'bash' }}
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
|
||||
- if: matrix.os == 'windows-latest'
|
||||
uses: msys2/setup-msys2@v2
|
||||
with:
|
||||
msystem: UCRT64
|
||||
update: false
|
||||
# inherit: the Package step calls the runner's 7z.exe, which lives on
|
||||
# the Windows PATH; the default minimal path would not see it.
|
||||
path-type: inherit
|
||||
install: >-
|
||||
make
|
||||
mingw-w64-ucrt-x86_64-gcc
|
||||
|
||||
- if: matrix.os == 'macos-latest'
|
||||
run: brew install libomp
|
||||
|
||||
- name: Build engine
|
||||
run: |
|
||||
cd c
|
||||
make glm ${{ matrix.make_args }}
|
||||
ls -lh glm${{ matrix.ext }}
|
||||
|
||||
- name: Package
|
||||
run: |
|
||||
TAG=${GITHUB_REF#refs/tags/}
|
||||
mkdir -p dist
|
||||
cp c/glm${{ matrix.ext }} dist/colibri-${TAG}-${{ matrix.name }}${{ matrix.ext }}
|
||||
cp c/coli dist/
|
||||
cp c/version.py dist/
|
||||
cp c/openai_server.py dist/
|
||||
cp c/resource_plan.py dist/
|
||||
cp c/doctor.py dist/
|
||||
cp LICENSE dist/
|
||||
cd dist
|
||||
if [ "${{ matrix.ext }}" = ".exe" ]; then
|
||||
7z a colibri-${TAG}-${{ matrix.name }}.zip *
|
||||
else
|
||||
tar czf colibri-${TAG}-${{ matrix.name }}.tar.gz *
|
||||
fi
|
||||
|
||||
- uses: actions/upload-artifact@v4
|
||||
with:
|
||||
name: colibri-${{ matrix.name }}
|
||||
path: dist/colibri-*.*
|
||||
|
||||
release:
|
||||
needs: build
|
||||
runs-on: ubuntu-latest
|
||||
steps:
|
||||
- uses: actions/checkout@v4
|
||||
|
||||
- uses: actions/download-artifact@v4
|
||||
with:
|
||||
path: artifacts
|
||||
merge-multiple: true
|
||||
|
||||
- name: Create GitHub Release
|
||||
env:
|
||||
GH_TOKEN: ${{ github.token }}
|
||||
run: |
|
||||
TAG=${GITHUB_REF#refs/tags/}
|
||||
# Extract the latest version section from CHANGELOG
|
||||
NOTES=$(awk "/^## \\[${TAG#v}\\]/{found=1;next} /^## \\[/{if(found)exit} found{print}" CHANGELOG.md)
|
||||
if [ -z "$NOTES" ]; then
|
||||
NOTES="Release ${TAG}"
|
||||
fi
|
||||
gh release create "$TAG" artifacts/* \
|
||||
--title "colibrì ${TAG}" \
|
||||
--notes "$NOTES"
|
||||
@@ -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
|
||||
+7
-1
@@ -166,7 +166,7 @@ else
|
||||
PYTHON ?= python3
|
||||
endif
|
||||
CUDA_OBJ =
|
||||
TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE)
|
||||
TEST_BINS = tests/test_json$(EXE) tests/test_st$(EXE) tests/test_st_pread$(EXE) tests/test_tier$(EXE) tests/test_grammar$(EXE) tests/test_schema_gbnf$(EXE) tests/test_decode_batch$(EXE) tests/test_idot$(EXE) tests/test_i4_grouped$(EXE) tests/test_stops$(EXE) tests/test_topp$(EXE) tests/test_sample_nan$(EXE) tests/test_kv_alloc$(EXE) tests/test_i4_acc512$(EXE) tests/test_compat_direct$(EXE) tests/test_dsa_select$(EXE) tests/test_logit_nan$(EXE)
|
||||
ifneq (,$(LINUX))
|
||||
TEST_BINS += tests/test_uring$(EXE)
|
||||
endif
|
||||
@@ -328,9 +328,15 @@ tests/test_topp$(EXE): tests/test_topp.c glm.c st.h uring.h json.h tok.h tok_uni
|
||||
tests/bench_topp$(EXE): tests/bench_topp.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h
|
||||
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
|
||||
|
||||
tests/test_sample_nan$(EXE): tests/test_sample_nan.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h
|
||||
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
|
||||
|
||||
tests/test_kv_alloc$(EXE): tests/test_kv_alloc.c glm.c st.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h
|
||||
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
|
||||
|
||||
tests/test_logit_nan$(EXE): tests/test_logit_nan.c glm.c st.h uring.h json.h tok.h tok_unicode.h compat.h grammar.h tier.h
|
||||
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
|
||||
|
||||
tests/test_i4_acc512$(EXE): tests/test_i4_acc512.c
|
||||
$(CC) $(CFLAGS) $< -o $@ $(LDFLAGS)
|
||||
|
||||
|
||||
+143
-3
@@ -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;
|
||||
@@ -340,6 +349,49 @@ __global__ static void attention_absorb_batch_kernel(float *ctx,const float *q,
|
||||
ctx[((size_t)s*H+h)*V+v]=a*(fmt?wscale[row]:1.f);}
|
||||
}
|
||||
|
||||
/* Independent device-resident KV sequence per row. lengths selects the valid
|
||||
* prefix; latent/rope point at paged caches updated by the host wrapper. */
|
||||
__global__ static void attention_absorb_ragged_kernel(float *ctx,const float *q,
|
||||
const float *const *latent,const float *const *rope,const int *lengths,
|
||||
const void *weights,const float *wscale,int fmt,int S,int H,int Q,int R,
|
||||
int V,int K,int T,float scale){
|
||||
int s=blockIdx.y,h=blockIdx.x,tid=threadIdx.x,nt=lengths[s],rbase=h*(Q+V);
|
||||
if(s>=S||nt<1||nt>T)return;
|
||||
extern __shared__ float sm[];float *qa=sm,*cl=qa+K,*scores=cl+K,*red=scores+T;
|
||||
const float *qs=q+((size_t)s*H+h)*(Q+R);
|
||||
const float *ls=latent[s],*rs=rope[s];
|
||||
for(int k=tid;k<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);}
|
||||
}
|
||||
|
||||
__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);
|
||||
@@ -408,16 +460,17 @@ extern "C" void coli_cuda_shutdown(void) {
|
||||
for(int b=0;b<24;b++) if(ctx->pipe_buf[b]) cudaFree(ctx->pipe_buf[b]);
|
||||
if (ctx->host_x) cudaFreeHost(ctx->host_x);
|
||||
if (ctx->host_y) cudaFreeHost(ctx->host_y);
|
||||
if (ctx->host_kv) cudaFreeHost(ctx->host_kv);
|
||||
if (ctx->stream) cudaStreamDestroy(ctx->stream);
|
||||
if (ctx->group_desc) cudaFree(ctx->group_desc);
|
||||
ctx->x = ctx->y = ctx->gate = ctx->up = nullptr;
|
||||
ctx->qx=nullptr; ctx->qscale=nullptr;
|
||||
ctx->aq=ctx->al=ctx->ar=ctx->ac=nullptr;
|
||||
ctx->host_x=ctx->host_y=nullptr;ctx->stream=nullptr;
|
||||
ctx->host_x=ctx->host_y=ctx->host_kv=nullptr;ctx->stream=nullptr;
|
||||
ctx->x_cap = ctx->y_cap = ctx->gate_cap = ctx->up_cap = 0;
|
||||
ctx->qx_cap=ctx->qscale_cap=0;
|
||||
ctx->aq_cap=ctx->al_cap=ctx->ar_cap=ctx->ac_cap=0;
|
||||
ctx->host_x_cap=ctx->host_y_cap=0;
|
||||
ctx->host_x_cap=ctx->host_y_cap=ctx->host_kv_cap=0;
|
||||
ctx->group_desc=nullptr; ctx->group_desc_cap=0;
|
||||
}
|
||||
g_nctx = 0;
|
||||
@@ -795,6 +848,89 @@ extern "C" int coli_cuda_attention_project_batch(ColiCudaTensor *w,ColiCudaTenso
|
||||
return attention_absorb_batch_run(w,proj,out,q,latent,rope,S,H,Q,R,V,K,T,scale);
|
||||
}
|
||||
|
||||
extern "C" int coli_cuda_attention_project_ragged(ColiCudaTensor *w,ColiCudaTensor *proj,
|
||||
float *out,const float *q,const void *const *keys,
|
||||
const float *const *latent,const float *const *rope,
|
||||
const int *lengths,int S,int H,int Q,int R,int V,int K,int T,float scale){
|
||||
if(!w||!proj||!out||!q||!keys||!latent||!rope||!lengths||S<1||S>512||T<1||T>8192||
|
||||
H<1||Q<1||R<1||V<1||K<1||K>512||w->I!=K||w->O!=H*(Q+V)||
|
||||
proj->device!=w->device||proj->I!=H*V)return 0;
|
||||
DeviceContext *dc=find_ctx(w->device);
|
||||
if(!select_ctx(dc))return 0;
|
||||
float **dl=(float**)std::malloc((size_t)S*sizeof(*dl));
|
||||
float **dr=(float**)std::malloc((size_t)S*sizeof(*dr));
|
||||
int *old=(int*)std::malloc((size_t)S*sizeof(*old));
|
||||
int *add=(int*)std::malloc((size_t)S*sizeof(*add));
|
||||
int *off=(int*)std::malloc((size_t)S*sizeof(*off));int packed_n=0;
|
||||
if(!dl||!dr||!old||!add||!off){std::free(dl);std::free(dr);std::free(old);std::free(add);std::free(off);return 0;}
|
||||
for(int s=0;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);
|
||||
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(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,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")&&
|
||||
cuda_ok(cudaMemcpyAsync(out,dc->y,ob,cudaMemcpyDeviceToHost,dc->stream),"ragged output download")&&
|
||||
cuda_ok(cudaStreamSynchronize(dc->stream),"ragged attention synchronize");
|
||||
}
|
||||
|
||||
extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) {
|
||||
if (!tensor) return;
|
||||
DeviceContext *ctx = find_ctx(tensor->device);
|
||||
@@ -806,6 +942,10 @@ extern "C" void coli_cuda_tensor_free(ColiCudaTensor *tensor) {
|
||||
}
|
||||
if (tensor->weights) cudaFree(tensor->weights);
|
||||
if (tensor->scales) cudaFree(tensor->scales);
|
||||
for(int i=0;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);
|
||||
}
|
||||
|
||||
|
||||
+6
-1
@@ -14,6 +14,7 @@
|
||||
#define COLI_CUDA_DLLEXPORT
|
||||
#endif
|
||||
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
@@ -92,6 +93,11 @@ COLI_CUDA_DLLEXPORT int coli_cuda_attention_project_batch(ColiCudaTensor *kv_b,C
|
||||
const float *rope,int S,int H,int Q,int R,
|
||||
int V,int K,int T,float attention_scale);
|
||||
|
||||
COLI_CUDA_DLLEXPORT int coli_cuda_attention_project_ragged(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj,
|
||||
float *out,const float *q,const void *const *keys,
|
||||
const float *const *latent,const float *const *rope,
|
||||
const int *lengths,int S,int H,int Q,int R,int V,int K,int max_t,float attention_scale);
|
||||
|
||||
COLI_CUDA_DLLEXPORT void coli_cuda_tensor_free(ColiCudaTensor *tensor);
|
||||
COLI_CUDA_DLLEXPORT size_t coli_cuda_tensor_bytes(const ColiCudaTensor *tensor);
|
||||
COLI_CUDA_DLLEXPORT int coli_cuda_tensor_device(const ColiCudaTensor *tensor);
|
||||
@@ -143,4 +149,3 @@ COLI_CUDA_DLLEXPORT int coli_cuda_pipe_sync(int device);
|
||||
#endif
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
@@ -61,6 +61,10 @@ typedef int (*fn_attention_absorb_batch)(ColiCudaTensor *kv_b,float *ctx,const f
|
||||
typedef int (*fn_attention_absorb_batch_dev)(ColiCudaTensor *kv_b_shard,float *ctx_dev, const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale);
|
||||
typedef int (*fn_attention_absorb_kvdev)(ColiCudaTensor *kv_b,float *ctx,const float *q, const float *latent_dev,const float *rope_dev,int H,int Q,int R,int V,int K,int T, float scale);
|
||||
typedef int (*fn_attention_project_batch)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q,const float *latent, const float *rope,int S,int H,int Q,int R, int V,int K,int T,float attention_scale);
|
||||
typedef int (*fn_attention_project_ragged)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj,
|
||||
float *out,const float *q,const void *const *keys,
|
||||
const float *const *latent,const float *const *rope,
|
||||
const int *lengths,int S,int H,int Q,int R,int V,int K,int max_t,float attention_scale);
|
||||
typedef int (*fn_attention_project_batch_dev)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale);
|
||||
typedef int (*fn_attention_project_batch_dev_out)(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out_dev,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale);
|
||||
typedef int (*fn_pipe_add)(int device,float *x_dev,const float *t_dev,size_t n);
|
||||
@@ -107,6 +111,7 @@ static struct {
|
||||
fn_attention_absorb_batch_dev attention_absorb_batch_dev;
|
||||
fn_attention_absorb_kvdev attention_absorb_kvdev;
|
||||
fn_attention_project_batch attention_project_batch;
|
||||
fn_attention_project_ragged attention_project_ragged;
|
||||
fn_attention_project_batch_dev attention_project_batch_dev;
|
||||
fn_attention_project_batch_dev_out attention_project_batch_dev_out;
|
||||
fn_pipe_add pipe_add;
|
||||
@@ -200,6 +205,7 @@ static int coli_cuda_load(void){
|
||||
RESOLVE(attention_absorb_batch_dev, fn_attention_absorb_batch_dev)
|
||||
RESOLVE(attention_absorb_kvdev, fn_attention_absorb_kvdev)
|
||||
RESOLVE(attention_project_batch, fn_attention_project_batch)
|
||||
RESOLVE(attention_project_ragged, fn_attention_project_ragged)
|
||||
RESOLVE(attention_project_batch_dev, fn_attention_project_batch_dev)
|
||||
RESOLVE(attention_project_batch_dev_out, fn_attention_project_batch_dev_out)
|
||||
RESOLVE(pipe_add, fn_pipe_add)
|
||||
@@ -342,6 +348,15 @@ int coli_cuda_attention_project_batch(ColiCudaTensor *kv_b,ColiCudaTensor *o_pro
|
||||
return g_cuda.attention_project_batch(kv_b, o_proj, out, q, latent, rope, S, H, Q, R, V, K, T, attention_scale);
|
||||
}
|
||||
|
||||
int coli_cuda_attention_project_ragged(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj,
|
||||
float *out,const float *q,const void *const *keys,
|
||||
const float *const *latent,const float *const *rope,
|
||||
const int *lengths,int S,int H,int Q,int R,int V,int K,int max_t,float attention_scale){
|
||||
if(!coli_cuda_load()) return 0;
|
||||
return g_cuda.attention_project_ragged(kv_b,o_proj,out,q,keys,latent,rope,lengths,
|
||||
S,H,Q,R,V,K,max_t,attention_scale);
|
||||
}
|
||||
|
||||
int coli_cuda_attention_project_batch_dev(ColiCudaTensor *kv_b,ColiCudaTensor *o_proj, float *out,const float *q_dev,const float *latent_dev,const float *rope_dev, int S,int H,int Q,int R,int V,int K,int T,float scale){
|
||||
if(!g_cuda.available){ return 0; }
|
||||
return g_cuda.attention_project_batch_dev(kv_b, o_proj, out, q_dev, latent_dev, rope_dev, S, H, Q, R, V, K, T, scale);
|
||||
|
||||
@@ -40,6 +40,8 @@ if sys.platform == "win32":
|
||||
except (AttributeError, OSError): pass
|
||||
|
||||
HERE = os.path.dirname(os.path.abspath(__file__))
|
||||
sys.path.insert(0, HERE)
|
||||
from version import __version__ as _version
|
||||
|
||||
# Run-in-place (source checkout, "cd c && ./coli ..."): the engine, the
|
||||
# support modules (resource_plan.py, doctor.py, openai_server.py) and
|
||||
@@ -115,7 +117,7 @@ def sprite_lines():
|
||||
def banner(sub=""):
|
||||
sp=sprite_lines()
|
||||
txt=[
|
||||
f"{C.teal}{C.b}colibrì{C.r} {C.dim}v1.0{C.r}",
|
||||
f"{C.teal}{C.b}colibrì{C.r} {C.dim}v{_version}{C.r}",
|
||||
f"{C.dim}tiny engine, immense model{C.r}",
|
||||
f"{C.gray}GLM-5.2 · 744B MoE · int4 · streaming CPU{C.r}",
|
||||
f"{C.dgray}{sub}{C.r}" if sub else "",
|
||||
@@ -141,12 +143,24 @@ def need_model(model):
|
||||
sys.exit(f"{C.yel}engine is not built.{C.r} Run: coli build")
|
||||
|
||||
def cuda_binary():
|
||||
if not os.path.exists(GLM) or sys.platform != "linux": return False
|
||||
try:
|
||||
linked=subprocess.run(["ldd",GLM],capture_output=True,text=True,timeout=3)
|
||||
return any("libcudart" in line and "not found" not in line
|
||||
for line in linked.stdout.splitlines())
|
||||
except (OSError,subprocess.SubprocessError): return False
|
||||
if not os.path.exists(GLM): return False
|
||||
if sys.platform == "linux":
|
||||
try:
|
||||
linked=subprocess.run(["ldd",GLM],capture_output=True,text=True,timeout=3)
|
||||
return any("libcudart" in line and "not found" not in line
|
||||
for line in linked.stdout.splitlines())
|
||||
except (OSError,subprocess.SubprocessError): return False
|
||||
if sys.platform == "win32":
|
||||
# Windows CUDA_DLL=1 builds never link libcudart directly: glm.exe loads
|
||||
# coli_cuda.dll at runtime via LoadLibrary (backend_loader.c), so there's no
|
||||
# import-table entry for ldd/dumpbin to see. Detect the COLI_CUDA build via a
|
||||
# marker string baked into glm.c's #ifdef COLI_CUDA block instead, and require
|
||||
# coli_cuda.dll to actually sit next to glm.exe (else CUDA init fails at startup).
|
||||
try:
|
||||
with open(GLM,"rb") as f: built=b"[CUDA] mode: routed experts" in f.read()
|
||||
except OSError: return False
|
||||
return built and os.path.exists(os.path.join(os.path.dirname(GLM),"coli_cuda.dll"))
|
||||
return False
|
||||
|
||||
def resource_request(a, env):
|
||||
ctx=a.ctx or int(env.get("CTX",4096))
|
||||
@@ -480,8 +494,11 @@ def cmd_run(a):
|
||||
need_model(a.model)
|
||||
prompt=" ".join(a.prompt) if a.prompt else sys.exit('usage: coli run "your prompt"')
|
||||
banner("run")
|
||||
# template ufficiale GLM-5.2: niente \n dopo i ruoli; <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):
|
||||
@@ -890,6 +907,7 @@ def main():
|
||||
common.add_argument("--topp", type=float, default=0); common.add_argument("--topk", type=int, default=0)
|
||||
common.add_argument("--temp", type=float, default=None) # temperatura token (0=greedy, default 1.0+nucleus .95)
|
||||
ap=argparse.ArgumentParser(prog="coli", parents=[common], description="colibrì — run GLM-5.2 locally")
|
||||
ap.add_argument("--version", action="version", version=f"colibrì {_version}")
|
||||
sub=ap.add_subparsers(dest="cmd")
|
||||
sub.add_parser("build", parents=[common]); sub.add_parser("info", parents=[common])
|
||||
pp=sub.add_parser("plan",parents=[common])
|
||||
|
||||
+45
@@ -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>
|
||||
@@ -313,8 +314,52 @@ static inline int compat_setenv(const char *name, const char *value, int overwri
|
||||
}
|
||||
#define setenv(name,value,overwrite) compat_setenv(name,value,overwrite)
|
||||
|
||||
/* --- getenv_utf8: read an env var as UTF-8, not through the ANSI codepage ---
|
||||
* Plain getenv()/_environ are populated by the CRT from the ANSI-codepage view
|
||||
* of the process environment block, not UTF-8. A parent that hands the child a
|
||||
* Unicode value via CreateProcessW's wide env block (e.g. Python's subprocess
|
||||
* module, which coli uses to pass the chat prompt) round-trips correctly only
|
||||
* through GetEnvironmentVariableW; going through narrow getenv() re-encodes it
|
||||
* via CP_ACP first, so any non-ASCII prompt text (Cyrillic, CJK, ...) comes out
|
||||
* corrupted before the byte-level tokenizer ever sees it. Read the wide value
|
||||
* directly and convert straight to UTF-8, bypassing the ANSI codepage entirely.
|
||||
* Returned buffer is intentionally leaked: called a handful of times at
|
||||
* startup, lives for the process. */
|
||||
static inline const char *compat_getenv_utf8(const char *name){
|
||||
wchar_t wname[64];
|
||||
if(MultiByteToWideChar(CP_UTF8, 0, name, -1, wname, 64) <= 0) return getenv(name);
|
||||
DWORD need = GetEnvironmentVariableW(wname, NULL, 0);
|
||||
if(!need) return NULL;
|
||||
wchar_t *wval = (wchar_t*)malloc(need * sizeof(wchar_t));
|
||||
if(!wval) return NULL;
|
||||
GetEnvironmentVariableW(wname, wval, need);
|
||||
int blen = WideCharToMultiByte(CP_UTF8, 0, wval, -1, NULL, 0, NULL, NULL);
|
||||
char *val = blen>0 ? (char*)malloc((size_t)blen) : NULL;
|
||||
if(val) WideCharToMultiByte(CP_UTF8, 0, wval, -1, val, blen, NULL, NULL);
|
||||
free(wval);
|
||||
return val;
|
||||
}
|
||||
#define getenv_utf8(name) compat_getenv_utf8(name)
|
||||
|
||||
/* --- mkdtemp -> _mktemp + _mkdir (POSIX mkdtemp assente su Windows) ---
|
||||
* Test binaries (test_stops.c) create a scratch dir in the CWD via a
|
||||
* "name_XXXXXX" template; POSIX mkdtemp fills the X's and mkdirs 0700. The
|
||||
* Windows CRT has _mktemp (in-place, same XXXXXX contract) so we compose it.
|
||||
* Returns the template pointer on success, NULL on failure — matching POSIX. */
|
||||
static inline char *compat_mkdtemp(char *tmpl){
|
||||
if(!tmpl) return NULL;
|
||||
if(!_mktemp(tmpl)) return NULL; /* fills the trailing X's in place */
|
||||
if(_mkdir(tmpl) != 0) return NULL; /* EEXIST is impossible post-_mktemp */
|
||||
return tmpl;
|
||||
}
|
||||
#define mkdtemp(tmpl) compat_mkdtemp(tmpl)
|
||||
|
||||
#endif /* _WIN32 */
|
||||
|
||||
#ifndef getenv_utf8
|
||||
#define getenv_utf8(name) getenv(name)
|
||||
#endif
|
||||
|
||||
/* --- compat_aligned_free su piattaforme diverse da Windows ---
|
||||
* Su Linux/macOS, posix_memalign usa free() normale. */
|
||||
#ifndef compat_aligned_free
|
||||
|
||||
+26
-9
@@ -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, *,
|
||||
|
||||
@@ -196,12 +196,16 @@ typedef struct {
|
||||
uint64_t mtp_prop, mtp_acc; /* statistica acceptance */
|
||||
int **eroute; int *enr; /* metodo C: routing dell'ULTIMO token per layer */
|
||||
uint64_t eclock, hits, miss, ereq;
|
||||
uint64_t hit_pin, hit_ecache; /* split di hits per tier (#336): pin vs LRU ecache */
|
||||
uint64_t gpu_expert_calls; int gpu_expert_count; int64_t gpu_expert_bytes;
|
||||
uint64_t n_fw, n_emit; /* metodo E: forward di decode / token emessi */
|
||||
uint64_t route_slots, route_swaps; /* CACHE_ROUTE: slots chosen / substituted vs true top-K */
|
||||
uint64_t route_agree_hit, route_agree_tot; /* ROUTE_AGREE: |chosen ∩ true top-K| / K */
|
||||
double route_kl_sum; uint64_t route_kl_n; /* mean KL(true||chosen) on gate mass */
|
||||
double t_ewait, t_emm, t_attn, t_kvb, t_head;/* profiling: dove va il tempo (wall del
|
||||
double t_ewait, t_emm, t_ecpu, t_egpu, t_route, t_p2p, t_attn, t_kvb, t_head;
|
||||
uint64_t n_p2p; /* P0 execution profile: tier split + residual hops */
|
||||
uint64_t cpu_expert_rows; int64_t cpu_expert_bytes;
|
||||
/* profiling: dove va il tempo (wall del
|
||||
* thread di compute; il servizio disco
|
||||
* overlappato vive in g_edisk_ns) */
|
||||
double t_aproj,t_acore,t_aout; /* attention breakdown */
|
||||
@@ -307,15 +311,19 @@ static uint64_t g_prof_nlat; /* forwards recorded (monotonic
|
||||
static void prof_lat(double s){ g_prof_lat[g_prof_nlat++ % PROF_LAT_CAP]=s; }
|
||||
/* snapshot for windowed reports (serve mode: one report per turn) */
|
||||
typedef struct {
|
||||
double edisk,ewait,emm,attn,head;
|
||||
int64_t io; uint64_t hits,miss,ereq,n_fw,n_emit,nlat;
|
||||
double edisk,ewait,emm,ecpu,egpu,route,p2p,attn,head;
|
||||
int64_t io,cpu_bytes; uint64_t hits,miss,ereq,n_fw,n_emit,nlat,n_p2p,cpu_rows;
|
||||
uint64_t hit_pin,hit_ecache;
|
||||
} ProfBase;
|
||||
static void prof_base(Model *m, ProfBase *b){
|
||||
b->edisk=edisk_s(); b->ewait=m->t_ewait; b->emm=m->t_emm;
|
||||
b->ecpu=m->t_ecpu; b->egpu=m->t_egpu; b->route=m->t_route; b->p2p=m->t_p2p;
|
||||
b->attn=m->t_attn; b->head=m->t_head;
|
||||
b->io=atomic_load_explicit(&g_prof_io,memory_order_relaxed);
|
||||
b->hits=m->hits; b->miss=m->miss; b->ereq=m->ereq;
|
||||
b->n_fw=m->n_fw; b->n_emit=m->n_emit; b->nlat=g_prof_nlat;
|
||||
b->hit_pin=m->hit_pin; b->hit_ecache=m->hit_ecache;
|
||||
b->n_fw=m->n_fw; b->n_emit=m->n_emit; b->nlat=g_prof_nlat; b->n_p2p=m->n_p2p;
|
||||
b->cpu_bytes=m->cpu_expert_bytes;b->cpu_rows=m->cpu_expert_rows;
|
||||
}
|
||||
|
||||
static float *falloc(int64_t n){
|
||||
@@ -1649,7 +1657,7 @@ static void model_init(Model *m, const char *snap, int cap, int ebits, int dbits
|
||||
c->index_topk, c->index_topk);
|
||||
}
|
||||
}
|
||||
m->hlast=falloc(D); m->h_all=falloc((int64_t)64*D);
|
||||
m->hlast=falloc(D); m->h_all=falloc((int64_t)512*D);
|
||||
|
||||
/* byte della parte DENSA residente (embed+lm_head+attn+mlp densa+shared+norme) */
|
||||
int64_t rb=qt_bytes(&m->embed)+qt_bytes(&m->lm_head);
|
||||
@@ -2284,6 +2292,18 @@ static void expert_prefetch(Model *m, int layer, int eid){
|
||||
static void qt_addrow(const QT *t, int row, float coef, float *acc){
|
||||
int I=t->I;
|
||||
if(t->fmt==0){ const float *w=t->qf+(int64_t)row*I; for(int i=0;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 +2322,13 @@ static void qt_matvec_rows(const QT *t, int r0, int n, const float *x, float *y)
|
||||
else if(t->fmt==2){ const uint8_t *w=t->q4+(int64_t)row*((I+1)/2); float s=t->s[row]; float acc=0;
|
||||
for(int i=0;i+1<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;
|
||||
@@ -2686,7 +2713,25 @@ static void attention_rows(Model *m, Layer *l, int layer, float *x, int S, int p
|
||||
float *sc_all = falloc((int64_t)omp_get_max_threads()*sc_cap);
|
||||
int cuda_core=0,cuda_projected=0;
|
||||
#ifdef COLI_CUDA
|
||||
if(cuda_absorb&&l->n_kv_b_shard>1){
|
||||
if(kvs&&g_cuda_enabled&&getenv("COLI_CUDA_ATTN")&&atoi(getenv("COLI_CUDA_ATTN"))&&
|
||||
!dnsel&&l->kv_b.cuda_eligible&&l->o.cuda_eligible&&
|
||||
qt_cuda_upload(&l->kv_b)&&qt_cuda_upload(&l->o)){
|
||||
const float **rl=malloc((size_t)S*sizeof(*rl)),**rr=malloc((size_t)S*sizeof(*rr));
|
||||
const void **rk=malloc((size_t)S*sizeof(*rk));
|
||||
int *rn=malloc((size_t)S*sizeof(*rn)); int mt=0;
|
||||
if(rk&&rl&&rr&&rn){
|
||||
for(int s=0;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,rk,rl,rr,rn,S,H,c->qk_nope,c->qk_rope,vh,kvl,mt,c->attn_scale);
|
||||
}
|
||||
free(rk);free(rl);free(rr);free(rn);
|
||||
} else if(cuda_absorb&&l->n_kv_b_shard>1){
|
||||
int n=l->n_kv_b_shard,st0=m->kv_start[layer],nt=pos_base+S-st0,ok=1;
|
||||
float *qs=falloc((int64_t)S*H*qh),*cs=falloc((int64_t)S*H*vh);
|
||||
for(int d=0;d<n;d++)for(int s=0;s<S;s++)memcpy(
|
||||
@@ -2848,6 +2893,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 +3043,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]);
|
||||
@@ -3106,9 +3153,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++; } }
|
||||
}
|
||||
@@ -3263,7 +3310,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
|
||||
coli_cuda_expert_mlp(e->g.cuda,e->u.cuda,e->d.cuda,hh,xg,nr)){
|
||||
for(int r=0;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 +3319,9 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
|
||||
matmul_qt(hh, gg, &e->d, nr);
|
||||
for(int r=0;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;
|
||||
m->cpu_expert_bytes+=qt_bytes(&e->g)+qt_bytes(&e->u)+qt_bytes(&e->d);
|
||||
m->cpu_expert_rows+=(uint64_t)nr;}
|
||||
}
|
||||
#ifdef COLI_CUDA
|
||||
ColiCudaTensor *dev_g[COLI_CUDA_MAX_DEVICES][64],*dev_u[COLI_CUDA_MAX_DEVICES][64];
|
||||
@@ -3280,6 +3329,7 @@ static void moe(Model *m, Layer *l, int layer, float *x, int S, float *out, int
|
||||
int dev_rows[COLI_CUDA_MAX_DEVICES][64],dev_which[COLI_CUDA_MAX_DEVICES][64];
|
||||
int dev_nc[COLI_CUDA_MAX_DEVICES]={0},dev_total[COLI_CUDA_MAX_DEVICES]={0};
|
||||
int dev_off[COLI_CUDA_MAX_DEVICES]={0},dev_ok[COLI_CUDA_MAX_DEVICES]={0};
|
||||
double dev_time[COLI_CUDA_MAX_DEVICES]={0};
|
||||
for(int di=0;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 +3346,28 @@ 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->cpu_expert_bytes+=qt_bytes(&e->g)+qt_bytes(&e->u)+qt_bytes(&e->d);
|
||||
m->cpu_expert_rows+=(uint64_t)nr;}
|
||||
}
|
||||
if(g_prof)m->t_ecpu+=now_s()-tc;
|
||||
}
|
||||
float *src=dev_ok[di]?group_y+(int64_t)off*D:hh;
|
||||
for(int r=0;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 +3375,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 +3983,11 @@ static void layers_forward_rows(Model *m, float *x, int S, int pos_base,
|
||||
float *dst=coli_cuda_pipe_scratch(dev,15,xb);
|
||||
if(dst){
|
||||
if(x_dev_on<0) ok=coli_cuda_pipe_upload(dev,dst,x,xb);
|
||||
else if(x_dev_on!=dev){ ok=coli_cuda_pipe_peer_copy(dev,dst,x_dev_on,x_dev,xb); }
|
||||
else if(x_dev_on!=dev){
|
||||
double tp=g_prof?now_s():0;
|
||||
ok=coli_cuda_pipe_peer_copy(dev,dst,x_dev_on,x_dev,xb);
|
||||
if(g_prof){m->t_p2p+=now_s()-tp;m->n_p2p++;}
|
||||
}
|
||||
else dst=x_dev;
|
||||
if(ok){
|
||||
x_dev=dst; x_dev_on=dev;
|
||||
@@ -4021,7 +4083,7 @@ static float *step_all(Model *m, const int *ids, int S, int pos_base){
|
||||
float *x=falloc((int64_t)S*D);
|
||||
for(int s=0;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);
|
||||
@@ -4034,8 +4096,8 @@ static float *step_all(Model *m, const int *ids, int S, int pos_base){
|
||||
static float *step_decode_batch(Model *m, const DecodeRow *rows, int S){
|
||||
Cfg *c=&m->c; int D=c->hidden;
|
||||
/* Ragged KV currently uses MLA absorption; the stack kernel is sized to 512. */
|
||||
if(!rows || S<1 || S>64 || c->kv_lora>512) return NULL;
|
||||
KVState *kvs[64]; int positions[64];
|
||||
if(!rows || S<1 || S>512 || c->kv_lora>512) return NULL;
|
||||
KVState *kvs[512]; int positions[512];
|
||||
float *x=falloc((int64_t)S*D);
|
||||
for(int s=0;s<S;s++){
|
||||
if(!rows[s].kv || !rows[s].kv->Lc || !rows[s].kv->Rc || !rows[s].kv->kv_start ||
|
||||
@@ -4151,7 +4213,11 @@ static void mtp_absorb(Model *m, const int *next_ids, const float *x, int S, int
|
||||
}
|
||||
|
||||
static inline int argmax_v(const float *lo, int V){
|
||||
int b=0; float bv=lo[0]; for(int i=1;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) ----
|
||||
@@ -4255,9 +4321,28 @@ static void topp_siftdown(int *h, int n, int i){
|
||||
* la coda troncata va AZZERATA in g_pbuf (dist_sample la legge direttamente per id). */
|
||||
static void dist_build(const float *lo, int V){
|
||||
if(!g_pbuf){ g_pbuf=falloc(V); g_pidx=malloc(V*sizeof(int)); }
|
||||
float mx=lo[0]; for(int i=1;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;
|
||||
@@ -4558,6 +4643,11 @@ 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 / %.2f GB/s (%llu row) | routed GPU critical %.3fs | router %.3fs | residual P2P %.3fs / %llu hop | orchestration %.3fs\n",
|
||||
m->t_ecpu,m->t_ecpu>0?m->cpu_expert_bytes/1e9/m->t_ecpu:0.0,
|
||||
(unsigned long long)m->cpu_expert_rows,m->t_egpu,m->t_route,m->t_p2p,(unsigned long long)m->n_p2p,
|
||||
elapsed-m->t_ewait-m->t_emm-m->t_attn-m->t_head-m->t_route-m->t_p2p>0?
|
||||
elapsed-m->t_ewait-m->t_emm-m->t_attn-m->t_head-m->t_route-m->t_p2p:0);
|
||||
#ifdef COLI_METAL
|
||||
if(g_metal_enabled){ uint64_t ok=0,fb=0,ex=0; double su=0,gp=0,sc=0;
|
||||
coli_metal_moe_counts(&ok,&fb,&ex); coli_metal_moe_times(&su,&gp,&sc);
|
||||
@@ -4571,6 +4661,8 @@ static void profile_print(Model *m, double elapsed){
|
||||
|
||||
static void profile_reset(Model *m){
|
||||
m->t_ewait=m->t_emm=m->t_attn=m->t_kvb=m->t_head=0;
|
||||
m->t_ecpu=m->t_egpu=m->t_route=m->t_p2p=0;m->n_p2p=0;
|
||||
m->cpu_expert_bytes=0;m->cpu_expert_rows=0;
|
||||
m->t_aproj=m->t_acore=m->t_aout=0;
|
||||
atomic_store_explicit(&g_edisk_ns,0,memory_order_relaxed);
|
||||
}
|
||||
@@ -4607,19 +4699,31 @@ static void prof_report(Model *m, const ProfBase *b, double elapsed, int tokens,
|
||||
for(int i=0;i<=c->n_layers;i++){ if(m->npin)pinned+=m->npin[i]; if(m->ecn)lru+=m->ecn[i]; }
|
||||
double io_w=m->t_ewait-b->ewait; /* stall the compute thread felt */
|
||||
double io_svc=edisk_s()-b->edisk; /* read service on the loading threads (overlaps compute) */
|
||||
uint64_t dhp=m->hit_pin-b->hit_pin, dhe=m->hit_ecache-b->hit_ecache; /* split #336 */
|
||||
fprintf(f,"[PROF] expert I/O: %.3f GB fetched (%.1f MB/token, %.2f GB/s over the run%s) | "
|
||||
"hit %.1f%% (%llu hit / %llu load) | %.1f loads/token | %.1fs read service / %.1fs felt wait\n",
|
||||
"hit %.1f%% (%llu pin + %llu lru / %llu load) | %.1f loads/token | %.1fs read service / %.1fs felt wait\n",
|
||||
io/1e9, tokens>0?io/1e6/tokens:0.0, io/1e9/elapsed,
|
||||
g_mmap?"; COLI_MMAP=1: page cache may serve part":"",
|
||||
hitp,(unsigned long long)dh,(unsigned long long)dm, tokens>0?(double)dq/tokens:0.0,
|
||||
hitp,(unsigned long long)dhp,(unsigned long long)dhe,(unsigned long long)dm, tokens>0?(double)dq/tokens:0.0,
|
||||
io_svc,io_w);
|
||||
fprintf(f,"[PROF] resident experts: %d pinned (%.1f GB) + %d in LRU (%.1f GB, cap %d/layer)\n",
|
||||
pinned,pinned*eb/1e9,lru,lru*eb/1e9,m->ecap);
|
||||
double emm=m->t_emm-b->emm, attn=m->t_attn-b->attn, head=m->t_head-b->head;
|
||||
double other=elapsed-io_w-emm-attn-head; if(other<0) other=0;
|
||||
double emm=m->t_emm-b->emm, ecpu=m->t_ecpu-b->ecpu, egpu=m->t_egpu-b->egpu;
|
||||
double route=m->t_route-b->route,p2p=m->t_p2p-b->p2p;
|
||||
uint64_t np2p=m->n_p2p-b->n_p2p;
|
||||
int64_t cpu_bytes=m->cpu_expert_bytes-b->cpu_bytes;
|
||||
uint64_t cpu_rows=m->cpu_expert_rows-b->cpu_rows;
|
||||
double attn=m->t_attn-b->attn, head=m->t_head-b->head;
|
||||
double other=elapsed-io_w-emm-attn-head-route-p2p; if(other<0) other=0;
|
||||
double f_io=io_w/elapsed, f_emm=emm/elapsed, f_attn=attn/elapsed;
|
||||
fprintf(f,"[PROF] time shares: expert-I/O %.0f%% | expert-matmul %.0f%% | attention %.0f%% | lm_head %.0f%% | other %.0f%%\n",
|
||||
100*f_io,100*f_emm,100*f_attn,100*head/elapsed,100*other/elapsed);
|
||||
double slow=ecpu>egpu?ecpu:egpu,fast=ecpu<egpu?ecpu:egpu;
|
||||
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,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);
|
||||
if(hitp<90) fprintf(f," More cache is the lever: raise RAM_GB (or add RAM).");
|
||||
@@ -4646,7 +4750,7 @@ static void run_replay(Model *m, const int *full, int nfull, int np){
|
||||
if(np<2||nfull<=np){ fprintf(stderr,"REPLAY requires a non-empty prompt and continuation\n"); return; }
|
||||
kv_alloc(m,nfull+2);
|
||||
float *logit=step(m,full,np-1,0); free(logit);
|
||||
m->hits=m->miss=m->ereq=m->gpu_expert_calls=0;
|
||||
m->hits=m->miss=m->ereq=m->gpu_expert_calls=0; m->hit_pin=m->hit_ecache=0;
|
||||
profile_reset(m);
|
||||
ProfBase pb; prof_base(m,&pb);
|
||||
double t0=now_s(); int steps=0;
|
||||
@@ -4696,7 +4800,7 @@ static void run_text(Model *m, const char *snap, const char *prompt, int ngen){
|
||||
}
|
||||
prefill_t=now_s()-prefill_t;
|
||||
printf("PROFILO PREFILL (%.2fs):\n",prefill_t); profile_print(m,prefill_t);
|
||||
m->hits=m->miss=m->ereq=m->gpu_expert_calls=0;
|
||||
m->hits=m->miss=m->ereq=m->gpu_expert_calls=0; m->hit_pin=m->hit_ecache=0;
|
||||
m->n_emit=m->n_fw=0;
|
||||
g_last_repin=0;
|
||||
profile_reset(m);
|
||||
@@ -4709,8 +4813,9 @@ static void run_text(Model *m, const char *snap, const char *prompt, int ngen){
|
||||
double tot=m->hits+m->miss;
|
||||
int nsp=0; for(int i=0;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,
|
||||
@@ -4819,7 +4924,68 @@ static int repin_pick(Model *m, RepinCand *out, int maxc){
|
||||
}
|
||||
return nb;
|
||||
}
|
||||
/* ---- RSS GUARD (#403) -----------------------------------------------------
|
||||
* La proiezione di cap_for_ram e' una STIMA: sul GB10 (#403) le generazioni
|
||||
* lunghe l'hanno sforata di ~40 GB (proiettato 74.4, reale 115.6 -> 3 kill del
|
||||
* kernel). La run D dell'issue prova che un cap piu' basso CONTIENE la crescita:
|
||||
* questa guardia lo fa da sola, sull'RSS MISURATO invece che sul proiettato.
|
||||
* Al safe point (stessa sede di repin: nessun moe in volo), ogni ~16 token
|
||||
* emessi: se l'RSS supera il budget, svuota gli slot LRU meno usati e abbassa
|
||||
* ecap perche' non ricrescano. Gli slab sono >128KB (mmap'd da glibc): la free
|
||||
* restituisce le pagine al kernel subito, quindi l'RSS scende davvero.
|
||||
* Lo slot NON viene compattato via: resta al suo posto con eid=-1/used=0 (primo
|
||||
* candidato al riuso), perche' con PILOT_REAL il worker tiene puntatori dentro
|
||||
* ecache[] durante i suoi pread e uno spostamento li invaliderebbe; per lo
|
||||
* stesso motivo gli slot eid<0 (riservati/in caricamento) non si toccano e la
|
||||
* selezione avviene sotto g_pilot_mx. resident_bytes resta invariato: gli slot
|
||||
* LRU non sono mai contati li' (solo pin e densa).
|
||||
* EN: evict = free the slab in place (eid=-1, used=0, never compact: PILOT_REAL
|
||||
* EN: holds pointers into ecache[] across its preads), skip eid<0 reservations,
|
||||
* EN: select under g_pilot_mx. RSS_GUARD_GB=<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;
|
||||
@@ -5168,7 +5334,7 @@ static void run_serve_mux(Model *m, const char *snap){
|
||||
g_draft=0; /* one scheduler owns every forward; MTP/speculation is not ragged-safe */
|
||||
int maxctx=getenv("CTX")?atoi(getenv("CTX")):4096;
|
||||
int nctx=getenv("KV_SLOTS")?atoi(getenv("KV_SLOTS")):1;
|
||||
if(nctx<1||nctx>16){fprintf(stderr,"KV_SLOTS deve essere tra 1 e 16\n");exit(2);}
|
||||
if(nctx<1||nctx>512){fprintf(stderr,"KV_SLOTS must be between 1 and 512\n");exit(2);}
|
||||
g_kvsave=getenv("KVSAVE")?atoi(getenv("KVSAVE")):1;
|
||||
KVState *initial=m->kv; free(initial->kv_start); free(initial);
|
||||
ServeCtx *ctx=calloc(nctx,sizeof(*ctx)); ServeReq *req=calloc(nctx,sizeof(*req));
|
||||
@@ -5224,7 +5390,7 @@ static void run_serve_mux(Model *m, const char *snap){
|
||||
}
|
||||
active=0; for(int i=0;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;
|
||||
}
|
||||
@@ -5919,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,
|
||||
@@ -6008,9 +6175,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; }
|
||||
@@ -6231,8 +6398,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"))){
|
||||
@@ -6303,11 +6471,22 @@ int main(int argc, char **argv){
|
||||
#endif
|
||||
}
|
||||
if(getenv("DSA_TOPK")) m.c.index_topk=atoi(getenv("DSA_TOPK")); /* override per test */
|
||||
/* Il path MUX (SERVE_BATCH=1, cioe' `coli serve`) forza g_draft=0 sotto —
|
||||
* la speculazione non e' ragged-safe nel batch multi-slot. Segnalarlo QUI,
|
||||
* altrimenti "MTP active (draft=8)" mentirebbe: il messaggio e' stampato
|
||||
* prima della scelta del path (run_serve_mux, sotto), e con DRAFT=8 diceva
|
||||
* "active" per poi disabilitarlo in silenzio (#358, LordMZTE). */
|
||||
int mux_will_disable_mtp = getenv("SERVE") && getenv("SERVE_BATCH") && atoi(getenv("SERVE_BATCH"));
|
||||
int eff_draft = mux_will_disable_mtp ? 0 : g_draft;
|
||||
printf("loaded in %.2fs | resident dense: %.2f MB | layers=%d experts=%d | MTP %s (draft=%d)\n",
|
||||
now_s()-t0, m.resident_bytes/(1024.0*1024.0), m.c.n_layers, m.c.n_experts,
|
||||
m.has_mtp?"ACTIVE":"absent", g_draft);
|
||||
m.has_mtp?(mux_will_disable_mtp?"DISABLED (multiplexed serve)":"ACTIVE"):"absent", eff_draft);
|
||||
/* anche su stderr: e' il canale che le UI (coli) mostrano all'utente */
|
||||
fprintf(stderr,"[MTP] %s (draft=%d)\n", m.has_mtp?"active: native speculative decoding":"absent", g_draft);
|
||||
if(mux_will_disable_mtp && m.has_mtp)
|
||||
fprintf(stderr,"[MTP] disabled in multiplexed serve (SERVE_BATCH=1): speculation is not "
|
||||
"ragged-safe across KV slots. Single-client interactive use (`coli chat`) keeps MTP.\n");
|
||||
else
|
||||
fprintf(stderr,"[MTP] %s (draft=%d)\n", m.has_mtp?"active: native speculative decoding":"absent", eff_draft);
|
||||
#ifdef __linux__
|
||||
{ /* Only warn for a GENUINE 9p mount (WSL Windows drives, magic 0x01021997), where
|
||||
* fadvise is a no-op. The old check was `snap` starting with "/mnt/", which
|
||||
@@ -6457,8 +6636,9 @@ int main(int argc, char **argv){
|
||||
double tot=m.hits+m.miss;
|
||||
printf("N-gram speculation (DRAFT=%d): %.2f tokens/forward (%llu forwards per %llu tokens)\n",
|
||||
g_draft, m.n_fw?(double)m.n_emit/m.n_fw:1.0, (unsigned long long)m.n_fw, (unsigned long long)m.n_emit);
|
||||
printf("Expert cache hit rate: %.1f%% (hit=%llu miss=%llu) | RSS: %.2f GB | %.1f tok/s\n",
|
||||
tot?100.0*m.hits/tot:0.0, (unsigned long long)m.hits, (unsigned long long)m.miss, rss_gb(), n_new/dt);
|
||||
printf("Expert cache hit rate: %.1f%% (%llu pin + %llu lru / %llu miss) | RSS: %.2f GB | %.1f tok/s\n",
|
||||
tot?100.0*m.hits/tot:0.0, (unsigned long long)m.hit_pin, (unsigned long long)m.hit_ecache,
|
||||
(unsigned long long)m.miss, rss_gb(), n_new/dt);
|
||||
profile_print(&m,dt);
|
||||
if(g_prof) prof_report(&m,&pb,dt,n_new,stdout);
|
||||
#ifdef COLI_CUDA
|
||||
|
||||
+12
-1
@@ -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]
|
||||
@@ -406,7 +414,10 @@ def generation_options(body, limit):
|
||||
maximum = body.get("max_tokens")
|
||||
maximum_param = "max_tokens"
|
||||
if maximum is None:
|
||||
maximum = min(256, limit)
|
||||
# Client omitted max_tokens: honor the operator's configured budget (--max-tokens /
|
||||
# --ngen), not an arbitrary 256 — `coli serve --ngen 32768` must mean 32768 (#382).
|
||||
# Generation still ends at EOS, so this is a cap, not a target.
|
||||
maximum = limit
|
||||
temperature = body.get("temperature")
|
||||
top_p = body.get("top_p")
|
||||
temperature = 0.7 if temperature is None else temperature
|
||||
|
||||
+23
-1
@@ -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"]
|
||||
|
||||
+47
-14
@@ -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;
|
||||
}
|
||||
|
||||
@@ -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 / 123.40 GB/s (456 row) | 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, 123.4, 456.0, 0.15, 0.2, 0.03, 75.0, 0.1])
|
||||
|
||||
|
||||
if __name__ == "__main__":
|
||||
unittest.main()
|
||||
|
||||
@@ -0,0 +1,70 @@
|
||||
/* 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"); }
|
||||
|
||||
/* --- NaN at index 0: poisons the max scan itself (the old mx=lo[0] seed),
|
||||
* the failure mode that starts before the sum (review note on #369) --- */
|
||||
{ float lo[8]={NAN,1.f,0.5f,6.f,0.2f,-1.f,0.f,0.3f}; /* max finite = idx3 (6.0) */
|
||||
dist_build(lo,8);
|
||||
double sum=0; int nan=0;
|
||||
for(int i=0;i<8;i++){ if(!(g_pbuf[i]==g_pbuf[i])) nan=1; sum+=g_pbuf[i]; }
|
||||
assert(!nan && "NaN at lo[0] must not poison via the mx seed");
|
||||
assert(approx1(sum) && "still normalizes to 1");
|
||||
assert(dist_sample(8,-1)==3 && "emits the max finite logit, not token 0"); }
|
||||
|
||||
/* --- all-NaN logits: worst case — must stay finite, no crash --- */
|
||||
{ float lo[8]; for(int i=0;i<8;i++) lo[i]=NAN;
|
||||
dist_build(lo,8);
|
||||
for(int i=0;i<8;i++) assert(g_pbuf[i]==g_pbuf[i] && "all-NaN: g_pbuf stays finite"); }
|
||||
|
||||
/* --- regression: clean logits still produce a valid distribution --- */
|
||||
{ float lo[8]={0.1f,0.2f,3.0f,0.4f,0.5f,0.6f,0.7f,0.8f}; /* peak = idx2 */
|
||||
dist_build(lo,8);
|
||||
double sum=0; int nan=0;
|
||||
for(int i=0;i<8;i++){ if(!(g_pbuf[i]==g_pbuf[i])) nan=1; sum+=g_pbuf[i]; }
|
||||
assert(!nan && "clean softmax stays finite");
|
||||
assert(approx1(sum) && "clean softmax must sum to 1");
|
||||
assert(g_pbuf[2]>=g_pbuf[0] && "peak token keeps the most mass"); }
|
||||
|
||||
printf("OK test_logit_nan: argmax_v NaN-skip + dist_build finite-collapse\n");
|
||||
return 0;
|
||||
}
|
||||
@@ -0,0 +1,182 @@
|
||||
"""End-to-end tool-calling test for the OpenAI gateway (#401).
|
||||
|
||||
Unlike the unit tests in test_openai_server.py (which call parse_tool_calls /
|
||||
render_chat directly), this suite runs openai_server.py as a real subprocess
|
||||
against a mock engine that speaks the actual SERVE wire protocol
|
||||
(READY / SUBMIT / DATA / DONE), then talks to it over real HTTP. It pins down
|
||||
the full path a coding client exercises: tool declaration rendering, marker
|
||||
suppression in streamed deltas (across chunk boundaries), tool_calls in both
|
||||
response shapes, and the <|observation|><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"]}}}]
|
||||
|
||||
|
||||
@unittest.skipUnless(os.name == "posix",
|
||||
"the mock engine is a shebang script the gateway execs directly; "
|
||||
"Windows CreateProcess cannot run it. The gateway logic under test "
|
||||
"is platform-independent and covered by the POSIX CI jobs.")
|
||||
class ToolCallingE2E(unittest.TestCase):
|
||||
@classmethod
|
||||
def setUpClass(cls):
|
||||
cls.tmp = tempfile.TemporaryDirectory()
|
||||
mock = Path(cls.tmp.name) / "mock_engine.py"
|
||||
mock.write_text(MOCK_ENGINE)
|
||||
mock.chmod(0o755)
|
||||
cls.mock_log = Path(cls.tmp.name) / "prompts.log"
|
||||
cls.mock_log.touch()
|
||||
with socket.socket() as probe: # free port, then hand it to the server
|
||||
probe.bind(("127.0.0.1", 0))
|
||||
cls.port = probe.getsockname()[1]
|
||||
env = dict(os.environ, MOCK_LOG=str(cls.mock_log))
|
||||
cls.server = subprocess.Popen(
|
||||
[sys.executable, str(SERVER), "--model", cls.tmp.name,
|
||||
"--engine", str(mock), "--port", str(cls.port)],
|
||||
env=env, stderr=subprocess.DEVNULL)
|
||||
cls.base = f"http://127.0.0.1:{cls.port}/v1"
|
||||
for _ in range(100):
|
||||
try:
|
||||
urllib.request.urlopen(cls.base + "/models", timeout=2)
|
||||
return
|
||||
except OSError:
|
||||
if cls.server.poll() is not None:
|
||||
raise RuntimeError("gateway exited during startup")
|
||||
import time
|
||||
time.sleep(0.1)
|
||||
raise RuntimeError("gateway did not come up")
|
||||
|
||||
@classmethod
|
||||
def tearDownClass(cls):
|
||||
cls.server.terminate()
|
||||
cls.server.wait(timeout=5)
|
||||
cls.tmp.cleanup()
|
||||
|
||||
def post(self, body, stream=False):
|
||||
req = urllib.request.Request(
|
||||
self.base + "/chat/completions", json.dumps(body).encode(),
|
||||
{"Content-Type": "application/json"})
|
||||
resp = urllib.request.urlopen(req, timeout=30)
|
||||
if not stream:
|
||||
return json.loads(resp.read())
|
||||
events = []
|
||||
for raw in resp:
|
||||
line = raw.decode().strip()
|
||||
if line.startswith("data: ") and line != "data: [DONE]":
|
||||
events.append(json.loads(line[6:]))
|
||||
return events
|
||||
|
||||
def test_tool_call_non_stream(self):
|
||||
r = self.post({"model": MODEL_ID, "tools": TOOLS,
|
||||
"messages": [{"role": "user",
|
||||
"content": "What is the weather in Rome?"}]})
|
||||
choice = r["choices"][0]
|
||||
self.assertEqual(choice["finish_reason"], "tool_calls")
|
||||
calls = choice["message"]["tool_calls"]
|
||||
self.assertEqual(len(calls), 1)
|
||||
self.assertEqual(calls[0]["function"]["name"], "get_weather")
|
||||
self.assertEqual(json.loads(calls[0]["function"]["arguments"]), {"city": "Rome"})
|
||||
self.assertNotIn("<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()
|
||||
@@ -0,0 +1,39 @@
|
||||
#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];
|
||||
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();keys[s]=&l[s];
|
||||
}
|
||||
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++){
|
||||
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;
|
||||
}
|
||||
@@ -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,
|
||||
|
||||
@@ -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;
|
||||
}
|
||||
@@ -17,6 +17,12 @@ 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 / ([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_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]]:
|
||||
@@ -30,11 +36,20 @@ def parse_output(stdout: str, stderr: str = "") -> tuple[float, list[float]]:
|
||||
return float(speed.group(1)), [disk] + [float(value) for value in rest]
|
||||
|
||||
|
||||
def execute(engine: str, env: dict[str, str]) -> tuple[float, list[float]]:
|
||||
def parse_p0(stdout: str) -> list[float]:
|
||||
"""Extract the optional PROF=1 execution-layer breakdown."""
|
||||
row = P0_RE.search(stdout)
|
||||
if not row:
|
||||
raise RuntimeError("benchmark output missing P0-EXEC profile")
|
||||
return [float(value) for value in row.groups()]
|
||||
|
||||
|
||||
def execute(engine: str, env: dict[str, str]) -> tuple[float, list[float], list[float]]:
|
||||
run = subprocess.run(
|
||||
[engine, "4", "4", "4"], env=env, text=True, capture_output=True, check=True
|
||||
)
|
||||
return parse_output(run.stdout, run.stderr)
|
||||
speed, profile = parse_output(run.stdout, run.stderr)
|
||||
return speed, profile, parse_p0(run.stdout)
|
||||
|
||||
|
||||
def main() -> None:
|
||||
@@ -63,6 +78,8 @@ def main() -> None:
|
||||
OMP_NUM_THREADS=str(args.threads),
|
||||
OMP_PROC_BIND="spread",
|
||||
OMP_PLACES="cores",
|
||||
DRAFT="0",
|
||||
PROF="1",
|
||||
)
|
||||
|
||||
execute(args.engine, base | {"STATS": str(stats)})
|
||||
@@ -86,13 +103,15 @@ def main() -> None:
|
||||
execute(args.engine, base | extra) # warm-up
|
||||
speeds = {name: [] for name in modes}
|
||||
profiles = {name: [] for name in modes}
|
||||
p0_profiles = {name: [] for name in modes}
|
||||
names = list(modes)
|
||||
for run_index in range(args.runs):
|
||||
order = names[run_index % len(names):] + names[:run_index % len(names)]
|
||||
for name in order:
|
||||
speed, profile = execute(args.engine, base | modes[name])
|
||||
speed, profile, p0 = execute(args.engine, base | modes[name])
|
||||
speeds[name].append(speed)
|
||||
profiles[name].append(profile)
|
||||
p0_profiles[name].append(p0)
|
||||
|
||||
result = {}
|
||||
for name in names:
|
||||
@@ -103,6 +122,10 @@ def main() -> None:
|
||||
key: statistics.median(row[index] for row in profiles[name])
|
||||
for index, key in enumerate(PROFILE_KEYS)
|
||||
},
|
||||
"median_p0": {
|
||||
key: statistics.median(row[index] for row in p0_profiles[name])
|
||||
for index, key in enumerate(P0_KEYS)
|
||||
},
|
||||
}
|
||||
print(json.dumps(result, indent=2))
|
||||
|
||||
|
||||
@@ -286,6 +286,17 @@ def main():
|
||||
# testa MTP a int4 = acceptance ~0-4% (misurato, issue #8): il draft sbaglia sempre
|
||||
# e la speculazione non parte mai. A int8: 39-59%, 2.2-2.8 token/forward.
|
||||
a.ebits = 8 if (a.mtp or a.indexer) else 4
|
||||
if a.mtp and a.ebits < 8 and a.group_size <= 0:
|
||||
# Non solo lossy: eh_proj ha ~20-30x di asimmetria di scala fra le due meta' di
|
||||
# colonna, quindi l'int4 per-riga (UNA scala per riga) arrotonda a ZERO l'intera
|
||||
# meta' embedding -> il draft non vede il token -> acceptance ~0% (issue #8).
|
||||
# EN: not merely lossy: eh_proj has ~20-30x column-scale asymmetry, so per-row
|
||||
# EN: int4 rounds its ENTIRE embedding half to exact zeros -> the draft cannot
|
||||
# EN: see the input token -> ~0% acceptance (issue #8). A container converted
|
||||
# EN: this way is repairable in place with tools/repair_mtp_int8.py.
|
||||
print(f"WARNING: --mtp with --ebits {a.ebits} and per-row scales ZEROES eh_proj's "
|
||||
"embedding half -> MTP acceptance ~0% (issue #8). Use the default --ebits 8, "
|
||||
"or add --group-size 128 for group-scaled int4.")
|
||||
if a.xbits is None: a.xbits = a.ebits
|
||||
|
||||
# Build per-type bits map. If a type-specific arg is set, use it; otherwise the
|
||||
@@ -299,6 +310,16 @@ def main():
|
||||
if bits_map:
|
||||
print(f"[MIXED] precision map: " + ", ".join(f"{k}={v}bit" for k,v in sorted(bits_map.items())))
|
||||
|
||||
# Il PIANO risolto, PRIMA di toccare qualunque cosa (#383): --mtp/--indexer cambiano il
|
||||
# default di ebits a 8 (testa int4 = acceptance ~0%, issue #8) e il ramo grouped e'
|
||||
# gated su bits<=4 — combinazioni sorprendenti devono mostrarsi al secondo 1 di un job
|
||||
# da ore, non nel size-check dopo. EN: print the RESOLVED plan before doing anything.
|
||||
mode = "MTP head only" if a.mtp else "DSA indexer only" if a.indexer else "main model"
|
||||
grp = f"grouped gs={a.group_size} (fmt=4)" if (a.group_size and a.ebits <= 4) else \
|
||||
(f"PER-ROW (grouped branch needs bits<=4; ebits={a.ebits} disables it)" if a.group_size else "per-row")
|
||||
print(f"[PLAN] mode: {mode} | source: {'local ' + a.indir if a.indir else 'download ' + a.repo} | "
|
||||
f"experts {a.ebits}-bit, embed/lm_head {a.io_bits}-bit, x {a.xbits}-bit | {grp}")
|
||||
|
||||
if a.selftest_nvfp4:
|
||||
import torch
|
||||
# 1) LUT e2m1: i 16 codici devono decodificare esattamente ai valori attesi.
|
||||
@@ -379,14 +400,99 @@ def main():
|
||||
if a.indir: # conversione locale (test)
|
||||
shards = sorted(glob.glob(os.path.join(a.indir, "*.safetensors")))
|
||||
from safetensors.numpy import save_file
|
||||
# #383: se l'indice c'e', i passaggi --mtp/--indexer convertono SOLO gli shard
|
||||
# che contengono i tensori richiesti (3 invece di scandire tutti i 141 — ogni
|
||||
# scansione a vuoto apre comunque uno shard da 5 GB). Senza indice: scansione
|
||||
# completa come prima.
|
||||
# EN: #383: when the index is present, the --mtp/--indexer passes convert ONLY
|
||||
# the shards that hold the requested tensors (3 instead of scanning all 141 —
|
||||
# every empty scan still opens a 5 GB shard). Without the index: full scan as
|
||||
# before.
|
||||
if a.mtp or a.indexer:
|
||||
idxp = os.path.join(a.indir, "model.safetensors.index.json")
|
||||
if os.path.exists(idxp):
|
||||
wmap = json.load(open(idxp))["weight_map"]
|
||||
if a.mtp:
|
||||
want = {v for k, v in wmap.items() if k.startswith(f"model.layers.{a.n_layers}.")}
|
||||
else:
|
||||
want = {v for k, v in wmap.items() if "indexer" in k and 0 <= layer_idx(k) < a.n_layers}
|
||||
keep = [sp for sp in shards if os.path.basename(sp) in want]
|
||||
print(f"[PLAN] index: {len(keep)}/{len(shards)} local shard(s) hold the requested tensors")
|
||||
shards = keep
|
||||
# BUG #355: questo ramo ignorava --mtp/--indexer. Con --mtp scriveva
|
||||
# out-NNNNN (gli STESSI nomi di una conversione normale) in ebits=8 e
|
||||
# keep_mtp=False -> il "secondo passaggio MTP" nella stessa outdir
|
||||
# SOVRASCRIVEVA il container gia' finito con una riconversione int8
|
||||
# completa, in silenzio (137/141 shard distrutti prima di accorgersene).
|
||||
# Ora il ramo locale rispecchia il download path: prefisso corretto,
|
||||
# flag passate, shard vuoti saltati.
|
||||
prefix = "out-mtp-" if a.mtp else "out-idx-" if a.indexer else "out-"
|
||||
# RIPRESA (#383): i nomi out-NNNNN contano gli shard EMESSI, non l'indice di
|
||||
# input (gli shard senza tensori rilevanti non producono file), quindi "il
|
||||
# file esiste" non basta per saltare il lavoro gia' fatto. Un manifest
|
||||
# sidecar ricorda input -> output (o "vuoto") e con quali parametri: la
|
||||
# ripresa salta solo cio' che combacia, e parametri diversi sulla stessa
|
||||
# outdir vengono rifiutati invece di mescolare container (il modo #355).
|
||||
# EN: RESUME (#383): out-NNNNN names count EMITTED shards, not the input
|
||||
# EN: index (shards with no relevant tensors emit no file), so "the file
|
||||
# EN: exists" is not enough to skip completed work. A sidecar manifest
|
||||
# EN: records input -> output (or "empty") plus the conversion parameters:
|
||||
# EN: resume skips only what matches, and different parameters on the same
|
||||
# EN: outdir are refused instead of mixing containers (the #355 failure mode).
|
||||
params = {"ebits": a.ebits, "io_bits": a.io_bits, "xbits": a.xbits,
|
||||
"group_size": a.group_size, "n_layers": a.n_layers, "bits_map": bits_map}
|
||||
prog_path = os.path.join(a.outdir, f".{prefix}progress.json")
|
||||
prog = {}
|
||||
if os.path.exists(prog_path):
|
||||
try: prog = json.loads(open(prog_path).read())
|
||||
except (OSError, ValueError): prog = {}
|
||||
if prog and prog.get("params") != params:
|
||||
print(f"ERROR: {prog_path} records a conversion with {prog.get('params')};\n"
|
||||
f" this run uses {params}. Refusing to mix conversions in the same "
|
||||
f"outdir — use a fresh --outdir (or delete the manifest and the "
|
||||
f"{prefix}*.safetensors shards to redo).")
|
||||
return
|
||||
done = prog.setdefault("shards", {}); prog["params"] = params
|
||||
n = 0; fresh = 0; skipped = 0
|
||||
for i, sp in enumerate(shards):
|
||||
out = {}; convert_shard(sp, out, a.n_layers, a.ebits, a.io_bits, a.xbits, group_size=a.group_size, bits_map=bits_map)
|
||||
save_file(out, os.path.join(a.outdir, f"out-{i:05d}.safetensors"))
|
||||
# copia config + tokenizer
|
||||
for fn in ["config.json"]:
|
||||
src = os.path.join(a.indir, fn)
|
||||
if os.path.exists(src): shutil.copy(src, a.outdir)
|
||||
print(f"converted {len(shards)} shards -> {a.outdir}")
|
||||
key = os.path.basename(sp)
|
||||
prev = done.get(key) # None = mai visto; "" = visto, vuoto; nome = emesso
|
||||
if prev is not None and (prev == "" or os.path.exists(os.path.join(a.outdir, prev))):
|
||||
if prev: n += 1
|
||||
skipped += 1
|
||||
continue
|
||||
out = {}
|
||||
convert_shard(sp, out, a.n_layers, a.ebits, a.io_bits, a.xbits,
|
||||
keep_mtp=a.mtp, keep_idx=a.indexer,
|
||||
group_size=a.group_size, bits_map=bits_map)
|
||||
if not out: # shard senza MTP/idx: niente file (come il download path)
|
||||
done[key] = ""
|
||||
else:
|
||||
name = f"{prefix}{n:05d}.safetensors"
|
||||
save_file(out, os.path.join(a.outdir, name))
|
||||
done[key] = name; n += 1; fresh += 1
|
||||
tmp_prog = prog_path + ".tmp" # scrittura atomica: una ripresa non vede mai un manifest mezzo scritto
|
||||
with open(tmp_prog, "w") as f: json.dump(prog, f, indent=1) # EN: atomic write: a resume never sees a half-written manifest
|
||||
os.replace(tmp_prog, prog_path)
|
||||
if skipped: print(f"[RESUME] {skipped} shard(s) already done in {a.outdir}, skipped")
|
||||
# metadati per la conversione principale: gli stessi quattro file del download
|
||||
# path — senza tokenizer.json chat/serve non partono. I passaggi mtp/idx vanno
|
||||
# nella stessa outdir di un container gia' completo di metadati.
|
||||
# EN: metadata for the main pass: the same four files as the download path —
|
||||
# EN: chat/serve won't start without tokenizer.json. The mtp/idx passes target
|
||||
# EN: an outdir whose container already has its metadata.
|
||||
if not a.mtp and not a.indexer:
|
||||
copied, missing = [], []
|
||||
for fn in ["config.json", "tokenizer.json", "tokenizer_config.json", "generation_config.json"]:
|
||||
src = os.path.join(a.indir, fn)
|
||||
if os.path.exists(src): shutil.copy(src, a.outdir); copied.append(fn)
|
||||
else: missing.append(fn)
|
||||
print(f"[META] copied from {a.indir}: {', '.join(copied) if copied else 'nothing'}")
|
||||
if missing:
|
||||
print(f"[META] WARNING: not found in {a.indir}: {', '.join(missing)}"
|
||||
+ (" — chat/serve need tokenizer.json" if "tokenizer.json" in missing else ""))
|
||||
tag = "MTP" if a.mtp else "indexer" if a.indexer else "main"
|
||||
print(f"converted {fresh} {tag} shard(s), {n} in container -> {a.outdir} ({prefix}NNNNN)")
|
||||
return
|
||||
|
||||
# reale: scarica shard per shard, converte, cancella
|
||||
|
||||
@@ -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()
|
||||
@@ -0,0 +1,3 @@
|
||||
"""Single source of truth for the colibrì version number."""
|
||||
|
||||
__version__ = "1.0.0"
|
||||
+1
-1
@@ -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. |
|
||||
|
||||
-100
@@ -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 |
|
||||
+76
@@ -49,6 +49,82 @@ errors before streaming headers are sent. `GET /health` exposes
|
||||
active/queued/completed/rejected counters, and successful generation responses
|
||||
include `x-colibri-queue-wait-ms`.
|
||||
|
||||
## Connect a coding CLI or editor
|
||||
|
||||
The API is OpenAI-compatible, so most coding CLIs and editor extensions work by
|
||||
pointing them at Colibri as an *OpenAI-compatible* provider. Three settings:
|
||||
|
||||
- **Base URL** — `http://localhost:8000/v1`
|
||||
- **Model** — `glm-5.2-colibri` (or whatever you pass to `--model-id`)
|
||||
- **API key** — any non-empty string, e.g. `local`
|
||||
|
||||
Colibri needs **no** API key by default, but many clients refuse to start without
|
||||
one — give them any dummy value. The key is only enforced if you set `COLI_API_KEY`.
|
||||
|
||||
Smoke-test the endpoint first (no key needed unless you set one):
|
||||
|
||||
```bash
|
||||
curl http://127.0.0.1:8000/v1/chat/completions \
|
||||
-H 'Content-Type: application/json' \
|
||||
-d '{"model":"glm-5.2-colibri","messages":[{"role":"user","content":"hi"}]}'
|
||||
```
|
||||
|
||||
**aider**
|
||||
|
||||
```bash
|
||||
export OPENAI_API_BASE=http://localhost:8000/v1
|
||||
export OPENAI_API_KEY=local
|
||||
aider --model openai/glm-5.2-colibri # the openai/ prefix routes to OPENAI_API_BASE
|
||||
```
|
||||
|
||||
**crush** — add a provider to `crush.json` (`~/.config/crush/crush.json`, or
|
||||
`%USERPROFILE%\AppData\Local\crush\crush.json` on Windows):
|
||||
|
||||
```json
|
||||
{
|
||||
"$schema": "https://charm.land/crush.json",
|
||||
"providers": {
|
||||
"colibri": {
|
||||
"name": "Colibri",
|
||||
"type": "openai-compat",
|
||||
"base_url": "http://localhost:8000/v1/",
|
||||
"api_key": "local",
|
||||
"models": [
|
||||
{ "name": "GLM-5.2 (Colibri)", "id": "glm-5.2-colibri",
|
||||
"context_window": 131072, "default_max_tokens": 1024 }
|
||||
]
|
||||
}
|
||||
}
|
||||
}
|
||||
```
|
||||
|
||||
The `"api_key": "local"` dummy is what satisfies clients that demand a key.
|
||||
`context_window` is only the client's budget display — set it to whatever your
|
||||
KV configuration actually allows.
|
||||
|
||||
**Continue, Cline / Roo, `llm`, the OpenAI SDKs, …** — set the provider's base
|
||||
URL to `http://localhost:8000/v1`, the model to `glm-5.2-colibri`, and any dummy
|
||||
key (`OPENAI_API_KEY` / `OPENAI_BASE_URL` for env-based tools).
|
||||
|
||||
> **Set your expectations before connecting an agentic CLI.** Two costs dominate,
|
||||
> and the first one is invisible until you know it's there:
|
||||
>
|
||||
> 1. **Prefill.** Coding agents (crush, aider in repo-map mode, Cline, …) send a
|
||||
> large system prompt plus tool definitions — often 10–20k tokens — *before
|
||||
> your first word*. Prefill on the CPU-streaming path runs at a few tokens per
|
||||
> second (it is attention-bound, see #153), so a 15k-token agent preamble is
|
||||
> **an hour of silent "thinking" before the first output token**. The client
|
||||
> looks hung; it isn't. Smoke-test with the tiny `curl` above first — if that
|
||||
> answers in about a minute, the pipeline works and what you're paying for is
|
||||
> prompt size.
|
||||
> 2. **Decode.** Roughly 1 tok/s for a large model, so multi-turn agent loops
|
||||
> (which re-pay the growing context every turn) compound the cost.
|
||||
>
|
||||
> Practical guidance: single surgical asks with a short context work; iterative
|
||||
> agent sessions against a disk-streaming 744B model do not resemble a hosted
|
||||
> API and mostly won't be worth the wait. If your client lets you trim or disable
|
||||
> its system preamble and tool catalog, do it.
|
||||
|
||||
## Isolated KV contexts
|
||||
|
||||
`coli serve --kv-slots N` allocates up to 16 independent sequence contexts.
|
||||
|
||||
+6
-2
@@ -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
|
||||
|
||||
|
||||
+99
-17
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user